[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

4559

Please break the line to make it easier to read.

[tra]

test/OpenMP/target_map_codegen.cpp

4559

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

1823

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/target_map_codegen.cpp

void explicit_maps_single (int ii){
   // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
   // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
   // CK19-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
   // CK19-DAG: store i32* [[VAR0]], i32** [[CP0]]
 
   // CK19: call void [[CALL23:@.+]](i32* {{[^,]+}})
   #pragma omp target map(always, tofrom: a)
   {
    a++;
   }
 
   // Multidimensional arrays.
   int marr[4][5][6];
   int ***mptr;
 

 void explicit_maps_member_pointer_references(SSA *sap) {
   double *d;
   SSA sa(d);
   SSB sb(sap);
   sb.foo();
 }
 #endif
 #endif
+
+///==========================================================================///
+// RUN: %clang -DCK30 -std=c++11 -fopenmp -S -emit-llvm -fopenmp-targets=nvptx64-nvidia-cuda \
+// RUN: -nocudalib %s -o - 2>&1 | FileCheck -check-prefix=CK30 %s
+
+#ifdef CK30
+
+void target_maps_parallel_integer(int a){
+  int ParamToKernel = a;
+#pragma omp target map(tofrom: ParamToKernel)
+  {
+    ParamToKernel += 1;
+  }
+}
+
+// CK30: {{.*}}void @__omp_offloading_{{.*}}(i32* dereferenceable(4) %ParamToKernel)
+
+// CK30: {{.*}}void {{.*}}target_maps_parallel_integer{{.*}} {
+
+// CK30: [[GEPOBP:%.+]] = getelementptr inbounds {{.*}}offload_baseptrs
+// CK30: [[GEPOBPBIT:%.+]] = bitcast i8** [[GEPOBP]]
+// CK30: store i32* %ParamToKernel, i32** [[GEPOBPBIT]]
+// CK30: [[GEPOP:%.+]] = getelementptr inbounds {{.*}}offload_ptrs
+// CK30: [[GEPOPBIT:%.+]] = bitcast i8** [[GEPOP]]
+// CK30: store i32* %ParamToKernel, i32** [[GEPOPBIT]]
+// CK30: [[GEPOBPARG:%.+]] = getelementptr inbounds {{.*}}offload_baseptrs
+// CK30: [[GEPOPARG:%.+]] = getelementptr inbounds {{.*}}offload_ptrs
+// CK30: call {{.*}}tgt_target({{.*}}i8** [[GEPOBPARG]], i8** [[GEPOPARG]]
+
+#endif