maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[OpenMP] Adjust map type bits according to latest spec and use zero size array sections for pointers. 

Summary: This patch changes the bits used to specify the map types according to the latest version of the libomptarget document and add the support for zero size array section when pointers are being implicitly mapped. This completes the missing new 4.5 map semantics.

Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev

Subscribers: caomhin, cfe-commits

Differential Revision: http://reviews.llvm.org/D20111

llvm-svn: 270868
This commit is contained in:
Samuel Antao 2016-05-26 16:48:10 +00:00
parent 0a665a83da
commit 6782e944d2
8 changed files with 348 additions and 225 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

71
clang/lib/CodeGen/CGOpenMPRuntime.cpp
View File

@ -4909,8 +4909,6 @@ public:
/// \brief Values for bit flags used to specify the mapping type for
/// offloading.
enum OpenMPOffloadMappingFlags {
/// \brief Only allocate memory on the device,
OMP_MAP_ALLOC = 0x00,
/// \brief Allocate memory on the device and move data from host to device.
OMP_MAP_TO = 0x01,
/// \brief Allocate memory on the device and move data from device to host.
@ -4918,19 +4916,19 @@ public:
/// \brief Always perform the requested mapping action on the element, even
/// if it was already mapped before.
OMP_MAP_ALWAYS = 0x04,
/// \brief Decrement the reference count associated with the element without
/// executing any other action.
OMP_MAP_RELEASE = 0x08,
/// \brief Delete the element from the device environment, ignoring the
/// current reference count associated with the element.
OMP_MAP_DELETE = 0x10,
/// \brief The element passed to the device is a pointer.
OMP_MAP_PTR = 0x20,
/// \brief Signal the element as extra, i.e. is not argument to the target
/// region kernel.
OMP_MAP_EXTRA = 0x40,
OMP_MAP_DELETE = 0x08,
/// \brief The element being mapped is a pointer, therefore the pointee
/// should be mapped as well.
OMP_MAP_IS_PTR = 0x10,
/// \brief This flags signals that an argument is the first one relating to
/// a map/private clause expression. For some cases a single
/// map/privatization results in multiple arguments passed to the runtime
/// library.
OMP_MAP_FIRST_REF = 0x20,
/// \brief Pass the element to the device by value.
OMP_MAP_BYCOPY = 0x80,
OMP_MAP_PRIVATE_VAL = 0x100,
};
typedef SmallVector<llvm::Value *, 16> MapValuesArrayTy;
@ -4987,14 +4985,19 @@ private:
/// \brief Return the corresponding bits for a given map clause modifier. Add
/// a flag marking the map as a pointer if requested. Add a flag marking the
/// map as extra, meaning is not an argument of the kernel.
/// map as the first one of a series of maps that relate to the same map
/// expression.
unsigned getMapTypeBits(OpenMPMapClauseKind MapType,
OpenMPMapClauseKind MapTypeModifier, bool AddPtrFlag,
bool AddExtraFlag) const {
bool AddIsFirstFlag) const {
unsigned Bits = 0u;
switch (MapType) {
case OMPC_MAP_alloc:
Bits = OMP_MAP_ALLOC;
case OMPC_MAP_release:
// alloc and release is the default behavior in the runtime library, i.e.
// if we don't pass any bits alloc/release that is what the runtime is
// going to do. Therefore, we don't need to signal anything for these two
// type modifiers.
break;
case OMPC_MAP_to:
Bits = OMP_MAP_TO;
@ -5008,17 +5011,14 @@ private:
case OMPC_MAP_delete:
Bits = OMP_MAP_DELETE;
break;
case OMPC_MAP_release:
Bits = OMP_MAP_RELEASE;
break;
default:
llvm_unreachable("Unexpected map type!");
break;
}
if (AddPtrFlag)
Bits |= OMP_MAP_PTR;
if (AddExtraFlag)
Bits |= OMP_MAP_EXTRA;
Bits |= OMP_MAP_IS_PTR;
if (AddIsFirstFlag)
Bits |= OMP_MAP_FIRST_REF;
if (MapTypeModifier == OMPC_MAP_always)
Bits |= OMP_MAP_ALWAYS;
return Bits;
@ -5270,13 +5270,13 @@ private:
Pointers.push_back(LB);
Sizes.push_back(Size);
// We need to add a pointer flag for each map that comes from the the
// same expression except for the first one. We need to add the extra
// flag for each map that relates with the current capture, except for
// the first one (there is a set of entries for each capture).
// We need to add a pointer flag for each map that comes from the
// same expression except for the first one. We also need to signal
// this map is the first one that relates with the current capture
// (there is a set of entries for each capture).
Types.push_back(getMapTypeBits(MapType, MapTypeModifier,
!IsExpressionFirstInfo,
!IsCaptureFirstInfo));
IsCaptureFirstInfo));
// If we have a final array section, we are done with this expression.
if (IsFinalArraySection)
@ -5583,7 +5583,8 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
CurPointers.push_back(*CV);
CurSizes.push_back(CGF.getTypeSize(RI->getType()));
// Copy to the device as an argument. No need to retrieve it.
CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_BYCOPY);
CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_PRIVATE_VAL |
MappableExprsHandler::OMP_MAP_FIRST_REF);
} else {
// If we have any information in the map clause, we use it, otherwise we
// just do a default mapping.
@ -5602,10 +5603,10 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_TO |
MappableExprsHandler::OMP_MAP_FROM);
} else if (CI->capturesVariableByCopy()) {
CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_BYCOPY);
if (!RI->getType()->isAnyPointerType()) {
// If the field is not a pointer, we need to save the actual value
// and load it as a void pointer.
CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_PRIVATE_VAL);
auto DstAddr = CGF.CreateMemTemp(
Ctx.getUIntPtrType(),
Twine(CI->getCapturedVar()->getName()) + ".casted");
@ -5624,11 +5625,17 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
CurBasePointers.push_back(
CGF.EmitLoadOfLValue(DstLV, SourceLocation()).getScalarVal());
CurPointers.push_back(CurBasePointers.back());
// Get the size of the type to be used in the map.
CurSizes.push_back(CGF.getTypeSize(RI->getType()));
} else {
// Pointers are implicitly mapped with a zero size and no flags
// (other than first map that is added for all implicit maps).
CurMapTypes.push_back(0u);
CurBasePointers.push_back(*CV);
CurPointers.push_back(*CV);
CurSizes.push_back(llvm::Constant::getNullValue(CGM.SizeTy));
}
CurSizes.push_back(CGF.getTypeSize(RI->getType()));
} else {
assert(CI->capturesVariable() && "Expected captured reference.");
CurBasePointers.push_back(*CV);
@ -5640,13 +5647,15 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
CurSizes.push_back(CGF.getTypeSize(ElementType));
// The default map type for a scalar/complex type is 'to' because by
// default the value doesn't have to be retrieved. For an aggregate
// type,
// the default is 'tofrom'.
// type, the default is 'tofrom'.
CurMapTypes.push_back(ElementType->isAggregateType()
? (MappableExprsHandler::OMP_MAP_TO |
MappableExprsHandler::OMP_MAP_FROM)
: MappableExprsHandler::OMP_MAP_TO);
}
// Every default map produces a single argument, so, it is always the
// first one.
CurMapTypes.back() |= MappableExprsHandler::OMP_MAP_FIRST_REF;
}
}
// We expect to have at least an element of information for this capture.

