1 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \ 2 // RUN: -emit-llvm -o - -x hip %s | FileCheck \ 3 // RUN: -check-prefixes=COMMON,DEV,NORDC-D %s 4 5 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \ 6 // RUN: -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.dev 7 // RUN: cat %t.dev | FileCheck -check-prefixes=COMMON,DEV,RDC-D %s 8 9 // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \ 10 // RUN: -emit-llvm -o - -x hip %s | FileCheck \ 11 // RUN: -check-prefixes=COMMON,HOST,NORDC %s 12 13 // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \ 14 // RUN: -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.host 15 // RUN: cat %t.host | FileCheck -check-prefixes=COMMON,HOST,RDC %s 16 17 // Check device and host compilation use the same postfix for static 18 // variable name. 19 20 // RUN: cat %t.dev %t.host | FileCheck -check-prefix=POSTFIX %s 21 22 #include "Inputs/cuda.h" 23 24 struct vec { 25 float x,y,z; 26 }; 27 28 // DEV-DAG: @x.managed = addrspace(1) externally_initialized global i32 1, align 4 29 // DEV-DAG: @x = addrspace(1) externally_initialized global ptr addrspace(1) null 30 // NORDC-DAG: @x.managed = internal global i32 1 31 // RDC-DAG: @x.managed = global i32 1 32 // NORDC-DAG: @x = internal externally_initialized global ptr null 33 // RDC-DAG: @x = externally_initialized global ptr null 34 // HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"x\00" 35 __managed__ int x = 1; 36 37 // DEV-DAG: @v.managed = addrspace(1) externally_initialized global [100 x %struct.vec] zeroinitializer, align 4 38 // DEV-DAG: @v = addrspace(1) externally_initialized global ptr addrspace(1) null 39 __managed__ vec v[100]; 40 41 // DEV-DAG: @v2.managed = addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> <{ %struct.vec { float 1.000000e+00, float 1.000000e+00, float 1.000000e+00 }, [99 x %struct.vec] zeroinitializer }>, align 4 42 // DEV-DAG: @v2 = addrspace(1) externally_initialized global ptr addrspace(1) null 43 __managed__ vec v2[100] = {{1, 1, 1}}; 44 45 // DEV-DAG: @ex.managed = external addrspace(1) global i32, align 4 46 // DEV-DAG: @ex = external addrspace(1) externally_initialized global ptr addrspace(1) 47 // HOST-DAG: @ex.managed = external global i32 48 // HOST-DAG: @ex = external externally_initialized global ptr 49 extern __managed__ int ex; 50 51 // NORDC-D-DAG: @_ZL2sx.managed = addrspace(1) externally_initialized global i32 1, align 4 52 // NORDC-D-DAG: @_ZL2sx = addrspace(1) externally_initialized global ptr addrspace(1) null 53 // RDC-D-DAG: @_ZL2sx.static.[[HASH:.*]].managed = addrspace(1) externally_initialized global i32 1, align 4 54 // RDC-D-DAG: @_ZL2sx.static.[[HASH]] = addrspace(1) externally_initialized global ptr addrspace(1) null 55 // HOST-DAG: @_ZL2sx.managed = internal global i32 1 56 // HOST-DAG: @_ZL2sx = internal externally_initialized global ptr null 57 // NORDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx\00" 58 // RDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH:.*]]\00" 59 60 // POSTFIX: @_ZL2sx.static.[[HASH:.*]] = addrspace(1) externally_initialized global ptr addrspace(1) null 61 // POSTFIX: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH]]\00" 62 static __managed__ int sx = 1; 63 64 // DEV-DAG: @llvm.compiler.used 65 // DEV-SAME-DAG: @x.managed 66 // DEV-SAME-DAG: @x 67 // DEV-SAME-DAG: @v.managed 68 // DEV-SAME-DAG: @v 69 // DEV-SAME-DAG: @_ZL2sx.managed 70 // DEV-SAME-DAG: @_ZL2sx 71 72 // Force ex and sx mitted in device compilation. 73 __global__ void foo(int *z) { 74 *z = x + ex + sx; 75 v[1].x = 2; 76 } 77 78 // Force ex and sx emitted in host compilatioin. 79 int foo2() { 80 return ex + sx; 81 } 82 83 // COMMON-LABEL: define {{.*}}@_Z4loadv() 84 // DEV: %ld.managed = load ptr addrspace(1), ptr addrspace(1) @x, align 4 85 // DEV: %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr 86 // DEV: %1 = load i32, ptr %0, align 4 87 // DEV: ret i32 %1 88 // HOST: %ld.managed = load ptr, ptr @x, align 4 89 // HOST: %0 = load i32, ptr %ld.managed, align 4 90 // HOST: ret i32 %0 91 __device__ __host__ int load() { 92 return x; 93 } 94 95 // COMMON-LABEL: define {{.*}}@_Z5storev() 96 // DEV: %ld.managed = load ptr addrspace(1), ptr addrspace(1) @x, align 4 97 // DEV: %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr 98 // DEV: store i32 2, ptr %0, align 4 99 // HOST: %ld.managed = load ptr, ptr @x, align 4 100 // HOST: store i32 2, ptr %ld.managed, align 4 101 __device__ __host__ void store() { 102 x = 2; 103 } 104 105 // COMMON-LABEL: define {{.*}}@_Z10addr_takenv() 106 // DEV: %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr 107 // DEV: store ptr %0, ptr %p.ascast, align 8 108 // DEV: %1 = load ptr, ptr %p.ascast, align 8 109 // DEV: store i32 3, ptr %1, align 4 110 // HOST: %ld.managed = load ptr, ptr @x, align 4 111 // HOST: store ptr %ld.managed, ptr %p, align 8 112 // HOST: %0 = load ptr, ptr %p, align 8 113 // HOST: store i32 3, ptr %0, align 4 114 __device__ __host__ void addr_taken() { 115 int *p = &x; 116 *p = 3; 117 } 118 119 // HOST-LABEL: define {{.*}}@_Z5load2v() 120 // HOST: %ld.managed = load ptr, ptr @v, align 16 121 // HOST: %0 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed, i64 0, i64 1 122 // HOST: %1 = load float, ptr %0, align 4 123 // HOST: ret float %1 124 __device__ __host__ float load2() { 125 return v[1].x; 126 } 127 128 // HOST-LABEL: define {{.*}}@_Z5load3v() 129 // HOST: %ld.managed = load ptr, ptr @v2, align 16 130 // HOST: %0 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed, i64 0, i64 1 131 // HOST: %1 = getelementptr inbounds nuw %struct.vec, ptr %0, i32 0, i32 1 132 // HOST: %2 = load float, ptr %1, align 4 133 // HOST: ret float %2 134 float load3() { 135 return v2[1].y; 136 } 137 138 // HOST-LABEL: define {{.*}}@_Z11addr_taken2v() 139 // HOST: %ld.managed = load ptr, ptr @v, align 16 140 // HOST: %0 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed, i64 0, i64 1 141 // HOST: %1 = ptrtoint ptr %0 to i64 142 // HOST: %ld.managed1 = load ptr, ptr @v2, align 16 143 // HOST: %2 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed1, i64 0, i64 1 144 // HOST: %3 = getelementptr inbounds nuw %struct.vec, ptr %2, i32 0, i32 1 145 // HOST: %4 = ptrtoint ptr %3 to i64 146 // HOST: %5 = sub i64 %4, %1 147 // HOST: %sub.ptr.div = sdiv exact i64 %5, 4 148 // HOST: %conv = sitofp i64 %sub.ptr.div to float 149 // HOST: ret float %conv 150 float addr_taken2() { 151 return (float)reinterpret_cast<long>(&(v2[1].y)-&(v[1].x)); 152 } 153 154 // COMMON-LABEL: define {{.*}}@_Z5load4v() 155 // DEV: %ld.managed = load ptr addrspace(1), ptr addrspace(1) @ex, align 4 156 // DEV: %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr 157 // DEV: %1 = load i32, ptr %0, align 4 158 // DEV: ret i32 %1 159 // HOST: %ld.managed = load ptr, ptr @ex, align 4 160 // HOST: %0 = load i32, ptr %ld.managed, align 4 161 // HOST: ret i32 %0 162 __device__ __host__ int load4() { 163 return ex; 164 } 165 166 // HOST-DAG: __hipRegisterManagedVar({{.*}}, ptr @x, ptr @x.managed, ptr @[[DEVNAMEX]], i64 4, i32 4) 167 // HOST-DAG: __hipRegisterManagedVar({{.*}}, ptr @_ZL2sx, ptr @_ZL2sx.managed, ptr @[[DEVNAMESX]] 168 // HOST-NOT: __hipRegisterManagedVar({{.*}}, ptr @ex, ptr @ex.managed 169 // HOST-DAG: declare void @__hipRegisterManagedVar(ptr, ptr, ptr, ptr, i64, i32) 170