diff options
Diffstat (limited to 'test/OpenMP/target_is_device_ptr_codegen.cpp')
| -rw-r--r-- | test/OpenMP/target_is_device_ptr_codegen.cpp | 81 |
1 files changed, 38 insertions, 43 deletions
diff --git a/test/OpenMP/target_is_device_ptr_codegen.cpp b/test/OpenMP/target_is_device_ptr_codegen.cpp index 1491f41a787fd..90514acadf815 100644 --- a/test/OpenMP/target_is_device_ptr_codegen.cpp +++ b/test/OpenMP/target_is_device_ptr_codegen.cpp @@ -9,6 +9,14 @@ // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32 + +// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} #ifdef CK1 double *g; @@ -35,7 +43,7 @@ double *g; // CK1: [[SIZES06:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}] // CK1: [[TYPES06:@.+]] = {{.+}}constant [2 x i64] [i64 288, i64 288] -// CK1-LABEL: @_Z3foo +// CK1-LABEL: @_Z3foo{{.*}}( template<typename T> void foo(float *&lr, T *&tr) { float *l; @@ -186,18 +194,32 @@ void bar(float *&a, int *&b) { // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32 + +// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s +// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s +// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s +// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s +// SIMD-ONLY1-NOT: {{__kmpc|__tgt}} #ifdef CK2 // CK2: [[ST:%.+]] = type { double*, double** } +// CK2-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l244.region_id = weak constant i8 0 + // CK2: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}] -// CK2: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33] +// CK2: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 288] + +// CK2-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l259.region_id = weak constant i8 0 -// CK2: [[SIZE01:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}] -// CK2: [[MTYPE01:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 17] +// CK2: [[SIZE01:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] +// CK2: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 288] -// CK2: [[SIZE02:@.+]] = {{.+}}constant [3 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}, i[[sz]] {{8|4}}] -// CK2: [[MTYPE02:@.+]] = {{.+}}constant [3 x i64] [i64 33, i64 0, i64 17] +// CK2-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l274.region_id = weak constant i8 0 + +// CK2: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] +// CK2: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 288] template <typename T> struct ST { @@ -216,66 +238,39 @@ struct ST { // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** - // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double*** + // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]** // CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]] - // CK2-DAG: store double** [[SEC0:%.+]], double*** [[CP0]] - // CK2-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + // CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]] #pragma omp target is_device_ptr(a) { a++; } - // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** - // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double**** + // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]** // CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]] - // CK2-DAG: store double*** [[SEC0:%.+]], double**** [[CP0]] - // CK2-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 - - // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 - // CK2-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double**** - // CK2-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double*** - // CK2-DAG: store double*** [[SEC0]], double**** [[CBP1]] - // CK2-DAG: store double** [[SEC1:%.+]], double*** [[CP1]] - // CK2-DAG: [[SEC1]] = load double**, double*** [[SEC0]] + // CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]] #pragma omp target is_device_ptr(b) { b++; } - // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** - // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double**** + // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]** // CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]] - // CK2-DAG: store double*** [[SEC0:%.+]], double**** [[CP0]] - // CK2-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 - - // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 - // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 - // CK2-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double**** - // CK2-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double*** - // CK2-DAG: store double*** [[SEC0]], double**** [[CBP1]] - // CK2-DAG: store double** [[SEC1:%.+]], double*** [[CP1]] - // CK2-DAG: [[SEC1]] = load double**, double*** [[SEC0]] - - // CK2-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 - // CK2-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK2-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]** - // CK2-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to double*** - // CK2-DAG: store [[ST]]* [[VAR2:%.+]], [[ST]]** [[CBP2]] - // CK2-DAG: store double** [[SEC2:%.+]], double*** [[CP2]] - // CK2-DAG: [[SEC2]] = getelementptr {{.*}}[[ST]]* [[VAR2]], i{{.+}} 0, i{{.+}} 0 + // CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]] #pragma omp target is_device_ptr(a, b) { a++; |