12
clang/test/OpenMP/target_codegen.cpp
View File

@ -33,15 +33,15 @@
// sizes.
// CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ:32|64]] 2]
// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2]
// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i32] [i32 128, i32 128]
// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i32] [i32 128, i32 3, i32 128, i32 3, i32 3, i32 128, i32 128, i32 3, i32 3]
// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i32] [i32 288, i32 288]
// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 35, i32 288, i32 35, i32 35, i32 288, i32 288, i32 35, i32 35]
// CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 40]
// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 128, i32 128, i32 3]
// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 35]
// CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40]
// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i32] [i32 128, i32 128, i32 128, i32 3]
// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i32] [i32 3, i32 128, i32 128, i32 128, i32 3]
// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i32] [i32 288, i32 288, i32 288, i32 35]
// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 35]
// CHECK-DAG: @{{.*}} = private constant i8 0
// CHECK-DAG: @{{.*}} = private constant i8 0
// CHECK-DAG: @{{.*}} = private constant i8 0

24
clang/test/OpenMP/target_codegen_registration.cpp
View File

@ -61,40 +61,40 @@
// CHECK-DAG: {{@.+}} = private constant i8 0
// TCHECK-NOT: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: {{@.+}} = private constant i8 0
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-NTARGET-NOT: private constant i8 0
// CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i

