98 lines · plain
1// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \2// RUN: -aux-triple x86_64-unknown-linux-gnu -std=c++17 -fgpu-rdc \3// RUN: -emit-llvm -o - -x hip %s > %t.dev4 5// RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \6// RUN: -aux-triple amdgcn-amd-amdhsa -std=c++17 -fgpu-rdc \7// RUN: -emit-llvm -o - -x hip %s > %t.host8 9// RUN: cat %t.dev %t.host | FileCheck -check-prefixes=HIP,COMMON %s10// RUN: cat %t.dev %t.host | FileCheck -check-prefixes=COMNEG %s11 12// RUN: echo "GPU binary" > %t.fatbin13 14// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -cuid=abc \15// RUN: -aux-triple x86_64-unknown-linux-gnu -std=c++17 -fgpu-rdc \16// RUN: -emit-llvm -o - %s > %t.dev17 18// RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \19// RUN: -aux-triple nvptx -std=c++17 -fgpu-rdc -fcuda-include-gpubinary %t.fatbin \20// RUN: -emit-llvm -o - %s > %t.host21 22// RUN: cat %t.dev %t.host | FileCheck -check-prefixes=CUDA,COMMON %s23// RUN: cat %t.dev %t.host | FileCheck -check-prefixes=COMNEG %s24 25#include "Inputs/cuda.h"26 27// HIP-DAG: define weak_odr {{.*}}void @[[KERN:_ZN12_GLOBAL__N_16kernelEv\.intern\.b04fd23c98500190]](28// HIP-DAG: define weak_odr {{.*}}void @[[KTX:_Z2ktIN12_GLOBAL__N_11XEEvT_\.intern\.b04fd23c98500190]](29// HIP-DAG: define weak_odr {{.*}}void @[[KTL:_Z2ktIN12_GLOBAL__N_1UlvE_EEvT_\.intern\.b04fd23c98500190]](30// HIP-DAG: @[[VM:_ZN12_GLOBAL__N_12vmE\.static\.b04fd23c98500190]] = addrspace(1) externally_initialized global31// HIP-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE\.static\.b04fd23c98500190]] = addrspace(4) externally_initialized constant32// HIP-DAG: @[[VT:_Z2vtIN12_GLOBAL__N_11XEE\.static\.b04fd23c98500190]] = addrspace(1) externally_initialized global33 34// CUDA-DAG: define weak_odr {{.*}}void @[[KERN:_ZN12_GLOBAL__N_16kernelEv__intern__b04fd23c98500190]](35// CUDA-DAG: define weak_odr {{.*}}void @[[KTX:_Z2ktIN12_GLOBAL__N_11XEEvT___intern__b04fd23c98500190]](36// CUDA-DAG: define weak_odr {{.*}}void @[[KTL:_Z2ktIN12_GLOBAL__N_1UlvE_EEvT___intern__b04fd23c98500190]](37// CUDA-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE__static__b04fd23c98500190]] = addrspace(4) externally_initialized constant38// CUDA-DAG: @[[VT:_Z2vtIN12_GLOBAL__N_11XEE__static__b04fd23c98500190]] = addrspace(1) externally_initialized global39 40// COMMON-DAG: @_ZN12_GLOBAL__N_12vdE = internal addrspace(1) global41// COMNEG-NOT: @{{.*}} = {{.*}} c"_ZN12_GLOBAL__N_12vdE{{.*}}\00"42 43// HIP-DAG: @llvm.compiler.used = {{.*}}@[[VM]]{{.*}}@[[VT]]{{.*}}@[[VC]]44// CUDA-DAG: @llvm.compiler.used = {{.*}}@[[VT]]{{.*}}@[[VC]]45 46// COMMON-DAG: @[[KERNSTR:.*]] = {{.*}} c"[[KERN]]\00"47// COMMON-DAG: @[[KTXSTR:.*]] = {{.*}} c"[[KTX]]\00"48// COMMON-DAG: @[[KTLSTR:.*]] = {{.*}} c"[[KTL]]\00"49// HIP-DAG: @[[VMSTR:.*]] = {{.*}} c"[[VM]]\00"50// COMMON-DAG: @[[VCSTR:.*]] = {{.*}} c"[[VC]]\00"51// COMMON-DAG: @[[VTSTR:.*]] = {{.*}} c"[[VT]]\00"52 53// COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KERNSTR]]54// COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KTXSTR]]55// COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KTLSTR]]56// HIP-DAG: call void @__{{.*}}RegisterManagedVar({{.*}}@[[VMSTR]]57// COMMON-DAG: call void @__{{.*}}RegisterVar({{.*}}@[[VCSTR]]58// COMMON-DAG: call void @__{{.*}}RegisterVar({{.*}}@[[VTSTR]]59 60template <typename T>61__global__ void kt(T x) {}62 63template <typename T>64__device__ T vt;65 66namespace {67 struct X {};68 X x;69 auto lambda = [](){};70#if __HIP__71 __managed__ int vm = 1;72#endif73 __constant__ int vc = 2;74 75 // C should not be externalized since it is used by device code only.76 __device__ int vd = 3;77 __global__ void kernel() { vd = 4; }78}79 80template<typename T>81void getSymbol(T *x) {}82 83void test() {84 kernel<<<1, 1>>>();85 86 kt<<<1, 1>>>(x);87 88 kt<<<1, 1>>>(lambda);89 90 // A, B, and tempVar<X> should be externalized since they are91 // used by host code.92#if __HIP__93 getSymbol(&vm);94#endif95 getSymbol(&vc);96 getSymbol(&vt<X>);97}98