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) global 9 __device__ int i; 10 11 // CHECK: @j = addrspace(4) 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 0.000000e+00 28 29 __device__ void foo() { 30 // CHECK: load i32* addrspacecast (i32 addrspace(1)* @i to i32*) 31 i++; 32 33 // CHECK: load i32* addrspacecast (i32 addrspace(4)* @j to i32*) 34 j++; 35 36 // CHECK: load i32* addrspacecast (i32 addrspace(3)* @k to i32*) 37 k++; 38 39 static int li; 40 // CHECK: load i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*) 41 li++; 42 43 __constant__ int lj; 44 // CHECK: load i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*) 45 lj++; 46 47 __shared__ int lk; 48 // CHECK: load i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*) 49 lk++; 50 } 51 52 __device__ void func0() { 53 __shared__ MyStruct a; 54 MyStruct *ap = &a; // composite type 55 ap->data1 = 1; 56 ap->data2 = 2; 57 } 58 // CHECK: define void @_Z5func0v() 59 // CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %ap 60 61 __device__ void callee(float *ap) { 62 *ap = 1.0f; 63 } 64 65 __device__ void func1() { 66 __shared__ float a; 67 callee(&a); // implicit cast from parameters 68 } 69 // CHECK: define void @_Z5func1v() 70 // CHECK: call void @_Z6calleePf(float* addrspacecast (float addrspace(3)* @_ZZ5func1vE1a to float*)) 71 72 __device__ void func2() { 73 __shared__ float a[256]; 74 float *ap = &a[128]; // implicit cast from a decayed array 75 *ap = 1.0f; 76 } 77 // CHECK: define void @_Z5func2v() 78 // CHECK: store float* getelementptr inbounds ([256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i32 0, i32 128), float** %ap 79 80 __device__ void func3() { 81 __shared__ float a; 82 float *ap = reinterpret_cast<float *>(&a); // explicit cast 83 *ap = 1.0f; 84 } 85 // CHECK: define void @_Z5func3v() 86 // CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %ap 87 88 __device__ void func4() { 89 __shared__ float a; 90 float *ap = (float *)&a; // explicit c-style cast 91 *ap = 1.0f; 92 } 93 // CHECK: define void @_Z5func4v() 94 // CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %ap 95 96 __shared__ float b; 97 98 __device__ float *func5() { 99 return &b; // implicit cast from a return value 100 } 101 // CHECK: define float* @_Z5func5v() 102 // CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*) 103
