86 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// Device side kernel name.11// HOST: @[[KERN_CAPTURE:[0-9]+]] = {{.*}} c"_Z1gIZ12test_capturevEUlvE_EvT_\00"12// HOST: @[[KERN_RESOLVE:[0-9]+]] = {{.*}} c"_Z1gIZ12test_resolvevEUlvE_EvT_\00"13 14// Check functions emitted for test_capture in host compilation.15// Check lambda is not emitted in host compilation.16// HOST-LABEL: define{{.*}} void @_Z12test_capturev17// HOST: call void @_Z19test_capture_helperIZ12test_capturevEUlvE_EvT_18// HOST-LABEL: define internal void @_Z19test_capture_helperIZ12test_capturevEUlvE_EvT_19// HOST: call void @_Z16__device_stub__gIZ12test_capturevEUlvE_EvT_20// HOST-NOT: define{{.*}}@_ZZ4mainENKUlvE_clEv21 22// Check functions emitted for test_resolve in host compilation.23// Check host version of template function 'overloaded' is emitted and called24// by the lambda function.25// HOST-LABEL: define{{.*}} void @_Z12test_resolvev26// HOST: call void @_Z19test_resolve_helperIZ12test_resolvevEUlvE_EvT_()27// HOST-LABEL: define internal void @_Z19test_resolve_helperIZ12test_resolvevEUlvE_EvT_28// HOST: call void @_Z16__device_stub__gIZ12test_resolvevEUlvE_EvT_29// HOST: call void @_ZZ12test_resolvevENKUlvE_clEv30// HOST-LABEL: define internal void @_ZZ12test_resolvevENKUlvE_clEv31// HOST: call noundef i32 @_Z10overloadedIiET_v32// HOST-LABEL: define linkonce_odr noundef i32 @_Z10overloadedIiET_v33// HOST: ret i32 234 35// Check kernel is registered with correct device side kernel name.36// HOST: @__hipRegisterFunction({{.*}}@[[KERN_CAPTURE]]37// HOST: @__hipRegisterFunction({{.*}}@[[KERN_RESOLVE]]38 39// DEV: @a ={{.*}} addrspace(1) externally_initialized global i32 040 41// Check functions emitted for test_capture in device compilation.42// Check lambda is emitted in device compilation and accessing device variable.43// DEV-LABEL: define{{.*}} amdgpu_kernel void @_Z1gIZ12test_capturevEUlvE_EvT_44// DEV: call void @_ZZ12test_capturevENKUlvE_clEv45// DEV-LABEL: define internal void @_ZZ12test_capturevENKUlvE_clEv46// DEV: store i32 1, ptr addrspacecast (ptr addrspace(1) @a to ptr)47 48// Check functions emitted for test_resolve in device compilation.49// Check device version of template function 'overloaded' is emitted and called50// by the lambda function.51// DEV-LABEL: define{{.*}} amdgpu_kernel void @_Z1gIZ12test_resolvevEUlvE_EvT_52// DEV: call void @_ZZ12test_resolvevENKUlvE_clEv53// DEV-LABEL: define internal void @_ZZ12test_resolvevENKUlvE_clEv54// DEV: call noundef i32 @_Z10overloadedIiET_v55// DEV-LABEL: define linkonce_odr noundef i32 @_Z10overloadedIiET_v56// DEV: ret i32 157 58__device__ int a;59 60template<class T>61__device__ T overloaded() { return 1; }62 63template<class T>64__host__ T overloaded() { return 2; }65 66template<class F>67__global__ void g(F f) { f(); }68 69template<class F>70void test_capture_helper(F f) { g<<<1,1>>>(f); }71 72template<class F>73void test_resolve_helper(F f) { g<<<1,1>>>(f); f(); }74 75// Test capture of device variable in lambda function.76void test_capture(void) {77 test_capture_helper([](){ a = 1;});78}79 80// Test resolving host/device function in lambda function.81// Callee should resolve to correct host/device function based on where82// the lambda function is called, not where it is defined.83void test_resolve(void) {84 test_resolve_helper([](){ overloaded<int>();});85}86