diff options
Diffstat (limited to 'clang/test/OpenMP/nvptx_target_simd_codegen.cpp')
| -rw-r--r-- | clang/test/OpenMP/nvptx_target_simd_codegen.cpp | 34 |
1 files changed, 26 insertions, 8 deletions
diff --git a/clang/test/OpenMP/nvptx_target_simd_codegen.cpp b/clang/test/OpenMP/nvptx_target_simd_codegen.cpp index 9bb76178a46..001eb684b15 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 ftemplate(int n) { 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 @@ tx ftemplate(int n) { 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 @@ int bar(int n){ 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 @@ int bar(int n){ // 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 @@ int bar(int n){ // 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 @@ int bar(int n){ // 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 |
