[OpenMP] Add an additional test for D34888

[gtbercea]

Summary

Test for checking if the mapping is performed correctly. This is a test initially included in Patch https://reviews.llvm.org/D29905

[tra]

LGTM in general.

test/OpenMP/target_map_codegen.cpp

4845 ↗

(On Diff #115845)

Please break the line to make it easier to read.

[tra]

test/OpenMP/target_map_codegen.cpp

4845 ↗

(On Diff #115845)

Better, but the line still wraps. I'd split it again somewhere between -emit-llvm and -fopenmp.
BTW, you have -fopenmp specified twice. Is that intentional?

[Hahnfeld]

Hi Doru,

if I remember correctly I submitted D34888 for a crash when mapping a scalar value with nested regions.
I've marked another test in this file that the codegen for tofrom is correct. So I don't know if this test checks some other conditions?

Jonas

test/OpenMP/target_map_codegen.cpp

1773 ↗

(On Diff #115947)

tofrom is tested here...

[gtbercea]

In D38040#878090, @Hahnfeld wrote:

Hi Doru,

if I remember correctly I submitted D34888 for a crash when mapping a scalar value with nested regions.
I've marked another test in this file that the codegen for tofrom is correct. So I don't know if this test checks some other conditions?

Jonas

Hi Jonas,

The test is verifying whether the parameter is passed to the kernel correctly. I believe it was not passed as a reference before the patch.

In addition to that, something that was in my previous patch is related to this code:

DSAStack->checkMappableExprComponentListsForDeclAtLevel(
        D, Level, [&](OMPClauseMappableExprCommon::MappableExprComponentListRef

In particular with the Level variable. Should the Level variable actually be Level + 1 in this case?

Thanks,

--Doru

[Hahnfeld]

In D38040#878441, @gtbercea wrote:

The test is verifying whether the parameter is passed to the kernel correctly. I believe it was not passed as a reference before the patch.

Ah, right: This isn't checked anywhere before. Maybe add a comment about what's tested here?
Do we want to check the rest of the codegen with a focus that the variable is passed as a reference?

In addition to that, something that was in my previous patch is related to this code:

DSAStack->checkMappableExprComponentListsForDeclAtLevel(
        D, Level, [&](OMPClauseMappableExprCommon::MappableExprComponentListRef

In particular with the Level variable. Should the Level variable actually be Level + 1 in this case?

I'm not sure, the current public clang-ykt has Level: https://github.com/clang-ykt/clang/blob/d181aed/lib/Sema/SemaOpenMP.cpp#L1361

[gtbercea]

In D38040#878799, @Hahnfeld wrote:

In D38040#878441, @gtbercea wrote:

The test is verifying whether the parameter is passed to the kernel correctly. I believe it was not passed as a reference before the patch.

Ah, right: This isn't checked anywhere before. Maybe add a comment about what's tested here?
Do we want to check the rest of the codegen with a focus that the variable is passed as a reference?

In addition to that, something that was in my previous patch is related to this code:

DSAStack->checkMappableExprComponentListsForDeclAtLevel(
        D, Level, [&](OMPClauseMappableExprCommon::MappableExprComponentListRef

In particular with the Level variable. Should the Level variable actually be Level + 1 in this case?

I'm not sure, the current public clang-ykt has Level: https://github.com/clang-ykt/clang/blob/d181aed/lib/Sema/SemaOpenMP.cpp#L1361

You're right!

[gtbercea]

Fix test.

[Hahnfeld]

LGTM

[gtbercea]

Open

[gtbercea]

Add nocudalib flag.

[gtbercea]

Open

[gtbercea]

Fix test.

test/OpenMP/openmp_offload_codegen.cpp

+// Test device for mapping codegen.
+///==========================================================================///
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -S -emit-llvm %s -o - 2>&1 \
+// RUN:   | FileCheck -check-prefix=CK1 %s
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK1-DEVICE
+
+// expected-no-diagnostics
+
+#ifdef CK1
+
+void target_maps_parallel_integer(int a){
+  int ParamToKernel = a;
+#pragma omp target map(tofrom: ParamToKernel)
+  {
+    ParamToKernel += 1;
+  }
+}
+
+// CK1-DEVICE: {{.*}}void @__omp_offloading_{{.*}}(i32* dereferenceable(4){{.*}}
+
+// CK1: {{.*}}void {{.*}}target_maps_parallel_integer{{.*}} {
+
+// CK1: [[GEPOBP:%.+]] = getelementptr inbounds {{.*}}
+// CK1: [[GEPOBPBIT:%.+]] = bitcast i8** [[GEPOBP]]
+// CK1: store i32* %ParamToKernel, i32** [[GEPOBPBIT]]
+// CK1: [[GEPOP:%.+]] = getelementptr inbounds {{.*}}
+// CK1: [[GEPOPBIT:%.+]] = bitcast i8** [[GEPOP]]
+// CK1: store i32* %ParamToKernel, i32** [[GEPOPBIT]]
+// CK1: [[GEPOBPARG:%.+]] = getelementptr inbounds {{.*}}
+// CK1: [[GEPOPARG:%.+]] = getelementptr inbounds {{.*}}
+// CK1: call {{.*}}tgt_target({{.*}}i8** [[GEPOBPARG]], i8** [[GEPOPARG]]
+
+#endif
 No newline at end of file