maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/OpenMP/target_codegen_global_captu...

456 lines
25 KiB
C++
Raw Normal View History

[OpenMP] Use fopenmp prefix for all options introduced by the offloading implementation. Summary: This patch changes the options used by offloading to start with -fopenmp instead of -fomp. This makes the option naming more consistent and materializes a suggestion by Richard Smith in http://reviews.llvm.org/D9888. Reviewers: hfinkel, carlo.bertolli, arpith-jacob, ABataev Subscribers: kkwli0, cfe-commits, caomhin Differential Revision: http://reviews.llvm.org/D21841 llvm-svn: 274283
2016-07-01 05:22:08 +08:00
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
[OPENMP] Support for -fopenmp-simd option with compilation of simd loops only. Added support for -fopenmp-simd option that allows compilation of simd-based constructs without emission of OpenMP runtime calls. llvm-svn: 321560
2017-12-30 02:07:07 +08:00
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
[OpenMP] Capture global variables in target regions. All global variables that are not enclosed in a declare target region must be captured in the target region as local variables do. Currently, there is no support for declare target, so this patch adds support for capturing all the global variables used in a the target region. llvm-svn: 249154
2015-10-03 01:14:03 +08:00
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
// CHECK-DAG: [[GA:@.+]] = global double 1.000000e+00
// CHECK-DAG: [[GB:@.+]] = global double 2.000000e+00
// CHECK-DAG: [[GC:@.+]] = global double 3.000000e+00
// CHECK-DAG: [[GD:@.+]] = global double 4.000000e+00
// CHECK-DAG: [[FA:@.+]] = internal global float 5.000000e+00
// CHECK-DAG: [[FB:@.+]] = internal global float 6.000000e+00
// CHECK-DAG: [[FC:@.+]] = internal global float 7.000000e+00
// CHECK-DAG: [[FD:@.+]] = internal global float 8.000000e+00
// CHECK-DAG: [[BA:@.+]] = internal global float 9.000000e+00
// CHECK-DAG: [[BB:@.+]] = internal global float 1.000000e+01
// CHECK-DAG: [[BC:@.+]] = internal global float 1.100000e+01
// CHECK-DAG: [[BD:@.+]] = internal global float 1.200000e+01
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[TBA:@.+]] = {{.*}}global float 1.700000e+01
// CHECK-DAG: [[TBB:@.+]] = {{.*}}global float 1.800000e+01
// CHECK-DAG: [[TBC:@.+]] = {{.*}}global float 1.900000e+01
// CHECK-DAG: [[TBD:@.+]] = {{.*}}global float 2.000000e+01
[OpenMP] Capture global variables in target regions. All global variables that are not enclosed in a declare target region must be captured in the target region as local variables do. Currently, there is no support for declare target, so this patch adds support for capturing all the global variables used in a the target region. llvm-svn: 249154
2015-10-03 01:14:03 +08:00
double Ga = 1.0;
double Gb = 2.0;
double Gc = 3.0;
double Gd = 4.0;
// CHECK: define {{.*}} @{{.*}}foo{{.*}}(
// CHECK-SAME: i16 {{[^,]*}}[[A:%[^,]+]],
// CHECK-SAME: i16 {{[^,]*}}[[B:%[^,]+]],
// CHECK-SAME: i16 {{[^,]*}}[[C:%[^,]+]],
// CHECK-SAME: i16 {{[^,]*}}[[D:%[^,]+]])
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK: [[LA:%.+]] = alloca i16,
// CHECK: [[LB:%.+]] = alloca i16,
// CHECK: [[LC:%.+]] = alloca i16,
// CHECK: [[LD:%.+]] = alloca i16,
[OpenMP] Capture global variables in target regions. All global variables that are not enclosed in a declare target region must be captured in the target region as local variables do. Currently, there is no support for declare target, so this patch adds support for capturing all the global variables used in a the target region. llvm-svn: 249154
2015-10-03 01:14:03 +08:00
int foo(short a, short b, short c, short d){
static float Sa = 5.0;
static float Sb = 6.0;
static float Sc = 7.0;
static float Sd = 8.0;
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LB]],
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
// CHECK-DAG: [[VALFB:%.+]] = load float, float* [[FB]],
// CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LC]],
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[VALFC:%.+]] = load float, float* [[FC]],
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LD]],
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
// CHECK-DAG: [[VALFD:%.+]] = load float, float* [[FD]],
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// 3 local vars being captured.
// CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
// CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
// CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALLB]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLB:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
// CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
// CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALLC]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLC:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
// CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
// CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALLD]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLD:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// 3 static vars being captured.
// CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
// CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
// CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALFB]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFB:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
// CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
// CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALFC]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFC:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
// CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
// CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALFD]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFD:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// 3 static global vars being captured.
// CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
// CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
// CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-64-DAG: store i[[sz]] [[CVALGB]], i[[sz]]* [[CBP:%.+]],
// CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to i[[sz]]*
// CHECK-32-DAG: store double* @Gb, double** [[CBP:%.+]],
// CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to double**
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
// CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
// CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-64-DAG: store i[[sz]] [[CVALGC]], i[[sz]]* [[CBP:%.+]],
// CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to i[[sz]]*
// CHECK-32-DAG: store double* @Gc, double** [[CBP:%.+]],
// CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to double**
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
// CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
// CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-64-DAG: store i[[sz]] [[CVALGD]], i[[sz]]* [[CBP:%.+]],
// CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to i[[sz]]*
// CHECK-32-DAG: store double* @Gd, double** [[CBP:%.+]],
// CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to double**
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
[OpenMP] Capture global variables in target regions. All global variables that are not enclosed in a declare target region must be captured in the target region as local variables do. Currently, there is no support for declare target, so this patch adds support for capturing all the global variables used in a the target region. llvm-svn: 249154
2015-10-03 01:14:03 +08:00
// CHECK: call i32 @__tgt_target
// CHECK: call void [[OFFLOADF:@.+]](
// Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
#pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)
{
b += 1;
Gb += 1.0;
Sb += 1.0;
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
[OpenMP] Capture global variables in target regions. All global variables that are not enclosed in a declare target region must be captured in the target region as local variables do. Currently, there is no support for declare target, so this patch adds support for capturing all the global variables used in a the target region. llvm-svn: 249154
2015-10-03 01:14:03 +08:00
// The parallel region only uses 3 captures.
[OPENMP] General code improvements. llvm-svn: 330140
2018-04-17 01:59:34 +08:00
// CHECK: call {{.*}}@__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), {{.+}}* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}})
[OpenMP] Capture global variables in target regions. All global variables that are not enclosed in a declare target region must be captured in the target region as local variables do. Currently, there is no support for declare target, so this patch adds support for capturing all the global variables used in a the target region. llvm-svn: 249154
2015-10-03 01:14:03 +08:00
// CHECK: call void @.omp_outlined.(i32* %{{.+}}, i32* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}})
// Capture d, Gd, Sd,
// CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}},
#pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)
{
d += 1;
Gd += 1.0;
Sd += 1.0;
}
}
return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
}
// CHECK: define {{.*}} @{{.*}}bar{{.*}}(
// CHECK-SAME: i16 {{[^,]*}}[[A:%[^,]+]],
// CHECK-SAME: i16 {{[^,]*}}[[B:%[^,]+]],
// CHECK-SAME: i16 {{[^,]*}}[[C:%[^,]+]],
// CHECK-SAME: i16 {{[^,]*}}[[D:%[^,]+]])
// CHECK: [[LA:%.+]] = alloca i16
// CHECK: [[LB:%.+]] = alloca i16
// CHECK: [[LC:%.+]] = alloca i16
// CHECK: [[LD:%.+]] = alloca i16
int bar(short a, short b, short c, short d){
static float Sa = 9.0;
static float Sb = 10.0;
static float Sc = 11.0;
static float Sd = 12.0;
[OPENMP] General code improvements. llvm-svn: 330140
2018-04-17 01:59:34 +08:00
// CHECK: call void {{.*}}@__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}})
[OpenMP] Capture global variables in target regions. All global variables that are not enclosed in a declare target region must be captured in the target region as local variables do. Currently, there is no support for declare target, so this patch adds support for capturing all the global variables used in a the target region. llvm-svn: 249154
2015-10-03 01:14:03 +08:00
// CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, i16* dereferenceable(2) [[A:%.+]], i16* dereferenceable(2) [[B:%.+]], i16* dereferenceable(2) [[C:%.+]], i16* dereferenceable(2) [[D:%.+]])
// Capture a, b, c, d
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK: [[ALLOCLA:%.+]] = alloca i16
// CHECK: [[ALLOCLB:%.+]] = alloca i16
// CHECK: [[ALLOCLC:%.+]] = alloca i16
// CHECK: [[ALLOCLD:%.+]] = alloca i16
// CHECK: [[LLA:%.+]] = load i16*, i16** [[ALLOCLA]],
// CHECK: [[LLB:%.+]] = load i16*, i16** [[ALLOCLB]],
// CHECK: [[LLC:%.+]] = load i16*, i16** [[ALLOCLC]],
// CHECK: [[LLD:%.+]] = load i16*, i16** [[ALLOCLD]],
[OpenMP] Capture global variables in target regions. All global variables that are not enclosed in a declare target region must be captured in the target region as local variables do. Currently, there is no support for declare target, so this patch adds support for capturing all the global variables used in a the target region. llvm-svn: 249154
2015-10-03 01:14:03 +08:00
#pragma omp parallel
{
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LLB]],
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
// CHECK-DAG: [[VALFB:%.+]] = load float, float* [[BB]],
// CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
// CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LLC]],
// CHECK-DAG: [[VALFC:%.+]] = load float, float* [[BC]],
// CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LLD]],
// CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
// CHECK-DAG: [[VALFD:%.+]] = load float, float* [[BD]],
// 3 local vars being captured.
// CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
// CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
// CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALLB]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLB:%.+]] to i[[sz]]*
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
// CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
// CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALLC]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLC:%.+]] to i[[sz]]*
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
// CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
// CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALLD]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLD:%.+]] to i[[sz]]*
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// 3 static vars being captured.
// CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
// CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
// CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALFB]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFB:%.+]] to i[[sz]]*
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
// CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
// CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALFC]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFC:%.+]] to i[[sz]]*
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
// CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
// CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALFD]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFD:%.+]] to i[[sz]]*
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// 3 static global vars being captured.
// CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
// CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
// CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-64-DAG: store i[[sz]] [[CVALGB]], i[[sz]]* [[CBP:%.+]],
// CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to i[[sz]]*
// CHECK-32-DAG: store double* @Gb, double** [[CBP:%.+]],
// CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to double**
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
// CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
// CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-64-DAG: store i[[sz]] [[CVALGC]], i[[sz]]* [[CBP:%.+]],
// CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to i[[sz]]*
// CHECK-32-DAG: store double* @Gc, double** [[CBP:%.+]],
// CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to double**
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
// CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
// CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-64-DAG: store i[[sz]] [[CVALGD]], i[[sz]]* [[CBP:%.+]],
// CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to i[[sz]]*
// CHECK-32-DAG: store double* @Gd, double** [[CBP:%.+]],
// CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to double**
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK: call i32 @__tgt_target
// CHECK: call void [[OFFLOADF:@.+]](
// Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
#pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)
{
b += 1;
Gb += 1.0;
Sb += 1.0;
// CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
[OPENMP] General code improvements. llvm-svn: 330140
2018-04-17 01:59:34 +08:00
// CHECK: call void {{.*}}@__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}})
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}})
// Capture d, Gd, Sd
#pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)
{
d += 1;
Gd += 1.0;
Sd += 1.0;
}
}
}
return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
}
///
/// Tests with template functions.
///
// CHECK: define {{.*}} @{{.*}}tbar2{{.*}}(
// CHECK: define {{.*}} @{{.*}}tbar{{.*}}(
// CHECK-SAME: i16 {{[^,]*}}[[A:%[^,]+]],
// CHECK-SAME: i16 {{[^,]*}}[[B:%[^,]+]],
// CHECK-SAME: i16 {{[^,]*}}[[C:%[^,]+]],
// CHECK-SAME: i16 {{[^,]*}}[[D:%[^,]+]])
// CHECK: [[LA:%.+]] = alloca i16
// CHECK: [[LB:%.+]] = alloca i16
// CHECK: [[LC:%.+]] = alloca i16
// CHECK: [[LD:%.+]] = alloca i16
template<typename T>
int tbar(T a, T b, T c, T d){
static float Sa = 17.0;
static float Sb = 18.0;
static float Sc = 19.0;
static float Sd = 20.0;
[OPENMP] General code improvements. llvm-svn: 330140
2018-04-17 01:59:34 +08:00
// CHECK: call void {{.*}}@__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}})
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, i16* dereferenceable(2) [[A:%.+]], i16* dereferenceable(2) [[B:%.+]], i16* dereferenceable(2) [[C:%.+]], i16* dereferenceable(2) [[D:%.+]])
// Capture a, b, c, d
// CHECK: [[ALLOCLA:%.+]] = alloca i16
// CHECK: [[ALLOCLB:%.+]] = alloca i16
// CHECK: [[ALLOCLC:%.+]] = alloca i16
// CHECK: [[ALLOCLD:%.+]] = alloca i16
// CHECK: [[LLA:%.+]] = load i16*, i16** [[ALLOCLA]],
// CHECK: [[LLB:%.+]] = load i16*, i16** [[ALLOCLB]],
// CHECK: [[LLC:%.+]] = load i16*, i16** [[ALLOCLC]],
// CHECK: [[LLD:%.+]] = load i16*, i16** [[ALLOCLD]],
#pragma omp parallel
{
// CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LLB]],
// CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
// CHECK-DAG: [[VALFB:%.+]] = load float, float* [[TBB]],
// CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LLC]],
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-DAG: [[VALFC:%.+]] = load float, float* [[TBC]],
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LLD]],
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
// CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
// CHECK-DAG: [[VALFD:%.+]] = load float, float* [[TBD]],
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// 3 local vars being captured.
// CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
// CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
// CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALLB]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLB:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
// CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
// CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALLC]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLC:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
// CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
// CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALLD]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPLD:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// 3 static vars being captured.
// CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
// CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
// CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALFB]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFB:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
// CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
// CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALFC]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFC:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
// CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
// CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-DAG: store i[[sz]] [[CVALFD]], i[[sz]]* [[CBP:%.+]],
// CHECK-DAG: [[CBP]] = bitcast i8** [[GEPFD:%.+]] to i[[sz]]*
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// 3 static global vars being captured.
// CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
// CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
// CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-64-DAG: store i[[sz]] [[CVALGB]], i[[sz]]* [[CBP:%.+]],
// CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to i[[sz]]*
// CHECK-32-DAG: store double* @Gb, double** [[CBP:%.+]],
// CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGB:%.+]] to double**
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
// CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
// CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-64-DAG: store i[[sz]] [[CVALGC]], i[[sz]]* [[CBP:%.+]],
// CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to i[[sz]]*
// CHECK-32-DAG: store double* @Gc, double** [[CBP:%.+]],
// CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGC:%.+]] to double**
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
// CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
// CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
// CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
[OPENMP][DEBUG] Generate second function with correct arg types. Currently, if the some of the parameters are captured by value, this argument is converted to uintptr_t type and thus we loosing the debug info about real type of the argument (captured variable): ``` void @.outlined_function.(uintptr %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.outlined_function.(uintptr %a.casted.val) ... ``` To resolve this problem, in debug mode a speciall external wrapper function is generated, that calls the outlined function with the correct parameters types: ``` void @.wrapper.(uintptr %par) { %a = alloca i32 %cast = bitcast i32* %a to uintptr* store uintptr %par, uintptr *%cast %a.val = load i32, i32* %a call void @.outlined_function.(i32 %a) ret void } void @.outlined_function.(i32 %par); ... %a = alloca i32 %a.casted = alloca uintptr %cast = bitcast uintptr* %a.casted to i32* %a.val = load i32, i32 *%a store i32 %a.val, i32 *%cast %a.casted.val = load uintptr, uintptr* %a.casted call void @.wrapper.(uintptr %a.casted.val) ... ``` llvm-svn: 306697
2017-06-30 00:43:05 +08:00
// CHECK-64-DAG: store i[[sz]] [[CVALGD]], i[[sz]]* [[CBP:%.+]],
// CHECK-64-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to i[[sz]]*
// CHECK-32-DAG: store double* @Gd, double** [[CBP:%.+]],
// CHECK-32-DAG: [[CBP]] = bitcast i8** [[GEPGD:%.+]] to double**
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
[OpenMP] Capture global variables in target regions. All global variables that are not enclosed in a declare target region must be captured in the target region as local variables do. Currently, there is no support for declare target, so this patch adds support for capturing all the global variables used in a the target region. llvm-svn: 249154
2015-10-03 01:14:03 +08:00
// CHECK: call i32 @__tgt_target
// CHECK: call void [[OFFLOADF:@.+]](
// Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
#pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)
{
b += 1;
Gb += 1.0;
Sb += 1.0;
[OpenMP] Update target directive codegen to use 4.5 implicit data mappings. Summary: This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed. Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available. Reviewers: hfinkel, rjmccall, ABataev Subscribers: cfe-commits, carlo.bertolli Differential Revision: http://reviews.llvm.org/D14940 llvm-svn: 254521
2015-12-03 01:44:43 +08:00
// CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
[OPENMP] General code improvements. llvm-svn: 330140
2018-04-17 01:59:34 +08:00
// CHECK: call void {{.*}}@__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}})
[OpenMP] Capture global variables in target regions. All global variables that are not enclosed in a declare target region must be captured in the target region as local variables do. Currently, there is no support for declare target, so this patch adds support for capturing all the global variables used in a the target region. llvm-svn: 249154
2015-10-03 01:14:03 +08:00
// CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}})
// Capture d, Gd, Sd
#pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)
{
d += 1;
Gd += 1.0;
Sd += 1.0;
}
}
}
return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
}
[OpenMP] Fix SEMA bug in the capture of global variables in template functions. Summary: Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function. Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin Differential Revision: http://reviews.llvm.org/D18110 llvm-svn: 271001
2016-05-27 23:21:27 +08:00
int tbar2(short a, short b, short c, short d){
return tbar(a, b, c, d);
}
[OpenMP] Capture global variables in target regions. All global variables that are not enclosed in a declare target region must be captured in the target region as local variables do. Currently, there is no support for declare target, so this patch adds support for capturing all the global variables used in a the target region. llvm-svn: 249154
2015-10-03 01:14:03 +08:00
#endif