137 lines · plain
1// REQUIRES: x86-registered-target2// REQUIRES: amdgpu-registered-target3 4// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \5// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o %t.nocuid.dev -x hip %s6// RUN: cat %t.nocuid.dev | FileCheck -check-prefixes=DEV,INT-DEV %s7 8// RUN: %clang_cc1 -triple x86_64-gnu-linux \9// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o %t.nocuid.host -x hip %s10// RUN: cat %t.nocuid.host | FileCheck -check-prefixes=HOST,INT-HOST %s11 12// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \13// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev14// RUN: cat %t.dev | FileCheck -check-prefixes=DEV,EXT-DEV %s15 16// RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \17// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.host18// RUN: cat %t.host | FileCheck -check-prefixes=HOST,EXT-HOST %s19 20// Check host and device compilations use the same postfixes for static21// variable names.22 23// RUN: cat %t.dev %t.host | FileCheck -check-prefix=POSTFIX %s24// RUN: cat %t.nocuid.dev %t.nocuid.host | FileCheck -check-prefix=POSTFIX-ID %s25 26// Negative tests.27 28// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \29// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \30// RUN: -check-prefix=DEV-NEG %s31 32// RUN: %clang_cc1 -triple x86_64-gnu-linux \33// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \34// RUN: -check-prefix=HOST-NEG %s35 36// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \37// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev38// RUN: cat %t.dev | FileCheck -check-prefix=DEV-NEG %s39 40// RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \41// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.host42// RUN: cat %t.host | FileCheck -check-prefix=HOST-NEG %s43 44// Check postfix for CUDA.45 46// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -cuid=abc \47// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - %s | FileCheck \48// RUN: -check-prefixes=CUDA %s49 50#include "Inputs/cuda.h"51 52// Make sure we can still mangle with a line directive.53#line 0 "-"54 55// Test function scope static device variable, which should not be externalized.56// DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 157 58 59// HOST-DAG: @_ZL1x = internal global i32 undef60// HOST-DAG: @_ZL1y = internal global i32 undef61 62// Test normal static device variables63// INT-DEV-DAG: @_ZL1x[[FILEID:.*]] = addrspace(1) externally_initialized global i32 064// INT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x[[FILEID:.*]]\00"65 66// Test externalized static device variables67// EXT-DEV-DAG: @_ZL1x.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 068// EXT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH:.*]]\00"69// CUDA-DAG: @_ZL1x__static__[[HASH:.*]] = addrspace(1) externally_initialized global i32 070 71// POSTFIX: @_ZL1x.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 072// POSTFIX: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH]]\00"73// POSTFIX-ID: @_ZL1x.static.[[FILEID:.*]] = addrspace(1) externally_initialized global i32 074// POSTFIX-ID: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[FILEID]]\00"75 76static __device__ int x;77 78// Test static device variables not used by host code should not be externalized79// DEV-DAG: @_ZL2x2 = internal addrspace(1) global i32 080 81static __device__ int x2;82 83// Test normal static device variables84// INT-DEV-DAG: @_ZL1y[[FILEID:.*]] = addrspace(4) externally_initialized constant i32 085// INT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y[[FILEID:.*]]\00"86 87// Test externalized static device variables88// EXT-DEV-DAG: @_ZL1y.static.[[HASH]] = addrspace(4) externally_initialized constant i32 089// EXT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y.static.[[HASH]]\00"90 91static __constant__ int y;92 93// Test static host variable, which should not be externalized nor registered.94// HOST-DAG: @_ZL1z = internal global i32 095// DEV-NEG-NOT: @_ZL1z96static int z;97 98// Test non-ODR-use of static device variable is not emitted or registered.99// DEV-NEG-NOT: @_ZL1u100// HOST-NEG-NOT: @_ZL1u101static __device__ int u;102 103// Test static device variable in inline function, which should not be104// externalized nor registered.105// DEV-DAG: @_ZZ6devfunPPKiE1p = linkonce_odr addrspace(4) constant i32 2, comdat106 107inline __device__ void devfun(const int ** b) {108 const static int p = 2;109 b[0] = &p;110}111 112__global__ void kernel(int *a, const int **b) {113 const static int w = 1;114 a[0] = x;115 a[1] = y;116 a[2] = sizeof(u);117 b[0] = &w;118 b[1] = &x2;119 devfun(b);120}121 122int* getDeviceSymbol(int *x);123 124void foo() {125 getDeviceSymbol(&x);126 getDeviceSymbol(&y);127 z = 123;128 decltype(u) tmp;129}130 131// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1x, {{.*}}@[[DEVNAMEX]]132// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1y, {{.*}}@[[DEVNAMEY]]133// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL2x2134// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w135// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p136// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL1u137