forked from OSchip/llvm-project
[OpenMP] Add support for close map modifier in Clang
Summary: This patch adds support for the close map modifier in Clang. This ensures that the new map type is marked and passed to the OpenMP runtime appropriately. Additional regression tests have been merged from patch D55892 (author @saghir). Reviewers: ABataev, caomhin, jdoerfert, kkwli0 Reviewed By: ABataev Subscribers: kkwli0, Hahnfeld, saghir, guansong, cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D65341 llvm-svn: 368491
This commit is contained in:
Gheorghe-Teodor Bercea
parent
26b2c11451
commit
0fd073b1bf
|
|
@ -7116,6 +7116,9 @@ public:
|
|||
OMP_MAP_LITERAL = 0x100,
|
||||
/// Implicit map
|
||||
OMP_MAP_IMPLICIT = 0x200,
|
||||
/// Close is a hint to the runtime to allocate memory close to
|
||||
/// the target device.
|
||||
OMP_MAP_CLOSE = 0x400,
|
||||
/// The 16 MSBs of the flags indicate whether the entry is member of some
|
||||
/// struct/class.
|
||||
OMP_MAP_MEMBER_OF = 0xffff000000000000,
|
||||
|
|
@ -7296,6 +7299,9 @@ private:
|
|||
if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_always)
|
||||
!= MapModifiers.end())
|
||||
Bits |= OMP_MAP_ALWAYS;
|
||||
if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close)
|
||||
!= MapModifiers.end())
|
||||
Bits |= OMP_MAP_CLOSE;
|
||||
return Bits;
|
||||
}
|
||||
|
||||
|
|
@ -7724,10 +7730,10 @@ private:
|
|||
|
||||
if (!IsExpressionFirstInfo) {
|
||||
// If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
|
||||
// then we reset the TO/FROM/ALWAYS/DELETE flags.
|
||||
// then we reset the TO/FROM/ALWAYS/DELETE/CLOSE flags.
|
||||
if (IsPointer)
|
||||
Flags &= ~(OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_ALWAYS |
|
||||
OMP_MAP_DELETE);
|
||||
OMP_MAP_DELETE | OMP_MAP_CLOSE);
|
||||
|
||||
if (ShouldBeMemberOf) {
|
||||
// Set placeholder value MEMBER_OF=FFFF to indicate that the flag
|
||||
|
|
|
|||
|
|
@ -40,6 +40,10 @@ double gc[100];
|
|||
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
|
||||
// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
|
||||
|
||||
// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1057]
|
||||
|
||||
// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1061]
|
||||
|
||||
// CK1-LABEL: _Z3fooi
|
||||
void foo(int arg) {
|
||||
int la;
|
||||
|
|
@ -163,6 +167,64 @@ void foo(int arg) {
|
|||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
|
||||
#pragma omp target data map(to: gb.b[:3])
|
||||
{++arg;}
|
||||
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
{++arg;}
|
||||
|
||||
// Region 05
|
||||
// CK1-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}})
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
||||
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
|
||||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
|
||||
// CK1-DAG: store float* [[VAR0]], float** [[CP0]]
|
||||
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
|
||||
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
|
||||
// CK1-DAG: call void @__tgt_target_data_end(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}})
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
|
||||
#pragma omp target data map(close, to: lb)
|
||||
{++arg;}
|
||||
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
{++arg;}
|
||||
|
||||
// Region 06
|
||||
// CK1-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}})
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
||||
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
|
||||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
|
||||
// CK1-DAG: store float* [[VAR0]], float** [[CP0]]
|
||||
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
|
||||
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
|
||||
// CK1-DAG: call void @__tgt_target_data_end(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}})
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
|
||||
#pragma omp target data map(always close, to: lb)
|
||||
{++arg;}
|
||||
}
|
||||
#endif
|
||||
///==========================================================================///
|
||||
|
|
@ -283,4 +345,150 @@ void no_target_devices(int arg) {
|
|||
{++arg;}
|
||||
}
|
||||
#endif
|
||||
///==========================================================================///
|
||||
// RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
|
||||
// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
|
||||
// RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-32
|
||||
// RUN: %clang_cc1 -DCK4 -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 CK4 --check-prefix CK4-32
|
||||
|
||||
// RUN: %clang_cc1 -DCK4 -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 -DCK4 -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 -DCK4 -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 -DCK4 -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 CK4
|
||||
|
||||
// CK4: [[STT:%.+]] = type { i32, double* }
|
||||
template <typename T>
|
||||
struct STT {
|
||||
T a;
|
||||
double *b;
|
||||
|
||||
T foo(T arg) {
|
||||
// Region 00
|
||||
#pragma omp target data map(always, close to: b[1:3]) if(a>123) device(arg)
|
||||
{arg++;}
|
||||
return arg;
|
||||
}
|
||||
};
|
||||
|
||||
// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976711701]
|
||||
|
||||
// CK4-LABEL: _Z3bari
|
||||
int bar(int arg){
|
||||
STT<int> A;
|
||||
return A.foo(arg);
|
||||
}
|
||||
|
||||
// Region 00
|
||||
// CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
|
||||
// CK4: [[IFTHEN]]
|
||||
// CK4-DAG: call void @__tgt_target_data_begin(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}})
|
||||
// CK4-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
|
||||
// CK4-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
|
||||
// CK4-DAG: [[GEPBP]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%[^,]+]]
|
||||
// CK4-DAG: [[GEPP]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%[^,]+]]
|
||||
// CK4-DAG: [[GEPS]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[S:%[^,]+]]
|
||||
|
||||
// CK4-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK4-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK4-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK4-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[STT]]**
|
||||
// CK4-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
|
||||
// CK4-DAG: store [[STT]]* [[VAR0:%.+]], [[STT]]** [[CBP0]]
|
||||
// CK4-DAG: store double** [[SEC0:%.+]], double*** [[CP0]]
|
||||
// CK4-DAG: store i[[sz]] {{%.+}}, i[[sz]]* [[S0]]
|
||||
// CK4-DAG: [[SEC0]] = getelementptr inbounds {{.*}}[[STT]]* [[VAR0]], i32 0, i32 1
|
||||
|
||||
// CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
|
||||
// CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
|
||||
// CK4-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
|
||||
// CK4-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double***
|
||||
// CK4-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
|
||||
// CK4-DAG: store double** [[SEC0]], double*** [[CBP1]]
|
||||
// CK4-DAG: store double* [[SEC1:%.+]], double** [[CP1]]
|
||||
// CK4-DAG: store i[[sz]] 24, i[[sz]]* [[S1]]
|
||||
// CK4-DAG: [[SEC1]] = getelementptr inbounds {{.*}}double* [[SEC11:%[^,]+]], i{{.+}} 1
|
||||
// CK4-DAG: [[SEC11]] = load double*, double** [[SEC111:%[^,]+]],
|
||||
// CK4-DAG: [[SEC111]] = getelementptr inbounds {{.*}}[[STT]]* [[VAR0]], i32 0, i32 1
|
||||
|
||||
// CK4: br label %[[IFEND:[^,]+]]
|
||||
|
||||
// CK4: [[IFELSE]]
|
||||
// CK4: br label %[[IFEND]]
|
||||
// CK4: [[IFEND]]
|
||||
// CK4: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
// CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
|
||||
|
||||
// CK4: [[IFTHEN]]
|
||||
// CK4-DAG: call void @__tgt_target_data_end(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}})
|
||||
// CK4-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
|
||||
// CK4-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
|
||||
// CK4-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
|
||||
// CK4-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
|
||||
// CK4-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
|
||||
// CK4: br label %[[IFEND:[^,]+]]
|
||||
// CK4: [[IFELSE]]
|
||||
// CK4: br label %[[IFEND]]
|
||||
// CK4: [[IFEND]]
|
||||
#endif
|
||||
///==========================================================================///
|
||||
// RUN: %clang_cc1 -DCK5 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
|
||||
// RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
|
||||
// RUN: %clang_cc1 -DCK5 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-32
|
||||
// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-32
|
||||
|
||||
// RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
|
||||
// RUN: %clang_cc1 -DCK5 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
|
||||
// RUN: %clang_cc1 -DCK5 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
|
||||
// RUN: %clang_cc1 -DCK5 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
|
||||
// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}#ifdef CK5
|
||||
#ifdef CK5
|
||||
struct S1 {
|
||||
int i;
|
||||
};
|
||||
struct S2 {
|
||||
S1 s;
|
||||
struct S2 *ps;
|
||||
};
|
||||
|
||||
void test_close_modifier(int arg) {
|
||||
S2 *ps;
|
||||
// CK5: private unnamed_addr constant [5 x i64] [i64 1059, i64 32, {{.*}}, i64 16, i64 1043]
|
||||
#pragma omp target data map(close,tofrom: arg, ps->ps->ps->ps->s)
|
||||
{
|
||||
++(arg);
|
||||
}
|
||||
}
|
||||
#endif
|
||||
///==========================================================================///
|
||||
// RUN: %clang_cc1 -DCK6 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK6 --check-prefix CK6-64
|
||||
// RUN: %clang_cc1 -DCK6 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK6 --check-prefix CK6-64
|
||||
// RUN: %clang_cc1 -DCK6 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK6 --check-prefix CK6-32
|
||||
// RUN: %clang_cc1 -DCK6 -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 CK6 --check-prefix CK6-32
|
||||
|
||||
// RUN: %clang_cc1 -DCK6 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
|
||||
// RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
|
||||
// RUN: %clang_cc1 -DCK6 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
|
||||
// RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
|
||||
// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
|
||||
#ifdef CK6
|
||||
void test_close_modifier(int arg) {
|
||||
// CK6: private unnamed_addr constant [1 x i64] [i64 1059]
|
||||
#pragma omp target data map(close,tofrom: arg)
|
||||
{++arg;}
|
||||
}
|
||||
#endif
|
||||
#endif
|
||||
|
|
|
|||
|
|
@ -40,6 +40,10 @@ double gc[100];
|
|||
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
|
||||
// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
|
||||
|
||||
// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1057]
|
||||
|
||||
// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1061]
|
||||
|
||||
// CK1-LABEL: _Z3fooi
|
||||
void foo(int arg) {
|
||||
int la;
|
||||
|
|
@ -146,6 +150,56 @@ void foo(int arg) {
|
|||
// CK1-NOT: __tgt_target_data_end
|
||||
#pragma omp target enter data map(to: gb.b[:3])
|
||||
{++arg;}
|
||||
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
{++arg;}
|
||||
|
||||
// Region 05
|
||||
// CK1-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}})
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
||||
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
|
||||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
|
||||
// CK1-DAG: store float* [[VAR0]], float** [[CP0]]
|
||||
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
|
||||
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
// CK1-NOT: __tgt_target_data_end
|
||||
#pragma omp target enter data map(close, to: lb)
|
||||
{++arg;}
|
||||
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
{++arg;}
|
||||
|
||||
// Region 06
|
||||
// CK1-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}})
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
||||
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
|
||||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
|
||||
// CK1-DAG: store float* [[VAR0]], float** [[CP0]]
|
||||
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
|
||||
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
// CK1-NOT: __tgt_target_data_end
|
||||
#pragma omp target enter data map(always close, to: lb)
|
||||
{++arg;}
|
||||
}
|
||||
#endif
|
||||
///==========================================================================///
|
||||
|
|
@ -298,4 +352,81 @@ void device_side_scan(int arg) {
|
|||
{++arg;}
|
||||
}
|
||||
#endif
|
||||
///==========================================================================///
|
||||
// RUN: %clang_cc1 -DCK5 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
|
||||
// RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
|
||||
// RUN: %clang_cc1 -DCK5 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-32
|
||||
// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-32
|
||||
|
||||
// RUN: %clang_cc1 -DCK5 -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 -DCK5 -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 -DCK5 -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 -DCK5 -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 CK5
|
||||
|
||||
// CK5: [[STT:%.+]] = type { i32, double* }
|
||||
template <typename T>
|
||||
struct STT {
|
||||
T a;
|
||||
double *b;
|
||||
|
||||
T foo(T arg) {
|
||||
// Region 00
|
||||
#pragma omp target enter data map(always close to: b[1:3]) if(a>123) device(arg)
|
||||
{arg++;}
|
||||
return arg;
|
||||
}
|
||||
};
|
||||
|
||||
// CK5: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976711701]
|
||||
|
||||
// CK5-LABEL: _Z3bari
|
||||
int bar(int arg){
|
||||
STT<int> A;
|
||||
return A.foo(arg);
|
||||
}
|
||||
|
||||
// Region 00
|
||||
// CK5: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
|
||||
// CK5: [[IFTHEN]]
|
||||
// CK5-DAG: call void @__tgt_target_data_begin(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}})
|
||||
// CK5-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
|
||||
// CK5-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
|
||||
// CK5-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK5-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK5-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
||||
// CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK5-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK5-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[STT]]**
|
||||
// CK5-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
|
||||
// CK5-DAG: store [[STT]]* [[VAR0:%.+]], [[STT]]** [[CBP0]]
|
||||
// CK5-DAG: store double** [[SEC0:%.+]], double*** [[CP0]]
|
||||
// CK5-DAG: store i[[sz]] {{%.+}}, i[[sz]]* [[S0]]
|
||||
// CK5-DAG: [[SEC0]] = getelementptr inbounds {{.*}}[[STT]]* [[VAR0]], i32 0, i32 1
|
||||
|
||||
// CK5-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
|
||||
// CK5-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
|
||||
// CK5-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double***
|
||||
// CK5-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
|
||||
// CK5-DAG: store double** [[SEC0]], double*** [[CBP1]]
|
||||
// CK5-DAG: store double* [[SEC1:%.+]], double** [[CP1]]
|
||||
// CK5-DAG: [[SEC1]] = getelementptr inbounds {{.*}}double* [[SEC11:%[^,]+]], i{{.+}} 1
|
||||
// CK5-DAG: [[SEC11]] = load double*, double** [[SEC111:%[^,]+]],
|
||||
// CK5-DAG: [[SEC111]] = getelementptr inbounds {{.*}}[[STT]]* [[VAR0]], i32 0, i32 1
|
||||
|
||||
// CK5: br label %[[IFEND:[^,]+]]
|
||||
|
||||
// CK5: [[IFELSE]]
|
||||
// CK5: br label %[[IFEND]]
|
||||
// CK5: [[IFEND]]
|
||||
// CK5: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
// CK5-NOT: __tgt_target_data_end
|
||||
#endif
|
||||
#endif
|
||||
|
|
|
|||
|
|
@ -40,6 +40,10 @@ double gc[100];
|
|||
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
|
||||
// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710672]
|
||||
|
||||
// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1058]
|
||||
|
||||
// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1062]
|
||||
|
||||
// CK1-LABEL: _Z3fooi
|
||||
void foo(int arg) {
|
||||
int la;
|
||||
|
|
@ -146,6 +150,56 @@ void foo(int arg) {
|
|||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
#pragma omp target exit data map(release: gb.b[:3])
|
||||
{++arg;}
|
||||
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
{++arg;}
|
||||
|
||||
// Region 05
|
||||
// CK1-NOT: __tgt_target_data_begin
|
||||
// CK1-DAG: call void @__tgt_target_data_end(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}})
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
||||
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
|
||||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
|
||||
// CK1-DAG: store float* [[VAL0]], float** [[CP0]]
|
||||
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4
|
||||
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
#pragma omp target exit data map(close, from: lb)
|
||||
{++arg;}
|
||||
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
{++arg;}
|
||||
|
||||
// Region 06
|
||||
// CK1-NOT: __tgt_target_data_begin
|
||||
// CK1-DAG: call void @__tgt_target_data_end(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}})
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
||||
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
|
||||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
|
||||
// CK1-DAG: store float* [[VAL0]], float** [[CP0]]
|
||||
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4
|
||||
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
#pragma omp target exit data map(always close, from: lb)
|
||||
{++arg;}
|
||||
}
|
||||
#endif
|
||||
///==========================================================================///
|
||||
|
|
@ -252,4 +306,81 @@ void no_target_devices(int arg) {
|
|||
{++arg;}
|
||||
}
|
||||
#endif
|
||||
///==========================================================================///
|
||||
// RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
|
||||
// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
|
||||
// RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-32
|
||||
// RUN: %clang_cc1 -DCK4 -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 CK4 --check-prefix CK4-32
|
||||
|
||||
// RUN: %clang_cc1 -DCK4 -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 -DCK4 -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 -DCK4 -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 -DCK4 -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 CK4
|
||||
|
||||
// CK4: [[STT:%.+]] = type { i32, double* }
|
||||
template <typename T>
|
||||
struct STT {
|
||||
T a;
|
||||
double *b;
|
||||
|
||||
T foo(T arg) {
|
||||
// Region 00
|
||||
#pragma omp target exit data map(always close, release: b[1:3]) if(a>123) device(arg)
|
||||
{arg++;}
|
||||
return arg;
|
||||
}
|
||||
};
|
||||
|
||||
// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976711700]
|
||||
|
||||
// CK4-LABEL: _Z3bari
|
||||
int bar(int arg){
|
||||
STT<int> A;
|
||||
return A.foo(arg);
|
||||
}
|
||||
|
||||
// Region 00
|
||||
// CK4-NOT: __tgt_target_data_begin
|
||||
// CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
|
||||
// CK4: [[IFTHEN]]
|
||||
// CK4-DAG: call void @__tgt_target_data_end(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}})
|
||||
// CK4-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
|
||||
// CK4-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
|
||||
// CK4-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK4-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK4-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
||||
// CK4-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK4-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK4-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK4-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[STT]]**
|
||||
// CK4-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
|
||||
// CK4-DAG: store [[STT]]* [[VAR0:%[^,]+]], [[STT]]** [[CBP0]]
|
||||
// CK4-DAG: store double** [[SEC0:%[^,]+]], double*** [[CP0]]
|
||||
// CK4-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK4-DAG: [[SEC0]] = getelementptr inbounds {{.*}}[[STT]]* [[VAR0]], i32 0, i32 1
|
||||
|
||||
// CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
|
||||
// CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
|
||||
// CK4-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double***
|
||||
// CK4-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
|
||||
// CK4-DAG: store double** [[SEC0]], double*** [[CBP1]]
|
||||
// CK4-DAG: store double* [[SEC1:%[^,]+]], double** [[CP1]]
|
||||
// CK4-DAG: [[SEC1]] = getelementptr inbounds {{.*}}double* [[SEC11:%[^,]+]], i{{.+}} 1
|
||||
// CK4-DAG: [[SEC11]] = load double*, double** [[SEC111:%[^,]+]],
|
||||
// CK4-DAG: [[SEC111]] = getelementptr inbounds {{.*}}[[STT]]* [[VAR0]], i32 0, i32 1
|
||||
|
||||
// CK4: br label %[[IFEND:[^,]+]]
|
||||
|
||||
// CK4: [[IFELSE]]
|
||||
// CK4: br label %[[IFEND]]
|
||||
// CK4: [[IFEND]]
|
||||
// CK4: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
#endif
|
||||
#endif
|
||||
|
|
|
|||
|
|
@ -4046,7 +4046,7 @@ int explicit_maps_struct_fields(int a){
|
|||
// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
|
||||
// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}}
|
||||
|
||||
// CK24: call void [[CALL19:@.+]]([[SC]]* {{[^,]+}})
|
||||
// CK24: call void [[CALL19:@.+]]([[SC]]* {{[^,]+}})
|
||||
#pragma omp target map(p->s.sp[3]->a)
|
||||
{ p->a++; }
|
||||
|
||||
|
|
@ -5257,5 +5257,77 @@ void map_with_deep_copy() {
|
|||
}
|
||||
}
|
||||
|
||||
#endif
|
||||
///==========================================================================///
|
||||
// RUN: %clang_cc1 -DCK31 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK31 --check-prefix CK31-64
|
||||
// RUN: %clang_cc1 -DCK31 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK31 --check-prefix CK31-64
|
||||
// RUN: %clang_cc1 -DCK31 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK31 --check-prefix CK31-32
|
||||
// RUN: %clang_cc1 -DCK31 -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 -allow-deprecated-dag-overlap %s --check-prefix CK31 --check-prefix CK31-32
|
||||
|
||||
// RUN: %clang_cc1 -DCK31 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
|
||||
// RUN: %clang_cc1 -DCK31 -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 -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
|
||||
// RUN: %clang_cc1 -DCK31 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
|
||||
// RUN: %clang_cc1 -DCK31 -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 -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
|
||||
// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
|
||||
#ifdef CK31
|
||||
|
||||
// CK31-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l5305.region_id = weak constant i8 0
|
||||
// CK31: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
|
||||
// CK31: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 1059]
|
||||
|
||||
// CK31-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l5324.region_id = weak constant i8 0
|
||||
// CK31: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
|
||||
// CK31: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 1063]
|
||||
|
||||
// CK31-LABEL: explicit_maps_single{{.*}}(
|
||||
void explicit_maps_single (int ii){
|
||||
// Map of a scalar.
|
||||
int a = ii;
|
||||
|
||||
// Close.
|
||||
// Region 00
|
||||
// CK31-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
|
||||
// CK31-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK31-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
|
||||
// CK31-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK31-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK31-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
|
||||
// CK31-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
|
||||
// CK31-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
|
||||
// CK31-DAG: store i32* [[VAR0]], i32** [[CP0]]
|
||||
|
||||
// CK31: call void [[CALL00:@.+]](i32* {{[^,]+}})
|
||||
#pragma omp target map(close, tofrom: a)
|
||||
{
|
||||
a++;
|
||||
}
|
||||
|
||||
// Always Close.
|
||||
// Region 01
|
||||
// CK31-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
|
||||
// CK31-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK31-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
|
||||
// CK31-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK31-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK31-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
|
||||
// CK31-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
|
||||
// CK31-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
|
||||
// CK31-DAG: store i32* [[VAR0]], i32** [[CP0]]
|
||||
|
||||
// CK31: call void [[CALL01:@.+]](i32* {{[^,]+}})
|
||||
#pragma omp target map(always close tofrom: a)
|
||||
{
|
||||
a++;
|
||||
}
|
||||
}
|
||||
// CK31: define {{.+}}[[CALL00]]
|
||||
// CK31: define {{.+}}[[CALL01]]
|
||||
|
||||
#endif
|
||||
#endif
|
||||
|
|
|
|||
Loading…
Reference in New Issue