1; RUN: llc < %s -march=ptx32 | FileCheck %s 2 3;CHECK: .extern .global .b8 array_i16[20]; 4@array_i16 = external global [10 x i16] 5 6;CHECK: .extern .const .b8 array_constant_i16[20]; 7@array_constant_i16 = external addrspace(1) constant [10 x i16] 8 9;CHECK: .extern .shared .b8 array_shared_i16[20]; 10@array_shared_i16 = external addrspace(4) global [10 x i16] 11 12;CHECK: .extern .global .b8 array_i32[40]; 13@array_i32 = external global [10 x i32] 14 15;CHECK: .extern .const .b8 array_constant_i32[40]; 16@array_constant_i32 = external addrspace(1) constant [10 x i32] 17 18;CHECK: .extern .shared .b8 array_shared_i32[40]; 19@array_shared_i32 = external addrspace(4) global [10 x i32] 20 21;CHECK: .extern .global .b8 array_i64[80]; 22@array_i64 = external global [10 x i64] 23 24;CHECK: .extern .const .b8 array_constant_i64[80]; 25@array_constant_i64 = external addrspace(1) constant [10 x i64] 26 27;CHECK: .extern .shared .b8 array_shared_i64[80]; 28@array_shared_i64 = external addrspace(4) global [10 x i64] 29 30;CHECK: .extern .global .b8 array_float[40]; 31@array_float = external global [10 x float] 32 33;CHECK: .extern .const .b8 array_constant_float[40]; 34@array_constant_float = external addrspace(1) constant [10 x float] 35 36;CHECK: .extern .shared .b8 array_shared_float[40]; 37@array_shared_float = external addrspace(4) global [10 x float] 38 39;CHECK: .extern .global .b8 array_double[80]; 40@array_double = external global [10 x double] 41 42;CHECK: .extern .const .b8 array_constant_double[80]; 43@array_constant_double = external addrspace(1) constant [10 x double] 44 45;CHECK: .extern .shared .b8 array_shared_double[80]; 46@array_shared_double = external addrspace(4) global [10 x double] 47 48 49define ptx_device void @t1_u16(i16* %p, i16 %x) { 50entry: 51;CHECK: st.global.u16 [%r{{[0-9]+}}], %rh{{[0-9]+}}; 52;CHECK: ret; 53 store i16 %x, i16* %p 54 ret void 55} 56 57define ptx_device void @t1_u32(i32* %p, i32 %x) { 58entry: 59;CHECK: st.global.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}; 60;CHECK: ret; 61 store i32 %x, i32* %p 62 ret void 63} 64 65define ptx_device void @t1_u64(i64* %p, i64 %x) { 66entry: 67;CHECK: st.global.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}}; 68;CHECK: ret; 69 store i64 %x, i64* %p 70 ret void 71} 72 73define ptx_device void @t1_f32(float* %p, float %x) { 74entry: 75;CHECK: st.global.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}; 76;CHECK: ret; 77 store float %x, float* %p 78 ret void 79} 80 81define ptx_device void @t1_f64(double* %p, double %x) { 82entry: 83;CHECK: st.global.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}}; 84;CHECK: ret; 85 store double %x, double* %p 86 ret void 87} 88 89define ptx_device void @t2_u16(i16* %p, i16 %x) { 90entry: 91;CHECK: st.global.u16 [%r{{[0-9]+}}+2], %rh{{[0-9]+}}; 92;CHECK: ret; 93 %i = getelementptr i16* %p, i32 1 94 store i16 %x, i16* %i 95 ret void 96} 97 98define ptx_device void @t2_u32(i32* %p, i32 %x) { 99entry: 100;CHECK: st.global.u32 [%r{{[0-9]+}}+4], %r{{[0-9]+}}; 101;CHECK: ret; 102 %i = getelementptr i32* %p, i32 1 103 store i32 %x, i32* %i 104 ret void 105} 106 107define ptx_device void @t2_u64(i64* %p, i64 %x) { 108entry: 109;CHECK: st.global.u64 [%r{{[0-9]+}}+8], %rd{{[0-9]+}}; 110;CHECK: ret; 111 %i = getelementptr i64* %p, i32 1 112 store i64 %x, i64* %i 113 ret void 114} 115 116define ptx_device void @t2_f32(float* %p, float %x) { 117entry: 118;CHECK: st.global.f32 [%r{{[0-9]+}}+4], %f{{[0-9]+}}; 119;CHECK: ret; 120 %i = getelementptr float* %p, i32 1 121 store float %x, float* %i 122 ret void 123} 124 125define ptx_device void @t2_f64(double* %p, double %x) { 126entry: 127;CHECK: st.global.f64 [%r{{[0-9]+}}+8], %fd{{[0-9]+}}; 128;CHECK: ret; 129 %i = getelementptr double* %p, i32 1 130 store double %x, double* %i 131 ret void 132} 133 134define ptx_device void @t3_u16(i16* %p, i32 %q, i16 %x) { 135entry: 136;CHECK: shl.b32 %r[[R0:[0-9]+]], %r{{[0-9]+}}, 1; 137;CHECK: add.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r[[R0]]; 138;CHECK: st.global.u16 [%r{{[0-9]+}}], %rh{{[0-9]+}}; 139;CHECK: ret; 140 %i = getelementptr i16* %p, i32 %q 141 store i16 %x, i16* %i 142 ret void 143} 144 145define ptx_device void @t3_u32(i32* %p, i32 %q, i32 %x) { 146entry: 147;CHECK: shl.b32 %r[[R0:[0-9]+]], %r{{[0-9]+}}, 2; 148;CHECK: add.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r[[R0]]; 149;CHECK: st.global.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}; 150;CHECK: ret; 151 %i = getelementptr i32* %p, i32 %q 152 store i32 %x, i32* %i 153 ret void 154} 155 156define ptx_device void @t3_u64(i64* %p, i32 %q, i64 %x) { 157entry: 158;CHECK: shl.b32 %r[[R0:[0-9]+]], %r{{[0-9]+}}, 3; 159;CHECK: add.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r[[R0]]; 160;CHECK: st.global.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}}; 161;CHECK: ret; 162 %i = getelementptr i64* %p, i32 %q 163 store i64 %x, i64* %i 164 ret void 165} 166 167define ptx_device void @t3_f32(float* %p, i32 %q, float %x) { 168entry: 169;CHECK: shl.b32 %r[[R0:[0-9]+]], %r{{[0-9]+}}, 2; 170;CHECK: add.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r[[R0]]; 171;CHECK: st.global.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}; 172;CHECK: ret; 173 %i = getelementptr float* %p, i32 %q 174 store float %x, float* %i 175 ret void 176} 177 178define ptx_device void @t3_f64(double* %p, i32 %q, double %x) { 179entry: 180;CHECK: shl.b32 %r[[R0:[0-9]+]], %r{{[0-9]+}}, 3; 181;CHECK: add.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r[[R0]]; 182;CHECK: st.global.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}}; 183;CHECK: ret; 184 %i = getelementptr double* %p, i32 %q 185 store double %x, double* %i 186 ret void 187} 188 189define ptx_device void @t4_global_u16(i16 %x) { 190entry: 191;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i16; 192;CHECK: st.global.u16 [%r[[R0]]], %rh{{[0-9]+}}; 193;CHECK: ret; 194 %i = getelementptr [10 x i16]* @array_i16, i16 0, i16 0 195 store i16 %x, i16* %i 196 ret void 197} 198 199define ptx_device void @t4_global_u32(i32 %x) { 200entry: 201;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i32; 202;CHECK: st.global.u32 [%r[[R0]]], %r{{[0-9]+}}; 203;CHECK: ret; 204 %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 0 205 store i32 %x, i32* %i 206 ret void 207} 208 209define ptx_device void @t4_global_u64(i64 %x) { 210entry: 211;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i64; 212;CHECK: st.global.u64 [%r[[R0]]], %rd{{[0-9]+}}; 213;CHECK: ret; 214 %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 0 215 store i64 %x, i64* %i 216 ret void 217} 218 219define ptx_device void @t4_global_f32(float %x) { 220entry: 221;CHECK: mov.u32 %r[[R0:[0-9]+]], array_float; 222;CHECK: st.global.f32 [%r[[R0]]], %f{{[0-9]+}}; 223;CHECK: ret; 224 %i = getelementptr [10 x float]* @array_float, i32 0, i32 0 225 store float %x, float* %i 226 ret void 227} 228 229define ptx_device void @t4_global_f64(double %x) { 230entry: 231;CHECK: mov.u32 %r[[R0:[0-9]+]], array_double; 232;CHECK: st.global.f64 [%r[[R0]]], %fd{{[0-9]+}}; 233;CHECK: ret; 234 %i = getelementptr [10 x double]* @array_double, i32 0, i32 0 235 store double %x, double* %i 236 ret void 237} 238 239define ptx_device void @t4_shared_u16(i16 %x) { 240entry: 241;CHECK: mov.u32 %r[[R0:[0-9]+]], array_shared_i16; 242;CHECK: st.shared.u16 [%r[[R0]]], %rh{{[0-9]+}}; 243;CHECK: ret; 244 %i = getelementptr [10 x i16] addrspace(4)* @array_shared_i16, i32 0, i32 0 245 store i16 %x, i16 addrspace(4)* %i 246 ret void 247} 248 249define ptx_device void @t4_shared_u32(i32 %x) { 250entry: 251;CHECK: mov.u32 %r[[R0:[0-9]+]], array_shared_i32; 252;CHECK: st.shared.u32 [%r[[R0]]], %r{{[0-9]+}}; 253;CHECK: ret; 254 %i = getelementptr [10 x i32] addrspace(4)* @array_shared_i32, i32 0, i32 0 255 store i32 %x, i32 addrspace(4)* %i 256 ret void 257} 258 259define ptx_device void @t4_shared_u64(i64 %x) { 260entry: 261;CHECK: mov.u32 %r[[R0:[0-9]+]], array_shared_i64; 262;CHECK: st.shared.u64 [%r[[R0]]], %rd{{[0-9]+}}; 263;CHECK: ret; 264 %i = getelementptr [10 x i64] addrspace(4)* @array_shared_i64, i32 0, i32 0 265 store i64 %x, i64 addrspace(4)* %i 266 ret void 267} 268 269define ptx_device void @t4_shared_f32(float %x) { 270entry: 271;CHECK: mov.u32 %r[[R0:[0-9]+]], array_shared_float; 272;CHECK: st.shared.f32 [%r[[R0]]], %f{{[0-9]+}}; 273;CHECK: ret; 274 %i = getelementptr [10 x float] addrspace(4)* @array_shared_float, i32 0, i32 0 275 store float %x, float addrspace(4)* %i 276 ret void 277} 278 279define ptx_device void @t4_shared_f64(double %x) { 280entry: 281;CHECK: mov.u32 %r[[R0:[0-9]+]], array_shared_double; 282;CHECK: st.shared.f64 [%r[[R0]]], %fd{{[0-9]+}}; 283;CHECK: ret; 284 %i = getelementptr [10 x double] addrspace(4)* @array_shared_double, i32 0, i32 0 285 store double %x, double addrspace(4)* %i 286 ret void 287} 288 289define ptx_device void @t5_u16(i16 %x) { 290entry: 291;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i16; 292;CHECK: st.global.u16 [%r[[R0]]+2], %rh{{[0-9]+}}; 293;CHECK: ret; 294 %i = getelementptr [10 x i16]* @array_i16, i32 0, i32 1 295 store i16 %x, i16* %i 296 ret void 297} 298 299define ptx_device void @t5_u32(i32 %x) { 300entry: 301;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i32; 302;CHECK: st.global.u32 [%r[[R0]]+4], %r{{[0-9]+}}; 303;CHECK: ret; 304 %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 1 305 store i32 %x, i32* %i 306 ret void 307} 308 309define ptx_device void @t5_u64(i64 %x) { 310entry: 311;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i64; 312;CHECK: st.global.u64 [%r[[R0]]+8], %rd{{[0-9]+}}; 313;CHECK: ret; 314 %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 1 315 store i64 %x, i64* %i 316 ret void 317} 318 319define ptx_device void @t5_f32(float %x) { 320entry: 321;CHECK: mov.u32 %r[[R0:[0-9]+]], array_float; 322;CHECK: st.global.f32 [%r[[R0]]+4], %f{{[0-9]+}}; 323;CHECK: ret; 324 %i = getelementptr [10 x float]* @array_float, i32 0, i32 1 325 store float %x, float* %i 326 ret void 327} 328 329define ptx_device void @t5_f64(double %x) { 330entry: 331;CHECK: mov.u32 %r[[R0:[0-9]+]], array_double; 332;CHECK: st.global.f64 [%r[[R0]]+8], %fd{{[0-9]+}}; 333;CHECK: ret; 334 %i = getelementptr [10 x double]* @array_double, i32 0, i32 1 335 store double %x, double* %i 336 ret void 337} 338