maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/CodeGenCUDA/lambda-reference-var.cu

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

136 lines
4.9 KiB
Plaintext
Raw Normal View History

[CUDA][HIP] Fix capturing reference to host variable In C++ when a reference variable is captured by copy, the lambda is supposed to make a copy of the referenced variable in the captures and refer to the copy in the lambda. Therefore, it is valid to capture a reference to a host global variable in a device lambda since the device lambda will refer to the copy of the host global variable instead of access the host global variable directly. However, clang tries to avoid capturing of reference to a host global variable if it determines the use of the reference variable in the lambda function is not odr-use. Clang also tries to emit load of the reference to a global variable as load of the global variable if it determines that the reference variable is a compile-time constant. For a device lambda to capture a reference variable to host global variable and use the captured value, clang needs to be taught that in such cases the use of the reference variable is odr-use and the reference variable is not compile-time constant. This patch fixes that. Differential Revision: https://reviews.llvm.org/D91088
2020-11-04 04:24:41 +08:00
// RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
// RUN: -triple x86_64-linux-gnu \
// RUN: | FileCheck -check-prefix=HOST %s
// RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
// RUN: -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: | FileCheck -check-prefix=DEV %s
#include "Inputs/cuda.h"
// HOST: %[[T1:.*]] = type <{ i32*, i32, [4 x i8] }>
// HOST: %[[T2:.*]] = type { i32*, i32** }
// HOST: %[[T3:.*]] = type <{ i32*, i32, [4 x i8] }>
// DEV: %[[T1:.*]] = type { i32* }
// DEV: %[[T2:.*]] = type { i32** }
// DEV: %[[T3:.*]] = type <{ i32*, i32, [4 x i8] }>
int global_host_var;
__device__ int global_device_var;
template<class F>
__global__ void kern(F f) { f(); }
// DEV-LABEL: @_ZZ27dev_capture_dev_ref_by_copyPiENKUlvE_clEv(
// DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
// DEV: store i32 %[[VAL]]
__device__ void dev_capture_dev_ref_by_copy(int *out) {
int &ref = global_device_var;
[=](){ *out = ref;}();
}
Fix assertion in tryEmitAsConstant due to cd95338ee3022bffd658e52cd3eb9419b4c218ca Need to check if result is LValue before getLValueBase.
2020-12-03 07:35:52 +08:00
// DEV-LABEL: @_ZZ28dev_capture_dev_rval_by_copyPiENKUlvE_clEv(
// DEV: store i32 3
__device__ void dev_capture_dev_rval_by_copy(int *out) {
constexpr int a = 1;
constexpr int b = 2;
constexpr int c = a + b;
[=](){ *out = c;}();
}
[CUDA][HIP] Fix capturing reference to host variable In C++ when a reference variable is captured by copy, the lambda is supposed to make a copy of the referenced variable in the captures and refer to the copy in the lambda. Therefore, it is valid to capture a reference to a host global variable in a device lambda since the device lambda will refer to the copy of the host global variable instead of access the host global variable directly. However, clang tries to avoid capturing of reference to a host global variable if it determines the use of the reference variable in the lambda function is not odr-use. Clang also tries to emit load of the reference to a global variable as load of the global variable if it determines that the reference variable is a compile-time constant. For a device lambda to capture a reference variable to host global variable and use the captured value, clang needs to be taught that in such cases the use of the reference variable is odr-use and the reference variable is not compile-time constant. This patch fixes that. Differential Revision: https://reviews.llvm.org/D91088
2020-11-04 04:24:41 +08:00
// DEV-LABEL: @_ZZ26dev_capture_dev_ref_by_refPiENKUlvE_clEv(
// DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
// DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
// DEV: store i32 %[[VAL2]], i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
// DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
// DEV: store i32 %[[VAL]]
__device__ void dev_capture_dev_ref_by_ref(int *out) {
int &ref = global_device_var;
[&](){ ref++; *out = ref;}();
}
[test] Make ELF tests immune to dso_local/dso_preemptable/(none) differences ELF -cc1 -mrelocation-model pic will default to no semantic interposition plus setting dso_local on default visibility external linkage definitions, so that COFF, Mach-O and ELF output will be similar. This patch makes tests immune to the differences.
2021-01-01 05:37:35 +08:00
// DEV-LABEL: define{{.*}} void @_Z7dev_refPi(
[CUDA][HIP] Fix capturing reference to host variable In C++ when a reference variable is captured by copy, the lambda is supposed to make a copy of the referenced variable in the captures and refer to the copy in the lambda. Therefore, it is valid to capture a reference to a host global variable in a device lambda since the device lambda will refer to the copy of the host global variable instead of access the host global variable directly. However, clang tries to avoid capturing of reference to a host global variable if it determines the use of the reference variable in the lambda function is not odr-use. Clang also tries to emit load of the reference to a global variable as load of the global variable if it determines that the reference variable is a compile-time constant. For a device lambda to capture a reference variable to host global variable and use the captured value, clang needs to be taught that in such cases the use of the reference variable is odr-use and the reference variable is not compile-time constant. This patch fixes that. Differential Revision: https://reviews.llvm.org/D91088
2020-11-04 04:24:41 +08:00
// DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
// DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
// DEV: store i32 %[[VAL2]], i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
// DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
// DEV: store i32 %[[VAL]]
__device__ void dev_ref(int *out) {
int &ref = global_device_var;
ref++;
*out = ref;
}
// DEV-LABEL: @_ZZ14dev_lambda_refPiENKUlvE_clEv(
// DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
// DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
// DEV: store i32 %[[VAL2]], i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
// DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
// DEV: store i32 %[[VAL]]
__device__ void dev_lambda_ref(int *out) {
[=](){
int &ref = global_device_var;
ref++;
*out = ref;
}();
}
// HOST-LABEL: @_ZZ29host_capture_host_ref_by_copyPiENKUlvE_clEv(
// HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
// HOST: store i32 %[[VAL]]
void host_capture_host_ref_by_copy(int *out) {
int &ref = global_host_var;
[=](){ *out = ref;}();
}
// HOST-LABEL: @_ZZ28host_capture_host_ref_by_refPiENKUlvE_clEv(
// HOST: %[[CAP:.*]] = getelementptr inbounds %[[T2]], %[[T2]]* %this1, i32 0, i32 0
// HOST: %[[REF:.*]] = load i32*, i32** %[[CAP]]
// HOST: %[[VAL:.*]] = load i32, i32* %[[REF]]
// HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
// HOST: store i32 %[[VAL2]], i32* %[[REF]]
// HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
// HOST: store i32 %[[VAL]]
void host_capture_host_ref_by_ref(int *out) {
int &ref = global_host_var;
[&](){ ref++; *out = ref;}();
}
[test] Make ELF tests immune to dso_local/dso_preemptable/(none) differences ELF -cc1 -mrelocation-model pic will default to no semantic interposition plus setting dso_local on default visibility external linkage definitions, so that COFF, Mach-O and ELF output will be similar. This patch makes tests immune to the differences.
2021-01-01 05:37:35 +08:00
// HOST-LABEL: define{{.*}} void @_Z8host_refPi(
[CUDA][HIP] Fix capturing reference to host variable In C++ when a reference variable is captured by copy, the lambda is supposed to make a copy of the referenced variable in the captures and refer to the copy in the lambda. Therefore, it is valid to capture a reference to a host global variable in a device lambda since the device lambda will refer to the copy of the host global variable instead of access the host global variable directly. However, clang tries to avoid capturing of reference to a host global variable if it determines the use of the reference variable in the lambda function is not odr-use. Clang also tries to emit load of the reference to a global variable as load of the global variable if it determines that the reference variable is a compile-time constant. For a device lambda to capture a reference variable to host global variable and use the captured value, clang needs to be taught that in such cases the use of the reference variable is odr-use and the reference variable is not compile-time constant. This patch fixes that. Differential Revision: https://reviews.llvm.org/D91088
2020-11-04 04:24:41 +08:00
// HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
// HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
// HOST: store i32 %[[VAL2]], i32* @global_host_var
// HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
// HOST: store i32 %[[VAL]]
void host_ref(int *out) {
int &ref = global_host_var;
ref++;
*out = ref;
}
// HOST-LABEL: @_ZZ15host_lambda_refPiENKUlvE_clEv(
// HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
// HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
// HOST: store i32 %[[VAL2]], i32* @global_host_var
// HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
// HOST: store i32 %[[VAL]]
void host_lambda_ref(int *out) {
[=](){
int &ref = global_host_var;
ref++;
*out = ref;
}();
}
[test] Make ELF tests immune to dso_local/dso_preemptable/(none) differences ELF -cc1 -mrelocation-model pic will default to no semantic interposition plus setting dso_local on default visibility external linkage definitions, so that COFF, Mach-O and ELF output will be similar. This patch makes tests immune to the differences.
2021-01-01 05:37:35 +08:00
// HOST-LABEL: define{{.*}} void @_Z28dev_capture_host_ref_by_copyPi(
[CUDA][HIP] Fix capturing reference to host variable In C++ when a reference variable is captured by copy, the lambda is supposed to make a copy of the referenced variable in the captures and refer to the copy in the lambda. Therefore, it is valid to capture a reference to a host global variable in a device lambda since the device lambda will refer to the copy of the host global variable instead of access the host global variable directly. However, clang tries to avoid capturing of reference to a host global variable if it determines the use of the reference variable in the lambda function is not odr-use. Clang also tries to emit load of the reference to a global variable as load of the global variable if it determines that the reference variable is a compile-time constant. For a device lambda to capture a reference variable to host global variable and use the captured value, clang needs to be taught that in such cases the use of the reference variable is odr-use and the reference variable is not compile-time constant. This patch fixes that. Differential Revision: https://reviews.llvm.org/D91088
2020-11-04 04:24:41 +08:00
// HOST: %[[CAP:.*]] = getelementptr inbounds %[[T3]], %[[T3]]* %{{.*}}, i32 0, i32 1
// HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
// HOST: store i32 %[[VAL]], i32* %[[CAP]]
// DEV-LABEL: define internal void @_ZZ28dev_capture_host_ref_by_copyPiENKUlvE_clEv(
// DEV: %[[CAP:.*]] = getelementptr inbounds %[[T3]], %[[T3]]* %this1, i32 0, i32 1
// DEV: %[[VAL:.*]] = load i32, i32* %[[CAP]]
// DEV: store i32 %[[VAL]]
void dev_capture_host_ref_by_copy(int *out) {
int &ref = global_host_var;
kern<<<1, 1>>>([=]__device__() { *out = ref;});
}