10
clang/test/OpenMP/target_data_codegen.cpp
View File

@ -22,15 +22,15 @@ ST<int> gb;
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 2]
// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 34]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 1]
// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 33]
// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 5]
// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 37]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] 24]
// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 1, i32 97]
// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 33, i32 17]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
@ -173,7 +173,7 @@ struct ST {
};
// CK2: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 24]
// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 5, i32 101]
// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 37, i32 21]
// CK2-LABEL: _Z3bari
int bar(int arg){

10
clang/test/OpenMP/target_enter_data_codegen.cpp
View File

@ -22,15 +22,15 @@ ST<int> gb;
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] zeroinitializer
// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 32]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 1]
// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 33]
// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 5]
// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 37]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] 24]
// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 1, i32 97]
// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 33, i32 17]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
@ -156,7 +156,7 @@ struct ST {
};
// CK2: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 24]
// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 5, i32 101]
// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 37, i32 21]
// CK2-LABEL: _Z3bari
int bar(int arg){

10
clang/test/OpenMP/target_exit_data_codegen.cpp
View File

@ -22,15 +22,15 @@ ST<int> gb;
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 2]
// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 34]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 8]
// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 32]
// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 6]
// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 38]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] 24]
// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 8, i32 104]
// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 32, i32 16]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
@ -156,7 +156,7 @@ struct ST {
};
// CK2: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 24]
// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 12, i32 108]
// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 36, i32 20]
// CK2-LABEL: _Z3bari
int bar(int arg){

15
clang/test/OpenMP/target_firstprivate_codegen.cpp
View File

@ -33,16 +33,15 @@ struct TT{
// TCHECK: [[S1:%.+]] = type { double }
// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ:32|64]] 4]
// CHECK: [[MAPT:@.+]] = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [9 x i32] [i32 128, i32 3, i32 128, i32 3, i32 3, i32 128, i32 128, i32 3, i32 3]
// CHECK-64-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ]] 8]
// CHECK-32-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [1 x i32] [i[[SZ]] 4]
// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [1 x i32] [i32 128]
// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [5 x i32] [i32 3, i32 128, i32 128, i32 128, i32 3]
// CHECK: [[MAPT:@.+]] = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 35, i32 288, i32 35, i32 35, i32 288, i32 288, i32 35, i32 35]
// CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] zeroinitializer
// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [1 x i32] [i32 32]
// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 35]
// CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 1, i[[SZ]] 40]
// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 128, i32 128, i32 3]
// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 35]
// CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 40]
// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [2 x i32] [i32 128, i32 3]
// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [2 x i32] [i32 288, i32 35]
// CHECK: define {{.*}}[[FOO:@.+]](

421
clang/test/OpenMP/target_map_codegen.cpp
View File

@ -16,8 +16,8 @@
#ifdef CK1
// CK1-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
// Map types: OMP_MAP_BYCOPY = 128
// CK1-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
// CK1-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
// CK1-LABEL: implicit_maps_integer
void implicit_maps_integer (int a){
@ -61,8 +61,8 @@ void implicit_maps_integer (int a){
#ifdef CK2
// CK2-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
// Map types: OMP_MAP_BYCOPY = 128
// CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
// CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
// CK2-LABEL: implicit_maps_integer_reference
void implicit_maps_integer_reference (int a){
@ -110,8 +110,8 @@ void implicit_maps_integer_reference (int a){
#ifdef CK3
// CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
// Map types: OMP_MAP_BYCOPY = 128
// CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
// CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
// CK3-LABEL: implicit_maps_parameter
void implicit_maps_parameter (int a){
@ -154,8 +154,8 @@ void implicit_maps_parameter (int a){
#ifdef CK4
// CK4-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
// Map types: OMP_MAP_BYCOPY = 128
// CK4-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
// CK4-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
// CK4-LABEL: implicit_maps_nested_integer
void implicit_maps_nested_integer (int a){
@ -210,8 +210,8 @@ void implicit_maps_nested_integer (int a){
#ifdef CK5
// CK5-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
// Map types: OMP_MAP_BYCOPY = 128
// CK5-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
// CK5-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
// CK5-LABEL: implicit_maps_nested_integer_and_enum
void implicit_maps_nested_integer_and_enum (int a){
@ -261,8 +261,8 @@ void implicit_maps_nested_integer_and_enum (int a){
#ifdef CK6
// CK6-DAG: [[GBL:@Gi]] = global i32 0
// CK6-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
// Map types: OMP_MAP_BYCOPY = 128
// CK6-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
// CK6-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
// CK6-LABEL: implicit_maps_host_global
int Gi;
@ -310,10 +310,10 @@ void implicit_maps_host_global (int a){
// therefore it is passed by reference with a map 'to' specification.
// CK7-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8]
// Map types: OMP_MAP_BYCOPY = 128
// CK7-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// Map types: OMP_MAP_TO = 1
// CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1]
// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
// CK7-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
// Map types: OMP_MAP_TO | OMP_MAP_IS_FIRST = 33
// CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33]
// CK7-LABEL: implicit_maps_double
void implicit_maps_double (int a){
@ -369,8 +369,8 @@ void implicit_maps_double (int a){
#ifdef CK8
// CK8-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
// Map types: OMP_MAP_BYCOPY = 128
// CK8-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
// CK8-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
// CK8-LABEL: implicit_maps_float
void implicit_maps_float (int a){
@ -413,8 +413,8 @@ void implicit_maps_float (int a){
#ifdef CK9
// CK9-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16]
// Map types: OMP_MAP_TO + OMP_MAP_FROM = 2 + 1
// CK9-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 3]
// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
// CK9-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 35]
// CK9-LABEL: implicit_maps_array
void implicit_maps_array (int a){
@ -453,9 +453,9 @@ void implicit_maps_array (int a){
// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK10
#ifdef CK10
// CK10-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}]
// Map types: OMP_MAP_BYCOPY = 128
// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// CK10-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] zeroinitializer
// Map types: OMP_MAP_IS_FIRST = 32
// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 32]
// CK10-LABEL: implicit_maps_pointer
void implicit_maps_pointer (){
@ -496,8 +496,8 @@ void implicit_maps_pointer (){
#ifdef CK11
// CK11-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16]
// Map types: OMP_MAP_TO = 1
// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1]
// Map types: OMP_MAP_TO + OMP_MAP_IS_FIRST = 33
// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33]
// CK11-LABEL: implicit_maps_double_complex
void implicit_maps_double_complex (int a){
@ -539,10 +539,10 @@ void implicit_maps_double_complex (int a){
// therefore it is passed by reference with a map 'to' specification.
// CK12-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8]
// Map types: OMP_MAP_BYCOPY = 128
// CK12-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// Map types: OMP_MAP_TO = 1
// CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1]
// Map types: OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
// CK12-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
// Map types: OMP_MAP_TO + OMP_MAP_IS_FIRST = 33
// CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33]
// CK12-LABEL: implicit_maps_float_complex
void implicit_maps_float_complex (int a){
@ -598,10 +598,10 @@ void implicit_maps_float_complex (int a){
// We don't have a constant map size for VLAs.
// Map types:
// - OMP_MAP_BYCOPY = 128 (vla size)
// - OMP_MAP_BYCOPY = 128 (vla size)
// - OMP_MAP_TO + OMP_MAP_FROM = 2 + 1
// CK13-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 128, i32 128, i32 3]
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 (vla size)
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 (vla size)
// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
// CK13-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 288, i32 288, i32 35]
// CK13-LABEL: implicit_maps_variable_length_array
void implicit_maps_variable_length_array (int a){
@ -669,9 +669,9 @@ void implicit_maps_variable_length_array (int a){
// CK14-DAG: [[ST:%.+]] = type { i32, double }
// CK14-DAG: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}, i{{64|32}} 4]
// Map types:
// - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2
// - OMP_MAP_BYCOPY = 128
// CK14-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128]
// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
// CK14-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288]
class SSS {
public:
@ -743,15 +743,15 @@ void implicit_maps_class (int a){
// CK15: [[ST:%.+]] = type { i32, double, i32* }
// CK15: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4]
// Map types:
// - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2
// - OMP_MAP_BYCOPY = 128
// CK15: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128]
// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
// CK15: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288]
// CK15: [[SIZES2:@.+]] = {{.+}}constant [2 x i[[sz]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4]
// Map types:
// - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2
// - OMP_MAP_BYCOPY = 128
// CK15: [[TYPES2:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128]
// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
// CK15: [[TYPES2:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288]
template<int x>
class SSST {
@ -870,8 +870,8 @@ void implicit_maps_templated_class (int a){
// CK16-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
// Map types:
// - OMP_MAP_BYCOPY = 128
// CK16-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
// CK16-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
template<int y>
int foo(int d) {
@ -923,8 +923,8 @@ void implicit_maps_templated_function (int a){
// CK17-DAG: [[ST:%.+]] = type { i32, double }
// CK17-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}]
// Map types: OMP_MAP_TO + OMP_MAP_FROM = 2 + 1
// CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 3]
// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
// CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 35]
class SSS {
public:
@ -971,8 +971,8 @@ void implicit_maps_struct (int a){
// CK18-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
// Map types:
// - OMP_MAP_BYCOPY = 128
// CK18-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
// CK18-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
template<typename T>
int foo(T d) {
@ -1022,122 +1022,122 @@ void implicit_maps_template_type_capture (int a){
#ifdef CK19
// CK19: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
// CK19: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
// CK19: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
// CK19: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
// CK19: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
// CK19: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
// CK19: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
// CK19: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
// CK19: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
// CK19: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
// CK19: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
// CK19: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
// CK19: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
// CK19: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK19: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
// CK19: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
// CK19: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
// CK19: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
// CK19: [[SIZE08:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK19: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
// CK19: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
// CK19: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
// CK19: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
// CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[SIZE11:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
// CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
// CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
// CK19: [[SIZE12:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
// CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
// CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
// CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
// CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
// CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
// CK19: [[SIZE15:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
// CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
// CK19: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 1]
// CK19: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 33]
// CK19: [[SIZE17:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 240]
// CK19: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 2]
// CK19: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 34]
// CK19: [[SIZE18:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 240]
// CK19: [[MTYPE18:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3]
// CK19: [[MTYPE18:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 35]
// CK19: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 0]
// CK19: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 32]
// CK19: [[SIZE20:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
// CK19: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 1]
// CK19: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 33]
// CK19: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3]
// CK19: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 35]
// CK19: [[SIZE22:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
// CK19: [[MTYPE22:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3]
// CK19: [[MTYPE22:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 35]
// CK19: [[SIZE23:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK19: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i32] [i32 7]
// CK19: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i32] [i32 39]
// CK19: [[SIZE24:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 480]
// CK19: [[MTYPE24:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE24:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[SIZE25:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
// CK19: [[MTYPE25:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE25:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[SIZE26:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 24]
// CK19: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[SIZE27:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK19: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[SIZE28:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 16]
// CK19: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i32] [i32 3, i32 99, i32 99]
// CK19: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i32] [i32 35, i32 19, i32 19]
// CK19: [[SIZE29:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4]
// CK19: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i32] [i32 3, i32 99, i32 99]
// CK19: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i32] [i32 35, i32 19, i32 19]
// CK19: [[MTYPE30:@.+]] = private {{.*}}constant [4 x i32] [i32 128, i32 128, i32 128, i32 3]
// CK19: [[MTYPE30:@.+]] = private {{.*}}constant [4 x i32] [i32 288, i32 288, i32 288, i32 35]
// CK19: [[SIZE31:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 40]
// CK19: [[MTYPE31:@.+]] = private {{.*}}constant [4 x i32] [i32 128, i32 128, i32 128, i32 3]
// CK19: [[MTYPE31:@.+]] = private {{.*}}constant [4 x i32] [i32 288, i32 288, i32 288, i32 35]
// CK19: [[SIZE32:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 13728]
// CK19: [[MTYPE32:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE32:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[SIZE33:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 13728]
// CK19: [[MTYPE33:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE33:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[SIZE34:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 13728]
// CK19: [[MTYPE34:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE34:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[MTYPE35:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE35:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[SIZE36:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 208]
// CK19: [[MTYPE36:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE36:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19: [[MTYPE37:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3]
// CK19: [[MTYPE37:@.+]] = private {{.*}}constant [3 x i32] [i32 288, i32 288, i32 35]
// CK19: [[MTYPE38:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3]
// CK19: [[MTYPE38:@.+]] = private {{.*}}constant [3 x i32] [i32 288, i32 288, i32 35]
// CK19: [[MTYPE39:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3]
// CK19: [[MTYPE39:@.+]] = private {{.*}}constant [3 x i32] [i32 288, i32 288, i32 35]
// CK19: [[MTYPE40:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3]
// CK19: [[MTYPE40:@.+]] = private {{.*}}constant [3 x i32] [i32 288, i32 288, i32 35]
// CK19: [[SIZE41:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 208]
// CK19: [[MTYPE41:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3]
// CK19: [[MTYPE41:@.+]] = private {{.*}}constant [3 x i32] [i32 288, i32 288, i32 35]
// CK19: [[SIZE42:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 104]
// CK19: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i32] [i32 3, i32 99, i32 99]
// CK19: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i32] [i32 35, i32 19, i32 19]
// CK19: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK19: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK19-LABEL: explicit_maps_single
void explicit_maps_single (int ii){
@ -2397,16 +2397,16 @@ void explicit_maps_single (int ii){
#ifdef CK20
// CK20: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
// CK20: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
// CK20: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
// CK20: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
// CK20: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
// CK20: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
// CK20: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK20: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
// CK20: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
// CK20: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 12]
// CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
// CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
// CK20-LABEL: explicit_maps_references_and_function_args
void explicit_maps_references_and_function_args (int a, float b, int (&c)[10], float *d){
@ -2514,22 +2514,22 @@ void explicit_maps_references_and_function_args (int a, float b, int (&c)[10], f
// CK21: [[ST:%.+]] = type { i32, i32, float* }
// CK21: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
// CK21: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK21: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK21: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 492]
// CK21: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK21: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK21: [[SIZE02:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 500]
// CK21: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i32] [i32 2, i32 98]
// CK21: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i32] [i32 34, i32 18]
// CK21: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 492]
// CK21: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
// CK21: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
// CK21: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK21: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
// CK21: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
// CK21: [[SIZE05:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] 4, i[[Z]] 4]
// CK21: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 67]
// CK21: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 3]
// CK21-LABEL: explicit_maps_template_args_and_members
@ -2705,49 +2705,49 @@ int explicit_maps_template_args_and_members(int a){
// CK22-DAG: [[STT:%.+]] = type { i32 }
// CK22: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
// CK22: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
// CK22: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
// CK22: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
// CK22: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK22: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
// CK22: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE07:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
// CK22: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE08:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
// CK22: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK22: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE11:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
// CK22: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE12:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
// CK22: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
// CK22: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK22: [[SIZE14:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
int a;
int c[100];
@ -3025,22 +3025,22 @@ int explicit_maps_globals(void){
#ifdef CK23
// CK23: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
// CK23: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK23: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK23: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK23: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK23: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK23: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
// CK23: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK23: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK23: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
// CK23: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK23: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK23: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
// CK23: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK23: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK23: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
// CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK23-LABEL: explicit_maps_inside_captured
int explicit_maps_inside_captured(int a){
@ -3215,76 +3215,76 @@ struct SC{
};
// CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}]
// CK24: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK24: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
// CK24: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE05:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{3560|2880}}]
// CK24: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
// CK24: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
// CK24: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK24: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE07:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
// CK24: [[MTYPE07:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
// CK24: [[MTYPE07:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
// CK24: [[SIZE08:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
// CK24: [[MTYPE08:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
// CK24: [[MTYPE08:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
// CK24: [[SIZE09:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
// CK24: [[MTYPE09:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
// CK24: [[MTYPE09:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
// CK24: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 8]
// CK24: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE11:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}]
// CK24: [[MTYPE11:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
// CK24: [[MTYPE11:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
// CK24: [[SIZE12:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4]
// CK24: [[MTYPE12:@.+]] = private {{.*}}constant [4 x i32] [i32 3, i32 99, i32 99, i32 99]
// CK24: [[MTYPE12:@.+]] = private {{.*}}constant [4 x i32] [i32 35, i32 19, i32 19, i32 19]
// CK24: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE14:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}]
// CK24: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE15:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK24: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE16:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
// CK24: [[MTYPE16:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE16:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE17:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{3560|2880}}]
// CK24: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
// CK24: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
// CK24: [[SIZE18:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK24: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE19:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
// CK24: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
// CK24: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
// CK24: [[SIZE20:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
// CK24: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
// CK24: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
// CK24: [[SIZE21:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
// CK24: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
// CK24: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
// CK24: [[SIZE22:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
// CK24: [[MTYPE22:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
// CK24: [[MTYPE22:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK24: [[SIZE23:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}]
// CK24: [[MTYPE23:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
// CK24: [[MTYPE23:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
// CK24: [[SIZE24:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4]
// CK24: [[MTYPE24:@.+]] = private {{.*}}constant [4 x i32] [i32 3, i32 99, i32 99, i32 99]
// CK24: [[MTYPE24:@.+]] = private {{.*}}constant [4 x i32] [i32 35, i32 19, i32 19, i32 19]
// CK24-LABEL: explicit_maps_struct_fields
int explicit_maps_struct_fields(int a){
@ -3997,10 +3997,10 @@ int explicit_maps_struct_fields(int a){
// CK25: [[CA01:%.+]] = type { i32* }
// CK25: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
// CK25: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
// CK25: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
// CK25: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK25: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
// CK25: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
// CK25-LABEL: explicit_maps_with_inner_lambda
@ -4087,16 +4087,16 @@ int explicit_maps_with_inner_lambda(int a){
// CK26: [[ST:%.+]] = type { i32, float*, i32, float* }
// CK26: [[SIZE00:@.+]] = private {{.*}}constant [2 x i[[Z:64|32]]] [i[[Z:64|32]] {{32|16}}, i[[Z:64|32]] 4]
// CK26: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3]
// CK26: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 35]
// CK26: [[SIZE01:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{32|16}}, i[[Z]] 4]
// CK26: [[MTYPE01:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3]
// CK26: [[MTYPE01:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 35]
// CK26: [[SIZE02:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{32|16}}, i[[Z]] 4]
// CK26: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3]
// CK26: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 35]
// CK26: [[SIZE03:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{32|16}}, i[[Z]] 4]
// CK26: [[MTYPE03:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3]
// CK26: [[MTYPE03:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 35]
// CK26-LABEL: explicit_maps_with_private_class_members
@ -4259,5 +4259,120 @@ int explicit_maps_with_private_class_members(){
CC c(B);
return c.foo();
}
#endif
///==========================================================================///
// RUN: %clang_cc1 -DCK27 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK27 --check-prefix CK27-64
// RUN: %clang_cc1 -DCK27 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK27 --check-prefix CK27-64
// RUN: %clang_cc1 -DCK27 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK27 --check-prefix CK27-32
// RUN: %clang_cc1 -DCK27 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK27 --check-prefix CK27-32
#ifdef CK27
// CK27: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] zeroinitializer
// CK27: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
// CK27: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer
// CK27: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK27: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer
// CK27: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK27: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer
// CK27: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
// CK27-LABEL: zero_size_section_maps
void zero_size_section_maps (int ii){
// Map of a pointer.
int *pa;
// Region 00
// CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
// CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
// CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
// CK27-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8*
// CK27-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8*
// CK27: call void [[CALL00:@.+]](i32* {{[^,]+}})
#pragma omp target
{
pa[50]++;
}
// Region 01
// CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
// CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
// CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
// CK27-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
// CK27-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
// CK27-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
// CK27-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0
// CK27-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
// CK27: call void [[CALL01:@.+]](i32* {{[^,]+}})
#pragma omp target map(pa[:0])
{
pa[50]++;
}
// Region 02
// CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
// CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
// CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
// CK27-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
// CK27-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
// CK27-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
// CK27-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0
// CK27-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
// CK27: call void [[CALL02:@.+]](i32* {{[^,]+}})
#pragma omp target map(pa[0:0])
{
pa[50]++;
}
// Region 03
// CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}})
// CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
// CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
// CK27-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
// CK27-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
// CK27-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
// CK27-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.+}}
// CK27-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
// CK27: call void [[CALL03:@.+]](i32* {{[^,]+}})
#pragma omp target map(pa[ii:0])
{
pa[50]++;
}
}
// CK27: define {{.+}}[[CALL00]]
// CK27: define {{.+}}[[CALL01]]
// CK27: define {{.+}}[[CALL02]]
// CK27: define {{.+}}[[CALL03]]
#endif
#endif