Yaxun (Sam) Liu | 45f2a56 | 2020-08-04 12:13:16 -0400 | [diff] [blame] | 1 | // REQUIRES: x86-registered-target |
| 2 | // REQUIRES: amdgpu-registered-target |
| 3 | |
Yaxun (Sam) Liu | fb04d7b | 2020-08-10 16:38:10 -0400 | [diff] [blame] | 4 | // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \ |
Yaxun (Sam) Liu | 45f2a56 | 2020-08-04 12:13:16 -0400 | [diff] [blame] | 5 | // RUN: -emit-llvm -o - -x hip %s | FileCheck \ |
| 6 | // RUN: -check-prefixes=DEV %s |
| 7 | |
Yaxun (Sam) Liu | fb04d7b | 2020-08-10 16:38:10 -0400 | [diff] [blame] | 8 | // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \ |
Yaxun (Sam) Liu | 45f2a56 | 2020-08-04 12:13:16 -0400 | [diff] [blame] | 9 | // RUN: -emit-llvm -o - -x hip %s | FileCheck \ |
| 10 | // RUN: -check-prefixes=HOST %s |
| 11 | |
| 12 | #include "Inputs/cuda.h" |
| 13 | |
| 14 | // Test function scope static device variable, which should not be externalized. |
| 15 | // DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1 |
| 16 | |
| 17 | // Check a static device variable referenced by host function is externalized. |
| 18 | // DEV-DAG: @_ZL1x = addrspace(1) externally_initialized global i32 0 |
| 19 | // HOST-DAG: @_ZL1x = internal global i32 undef |
| 20 | // HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x\00" |
| 21 | |
| 22 | static __device__ int x; |
| 23 | |
| 24 | // Check a static device variables referenced only by device functions and kernels |
| 25 | // is not externalized. |
| 26 | // DEV-DAG: @_ZL2x2 = internal addrspace(1) global i32 0 |
| 27 | static __device__ int x2; |
| 28 | |
| 29 | // Check a static device variable referenced by host device function is externalized. |
| 30 | // DEV-DAG: @_ZL2x3 = addrspace(1) externally_initialized global i32 0 |
| 31 | static __device__ int x3; |
| 32 | |
| 33 | // Check a static device variable referenced in file scope is externalized. |
| 34 | // DEV-DAG: @_ZL2x4 = addrspace(1) externally_initialized global i32 0 |
| 35 | static __device__ int x4; |
| 36 | int& x4_ref = x4; |
| 37 | |
| 38 | // Check a static device variable in anonymous namespace. |
| 39 | // DEV-DAG: @_ZN12_GLOBAL__N_12x5E = addrspace(1) externally_initialized global i32 0 |
| 40 | namespace { |
| 41 | static __device__ int x5; |
| 42 | } |
| 43 | |
| 44 | // Check a static constant variable referenced by host is externalized. |
| 45 | // DEV-DAG: @_ZL1y = addrspace(4) externally_initialized global i32 0 |
| 46 | // HOST-DAG: @_ZL1y = internal global i32 undef |
| 47 | // HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y\00" |
| 48 | |
| 49 | static __constant__ int y; |
| 50 | |
| 51 | // Test static host variable, which should not be externalized nor registered. |
| 52 | // HOST-DAG: @_ZL1z = internal global i32 0 |
| 53 | // DEV-NOT: @_ZL1z |
| 54 | static int z; |
| 55 | |
Yaxun (Sam) Liu | fb04d7b | 2020-08-10 16:38:10 -0400 | [diff] [blame] | 56 | // Test implicit static constant variable, which should not be externalized. |
| 57 | // HOST-DAG: @_ZL2z2 = internal constant i32 456 |
| 58 | // DEV-DAG: @_ZL2z2 = internal addrspace(4) constant i32 456 |
| 59 | |
| 60 | static constexpr int z2 = 456; |
| 61 | |
Yaxun (Sam) Liu | 45f2a56 | 2020-08-04 12:13:16 -0400 | [diff] [blame] | 62 | // Test static device variable in inline function, which should not be |
| 63 | // externalized nor registered. |
| 64 | // DEV-DAG: @_ZZ6devfunPPKiE1p = linkonce_odr addrspace(4) constant i32 2, comdat |
| 65 | |
Yaxun (Sam) Liu | 301e233 | 2020-09-22 12:52:07 -0400 | [diff] [blame^] | 66 | // Check a static device variable referenced by host function only is externalized. |
| 67 | // DEV-DAG: @_ZL1w = addrspace(1) externally_initialized global i32 0 |
| 68 | // HOST-DAG: @_ZL1w = internal global i32 undef |
| 69 | // HOST-DAG: @[[DEVNAMEW:[0-9]+]] = {{.*}}c"_ZL1w\00" |
| 70 | |
| 71 | static __device__ int w; |
| 72 | |
Yaxun (Sam) Liu | 45f2a56 | 2020-08-04 12:13:16 -0400 | [diff] [blame] | 73 | inline __device__ void devfun(const int ** b) { |
| 74 | const static int p = 2; |
| 75 | b[0] = &p; |
| 76 | b[1] = &x2; |
| 77 | } |
| 78 | |
| 79 | __global__ void kernel(int *a, const int **b) { |
| 80 | const static int w = 1; |
| 81 | a[0] = x; |
| 82 | a[1] = y; |
| 83 | a[2] = x2; |
| 84 | a[3] = x3; |
| 85 | a[4] = x4; |
| 86 | a[5] = x5; |
| 87 | b[0] = &w; |
Yaxun (Sam) Liu | fb04d7b | 2020-08-10 16:38:10 -0400 | [diff] [blame] | 88 | b[1] = &z2; |
Yaxun (Sam) Liu | 45f2a56 | 2020-08-04 12:13:16 -0400 | [diff] [blame] | 89 | devfun(b); |
| 90 | } |
| 91 | |
| 92 | __host__ __device__ void hdf(int *a) { |
| 93 | a[0] = x3; |
| 94 | } |
| 95 | |
| 96 | int* getDeviceSymbol(int *x); |
| 97 | |
Yaxun (Sam) Liu | fb04d7b | 2020-08-10 16:38:10 -0400 | [diff] [blame] | 98 | void foo(const int **a) { |
Yaxun (Sam) Liu | 45f2a56 | 2020-08-04 12:13:16 -0400 | [diff] [blame] | 99 | getDeviceSymbol(&x); |
| 100 | getDeviceSymbol(&x5); |
| 101 | getDeviceSymbol(&y); |
Yaxun (Sam) Liu | 301e233 | 2020-09-22 12:52:07 -0400 | [diff] [blame^] | 102 | getDeviceSymbol(&w); |
Yaxun (Sam) Liu | 45f2a56 | 2020-08-04 12:13:16 -0400 | [diff] [blame] | 103 | z = 123; |
Yaxun (Sam) Liu | fb04d7b | 2020-08-10 16:38:10 -0400 | [diff] [blame] | 104 | a[0] = &z2; |
Yaxun (Sam) Liu | 45f2a56 | 2020-08-04 12:13:16 -0400 | [diff] [blame] | 105 | } |
| 106 | |
| 107 | // HOST: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]] |
| 108 | // HOST: __hipRegisterVar({{.*}}@_ZL1y {{.*}}@[[DEVNAMEY]] |
Yaxun (Sam) Liu | 301e233 | 2020-09-22 12:52:07 -0400 | [diff] [blame^] | 109 | // HOST: __hipRegisterVar({{.*}}@_ZL1w {{.*}}@[[DEVNAMEW]] |
Yaxun (Sam) Liu | 45f2a56 | 2020-08-04 12:13:16 -0400 | [diff] [blame] | 110 | // HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w |
| 111 | // HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p |