1 // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s 2 3 // Verifies Clang emits correct address spaces and addrspacecast instructions 4 // for CUDA code. 5 6 #include "Inputs/cuda.h" 7 8 // CHECK: @i = addrspace(1) externally_initialized global 9 __device__ int i; 10 11 // CHECK: @j = addrspace(4) externally_initialized global 12 __constant__ int j; 13 14 // CHECK: @k = addrspace(3) global 15 __shared__ int k; 16 17 struct MyStruct { 18 int data1; 19 int data2; 20 }; 21 22 // CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct zeroinitializer 23 // CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float 0.000000e+00 24 // CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] zeroinitializer 25 // CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00 26 // CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00 27 // CHECK: @b = addrspace(3) global float undef 28 // CHECK: @c = addrspace(3) global %struct.c undef 29 // CHECK @d = addrspace(3) global %struct.d undef 30 31 __device__ void foo() { 32 // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*) 33 i++; 34 35 // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @j to i32*) 36 j++; 37 38 // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*) 39 k++; 40 41 static int li; 42 // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*) 43 li++; 44 45 __constant__ int lj; 46 // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*) 47 lj++; 48 49 __shared__ int lk; 50 // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*) 51 lk++; 52 } 53 54 __device__ void func0() { 55 __shared__ MyStruct a; 56 MyStruct *ap = &a; // composite type 57 ap->data1 = 1; 58 ap->data2 = 2; 59 } 60 // CHECK: define void @_Z5func0v() 61 // CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %ap 62 63 __device__ void callee(float *ap) { 64 *ap = 1.0f; 65 } 66 67 __device__ void func1() { 68 __shared__ float a; 69 callee(&a); // implicit cast from parameters 70 } 71 // CHECK: define void @_Z5func1v() 72 // CHECK: call void @_Z6calleePf(float* addrspacecast (float addrspace(3)* @_ZZ5func1vE1a to float*)) 73 74 __device__ void func2() { 75 __shared__ float a[256]; 76 float *ap = &a[128]; // implicit cast from a decayed array 77 *ap = 1.0f; 78 } 79 // CHECK: define void @_Z5func2v() 80 // CHECK: store float* getelementptr inbounds ([256 x float], [256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i32 0, i32 128), float** %ap 81 82 __device__ void func3() { 83 __shared__ float a; 84 float *ap = reinterpret_cast<float *>(&a); // explicit cast 85 *ap = 1.0f; 86 } 87 // CHECK: define void @_Z5func3v() 88 // CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %ap 89 90 __device__ void func4() { 91 __shared__ float a; 92 float *ap = (float *)&a; // explicit c-style cast 93 *ap = 1.0f; 94 } 95 // CHECK: define void @_Z5func4v() 96 // CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %ap 97 98 __shared__ float b; 99 100 __device__ float *func5() { 101 return &b; // implicit cast from a return value 102 } 103 // CHECK: define float* @_Z5func5v() 104 // CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*) 105 106 struct StructWithCtor { 107 __device__ StructWithCtor(): data(1) {} 108 __device__ StructWithCtor(const StructWithCtor &second): data(second.data) {} 109 __device__ int getData() { return data; } 110 int data; 111 }; 112 113 __device__ int construct_shared_struct() { 114 // CHECK-LABEL: define i32 @_Z23construct_shared_structv() 115 __shared__ StructWithCtor s; 116 // CHECK: call void @_ZN14StructWithCtorC1Ev(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*)) 117 __shared__ StructWithCtor t(s); 118 // CHECK: call void @_ZN14StructWithCtorC1ERKS_(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*), %struct.StructWithCtor* dereferenceable(4) addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*)) 119 return t.getData(); 120 // CHECK: call i32 @_ZN14StructWithCtor7getDataEv(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*)) 121 } 122 123 // Make sure we allow __shared__ structures with default or empty constructors. 124 struct c { 125 int i; 126 }; 127 __shared__ struct c c; 128 129 struct d { 130 int i; 131 d() {} 132 }; 133 __shared__ struct d d; 134
