commit fac26cf4ca4abd9142a4ed8939c167a6856f656e
author Alexey Bataev <a.bataev@hotmail.com> Wed May 02 20:03:27 2018 +0000
committer Alexey Bataev <a.bataev@hotmail.com> Wed May 02 20:03:27 2018 +0000
tree 1ae4d74c9cae569a7aba6051551532b412194ee2
parent 2d77119f72fa286dfd1c62904fc8ba9e46983ca8

[OPENMP] Add support for reductions on simd directives in target
regions.

Added codegen for `simd reduction()` constructs in target directives.

llvm-svn: 331393

clang/test/OpenMP/nvptx_target_simd_codegen.cpp

diff --git a/clang/test/OpenMP/nvptx_target_simd_codegen.cpp b/clang/test/OpenMP/nvptx_target_simd_codegen.cpp
index 9bb7617..001eb68 100644
--- a/clang/test/OpenMP/nvptx_target_simd_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_simd_codegen.cpp
@@ -9,9 +9,10 @@
 #define HEADER
 
 // Check that the execution mode of all 2 target regions on the gpu is set to SPMD Mode.
-// CHECK-DAG: {{@__omp_offloading_.+l24}}_exec_mode = weak constant i8 0
-// CHECK-DAG: {{@__omp_offloading_.+l29}}_exec_mode = weak constant i8 0
-// CHECK-DAG: {{@__omp_offloading_.+l34}}_exec_mode = weak constant i8 0
+// CHECK-DAG: {{@__omp_offloading_.+l25}}_exec_mode = weak constant i8 0
+// CHECK-DAG: {{@__omp_offloading_.+l30}}_exec_mode = weak constant i8 0
+// CHECK-DAG: {{@__omp_offloading_.+l35}}_exec_mode = weak constant i8 0
+// CHECK-DAG: {{@__omp_offloading_.+l40}}_exec_mode = weak constant i8 0
 
 #define N 1000
 
@@ -20,14 +21,14 @@
   tx a[N];
   short aa[N];
   tx b[10];
-  
+
   #pragma omp target simd
   for(int i = 0; i < n; i++) {
     a[i] = 1;
   }
 
   #pragma omp target simd
-  for(int i = 0; i < n; i++) {  
+  for (int i = 0; i < n; i++) {
     aa[i] += 1;
   }
 
@@ -36,6 +37,11 @@
     b[i] += 1;
   }
 
+  #pragma omp target simd reduction(+:n)
+  for(int i = 0; i < 10; i++) {
+    b[i] += 1;
+  }
+
   return a[0];
 }
 
@@ -47,7 +53,7 @@
   return a;
 }
 
-// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l24}}(
+// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l25}}(
 // CHECK-DAG: [[THREAD_LIMIT:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
 // CHECK: call void @__kmpc_spmd_kernel_init(i32 [[THREAD_LIMIT]],
 // CHECK-NOT: call void @__kmpc_for_static_init
@@ -55,7 +61,7 @@
 // CHECK: call void @__kmpc_spmd_kernel_deinit()
 // CHECK: ret void
 
-// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l29}}(
+// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l30}}(
 // CHECK-DAG: [[THREAD_LIMIT:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
 // CHECK: call void @__kmpc_spmd_kernel_init(i32 [[THREAD_LIMIT]],
 // CHECK-NOT: call void @__kmpc_for_static_init
@@ -63,7 +69,7 @@
 // CHECK: call void @__kmpc_spmd_kernel_deinit()
 // CHECK: ret void
 
-// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l34}}(
+// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l35}}(
 // CHECK-DAG: [[THREAD_LIMIT:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
 // CHECK: call void @__kmpc_spmd_kernel_init(i32 [[THREAD_LIMIT]],
 // CHECK-NOT: call void @__kmpc_for_static_init
@@ -71,4 +77,16 @@
 // CHECK: call void @__kmpc_spmd_kernel_deinit()
 // CHECK: ret void
 
+// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l40}}(
+// CHECK-DAG: [[THREAD_LIMIT:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 [[THREAD_LIMIT]],
+// CHECK-NOT: call void @__kmpc_for_static_init
+// CHECK-NOT: call void @__kmpc_for_static_fini
+// CHECK: [[RES:%.+]] = call i32 @__kmpc_nvptx_simd_reduce_nowait(i32 %{{.+}}, i32 1, i{{64|32}} {{8|4}}, i8* %{{.+}}, void (i8*, i16, i16, i16)* @{{.+}}, void (i8*, i32)* @{{.+}})
+// CHECK: switch i32 [[RES]]
+// CHECK: call void @__kmpc_nvptx_end_reduce_nowait(i32 %{{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_deinit()
+// CHECK: ret void
+
+
 #endif