sfantao updated this revision to Diff 58252. sfantao added a comment. - Rebase after last changes related to first private.
http://reviews.llvm.org/D20111 Files: lib/CodeGen/CGOpenMPRuntime.cpp test/OpenMP/target_codegen.cpp test/OpenMP/target_codegen_registration.cpp test/OpenMP/target_data_codegen.cpp test/OpenMP/target_enter_data_codegen.cpp test/OpenMP/target_exit_data_codegen.cpp test/OpenMP/target_firstprivate_codegen.cpp test/OpenMP/target_map_codegen.cpp
Index: test/OpenMP/target_map_codegen.cpp =================================================================== --- test/OpenMP/target_map_codegen.cpp +++ test/OpenMP/target_map_codegen.cpp @@ -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 @@ #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 @@ #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 @@ #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 @@ #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 @@ #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 @@ // 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 @@ #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 @@ #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 @@ // 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 @@ #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 @@ // 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 @@ // 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 @@ // 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 @@ // 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 @@ // 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 @@ // 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 @@ // 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 @@ #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 @@ #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 @@ // 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 @@ // 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 @@ #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 @@ }; // 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 @@ // 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 @@ // 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 @@ -4260,4 +4260,119 @@ 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 Index: test/OpenMP/target_firstprivate_codegen.cpp =================================================================== --- test/OpenMP/target_firstprivate_codegen.cpp +++ test/OpenMP/target_firstprivate_codegen.cpp @@ -33,16 +33,15 @@ // 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:@.+]]( Index: test/OpenMP/target_exit_data_codegen.cpp =================================================================== --- test/OpenMP/target_exit_data_codegen.cpp +++ test/OpenMP/target_exit_data_codegen.cpp @@ -22,15 +22,15 @@ 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 @@ }; // 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){ Index: test/OpenMP/target_enter_data_codegen.cpp =================================================================== --- test/OpenMP/target_enter_data_codegen.cpp +++ test/OpenMP/target_enter_data_codegen.cpp @@ -22,15 +22,15 @@ 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 @@ }; // 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){ Index: test/OpenMP/target_data_codegen.cpp =================================================================== --- test/OpenMP/target_data_codegen.cpp +++ test/OpenMP/target_data_codegen.cpp @@ -22,15 +22,15 @@ 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 @@ }; // 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){ Index: test/OpenMP/target_codegen_registration.cpp =================================================================== --- test/OpenMP/target_codegen_registration.cpp +++ test/OpenMP/target_codegen_registration.cpp @@ -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 Index: test/OpenMP/target_codegen.cpp =================================================================== --- test/OpenMP/target_codegen.cpp +++ test/OpenMP/target_codegen.cpp @@ -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 Index: lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- lib/CodeGen/CGOpenMPRuntime.cpp +++ lib/CodeGen/CGOpenMPRuntime.cpp @@ -4858,28 +4858,26 @@ /// \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. OMP_MAP_FROM = 0x02, /// \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; @@ -4936,14 +4934,19 @@ /// \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; @@ -4957,17 +4960,14 @@ 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; @@ -5219,13 +5219,13 @@ 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) @@ -5532,7 +5532,8 @@ 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. @@ -5551,10 +5552,10 @@ 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"); @@ -5573,11 +5574,17 @@ 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); @@ -5589,13 +5596,15 @@ 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.
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits