136 lines · plain
1// RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \2// RUN: -triple x86_64-linux-gnu \3// RUN: | FileCheck -check-prefix=HOST %s4// RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \5// RUN: -triple amdgcn-amd-amdhsa -fcuda-is-device \6// RUN: | FileCheck -check-prefix=DEV %s7 8#include "Inputs/cuda.h"9 10// HOST: %[[T1:.*]] = type <{ ptr, i32, [4 x i8] }>11// HOST: %[[T2:.*]] = type { ptr, ptr }12// HOST: %[[T3:.*]] = type <{ ptr, i32, [4 x i8] }>13// DEV: %[[T1:.*]] = type { ptr }14// DEV: %[[T2:.*]] = type { ptr }15// DEV: %[[T3:.*]] = type <{ ptr, i32, [4 x i8] }>16int global_host_var;17__device__ int global_device_var;18 19template<class F>20__global__ void kern(F f) { f(); }21 22// DEV-LABEL: @_ZZ27dev_capture_dev_ref_by_copyPiENKUlvE_clEv(23// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)24// DEV: store i32 %[[VAL]]25__device__ void dev_capture_dev_ref_by_copy(int *out) {26 int &ref = global_device_var;27 [=](){ *out = ref;}();28}29 30// DEV-LABEL: @_ZZ28dev_capture_dev_rval_by_copyPiENKUlvE_clEv(31// DEV: store i32 332__device__ void dev_capture_dev_rval_by_copy(int *out) {33 constexpr int a = 1;34 constexpr int b = 2;35 constexpr int c = a + b;36 [=](){ *out = c;}();37}38 39// DEV-LABEL: @_ZZ26dev_capture_dev_ref_by_refPiENKUlvE_clEv(40// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)41// DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 142// DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)43// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)44// DEV: store i32 %[[VAL]]45__device__ void dev_capture_dev_ref_by_ref(int *out) {46 int &ref = global_device_var;47 [&](){ ref++; *out = ref;}();48}49 50// DEV-LABEL: define{{.*}} void @_Z7dev_refPi(51// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)52// DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 153// DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)54// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)55// DEV: store i32 %[[VAL]]56__device__ void dev_ref(int *out) {57 int &ref = global_device_var;58 ref++;59 *out = ref;60}61 62// DEV-LABEL: @_ZZ14dev_lambda_refPiENKUlvE_clEv(63// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)64// DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 165// DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)66// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)67// DEV: store i32 %[[VAL]]68__device__ void dev_lambda_ref(int *out) {69 [=](){70 int &ref = global_device_var;71 ref++;72 *out = ref;73 }();74}75 76// HOST-LABEL: @_ZZ29host_capture_host_ref_by_copyPiENKUlvE_clEv(77// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var78// HOST: store i32 %[[VAL]]79void host_capture_host_ref_by_copy(int *out) {80 int &ref = global_host_var;81 [=](){ *out = ref;}();82}83 84// HOST-LABEL: @_ZZ28host_capture_host_ref_by_refPiENKUlvE_clEv(85// HOST: %[[CAP:.*]] = getelementptr inbounds nuw %[[T2]], ptr %this1, i32 0, i32 086// HOST: %[[REF:.*]] = load ptr, ptr %[[CAP]]87// HOST: %[[VAL:.*]] = load i32, ptr %[[REF]]88// HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 189// HOST: store i32 %[[VAL2]], ptr %[[REF]]90// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var91// HOST: store i32 %[[VAL]]92void host_capture_host_ref_by_ref(int *out) {93 int &ref = global_host_var;94 [&](){ ref++; *out = ref;}();95}96 97// HOST-LABEL: define{{.*}} void @_Z8host_refPi(98// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var99// HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1100// HOST: store i32 %[[VAL2]], ptr @global_host_var101// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var102// HOST: store i32 %[[VAL]]103void host_ref(int *out) {104 int &ref = global_host_var;105 ref++;106 *out = ref;107}108 109// HOST-LABEL: @_ZZ15host_lambda_refPiENKUlvE_clEv(110// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var111// HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1112// HOST: store i32 %[[VAL2]], ptr @global_host_var113// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var114// HOST: store i32 %[[VAL]]115void host_lambda_ref(int *out) {116 [=](){117 int &ref = global_host_var;118 ref++;119 *out = ref;120 }();121}122 123// HOST-LABEL: define{{.*}} void @_Z28dev_capture_host_ref_by_copyPi(124// HOST: %[[CAP:.*]] = getelementptr inbounds nuw %[[T3]], ptr %{{.*}}, i32 0, i32 1125// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var126// HOST: store i32 %[[VAL]], ptr %[[CAP]]127// DEV-LABEL: define internal void @_ZZ28dev_capture_host_ref_by_copyPiENKUlvE_clEv(128// DEV: %[[CAP:.*]] = getelementptr inbounds nuw %[[T3]], ptr %this1, i32 0, i32 1129// DEV: %[[VAL:.*]] = load i32, ptr %[[CAP]]130// DEV: store i32 %[[VAL]]131void dev_capture_host_ref_by_copy(int *out) {132 int &ref = global_host_var;133 kern<<<1, 1>>>([=]__device__() { *out = ref;});134}135 136