diff options
| author | Alexey Bataev <a.bataev@hotmail.com> | 2018-12-04 15:03:25 +0000 |
|---|---|---|
| committer | Alexey Bataev <a.bataev@hotmail.com> | 2018-12-04 15:03:25 +0000 |
| commit | c3028cac2462ec22f695ac673b5980be411abf41 (patch) | |
| tree | bc48c1792ec14e81722775e4d79d7f3addfc88b2 /clang/test/OpenMP | |
| parent | 320cf5dde5c709f5f899aab1d343035fc774ebbc (diff) | |
| download | bcm5719-llvm-c3028cac2462ec22f695ac673b5980be411abf41.tar.gz bcm5719-llvm-c3028cac2462ec22f695ac673b5980be411abf41.zip | |
[OPENMP][NVPTX]Mark __kmpc_barrier functions as convergent.
__kmpc_barrier runtime functions must be marked as convergent to prevent
some dangerous optimizations. Also, for NVPTX target all barriers must
be emitted as simple barriers.
llvm-svn: 348271
Diffstat (limited to 'clang/test/OpenMP')
| -rw-r--r-- | clang/test/OpenMP/nvptx_parallel_codegen.cpp | 10 |
1 files changed, 8 insertions, 2 deletions
diff --git a/clang/test/OpenMP/nvptx_parallel_codegen.cpp b/clang/test/OpenMP/nvptx_parallel_codegen.cpp index 2fd837c92b2..08431fccc00 100644 --- a/clang/test/OpenMP/nvptx_parallel_codegen.cpp +++ b/clang/test/OpenMP/nvptx_parallel_codegen.cpp @@ -45,6 +45,7 @@ tx ftemplate(int n) { #pragma omp parallel if(n>1000) { int a = 45; +#pragma omp barrier } a += 1; aa += 1; @@ -317,10 +318,13 @@ int bar(int n){ // CHECK: define internal void [[PARALLEL_FN4]]( // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]], // CHECK: store i[[SZ]] 45, i[[SZ]]* %a, +// CHECK: call void @__kmpc_barrier(%struct.ident_t* @{{.+}}, i32 %{{.+}}) // CHECK: ret void -// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l54}}_worker() -// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l54}}( +// CHECK: declare void @__kmpc_barrier(%struct.ident_t*, i32) #[[BARRIER_ATTRS:.+]] + +// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l55}}_worker() +// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l55}}( // CHECK-32: [[A_ADDR:%.+]] = alloca i32, // CHECK-64: [[A_ADDR:%.+]] = alloca i64, // CHECK-64: [[CONV:%.+]] = bitcast i64* [[A_ADDR]] to i32* @@ -357,4 +361,6 @@ int bar(int n){ // CHECK: store i32 [[NEW_CC_VAL]], i32* [[CC]], // CHECK: br label +// CHECK: attributes #[[BARRIER_ATTRS]] = {{.*}} convergent {{.*}} + #endif |
