| 1 | // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s |
|---|
| 2 | // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple amdgcn | FileCheck %s |
|---|
| 3 | |
|---|
| 4 | // Verifies Clang emits correct address spaces and addrspacecast instructions |
|---|
| 5 | // for CUDA code. |
|---|
| 6 | |
|---|
| 7 | #include "Inputs/cuda.h" |
|---|
| 8 | |
|---|
| 9 | // CHECK: @i = addrspace(1) externally_initialized global |
|---|
| 10 | __device__ int i; |
|---|
| 11 | |
|---|
| 12 | // CHECK: @j = addrspace(4) externally_initialized global |
|---|
| 13 | __constant__ int j; |
|---|
| 14 | |
|---|
| 15 | // CHECK: @k = addrspace(3) global |
|---|
| 16 | __shared__ int k; |
|---|
| 17 | |
|---|
| 18 | struct MyStruct { |
|---|
| 19 | int data1; |
|---|
| 20 | int data2; |
|---|
| 21 | }; |
|---|
| 22 | |
|---|
| 23 | // CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct undef |
|---|
| 24 | // CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float undef |
|---|
| 25 | // CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] undef |
|---|
| 26 | // CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float undef |
|---|
| 27 | // CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float undef |
|---|
| 28 | // CHECK: @b = addrspace(3) global float undef |
|---|
| 29 | |
|---|
| 30 | __device__ void foo() { |
|---|
| 31 | // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*) |
|---|
| 32 | i++; |
|---|
| 33 | |
|---|
| 34 | // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @j to i32*) |
|---|
| 35 | j++; |
|---|
| 36 | |
|---|
| 37 | // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*) |
|---|
| 38 | k++; |
|---|
| 39 | |
|---|
| 40 | __shared__ int lk; |
|---|
| 41 | // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*) |
|---|
| 42 | lk++; |
|---|
| 43 | } |
|---|
| 44 | |
|---|
| 45 | __device__ void func0() { |
|---|
| 46 | __shared__ MyStruct a; |
|---|
| 47 | MyStruct *ap = &a; // composite type |
|---|
| 48 | ap->data1 = 1; |
|---|
| 49 | ap->data2 = 2; |
|---|
| 50 | } |
|---|
| 51 | // CHECK: define void @_Z5func0v() |
|---|
| 52 | // CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %{{.*}} |
|---|
| 53 | |
|---|
| 54 | __device__ void callee(float *ap) { |
|---|
| 55 | *ap = 1.0f; |
|---|
| 56 | } |
|---|
| 57 | |
|---|
| 58 | __device__ void func1() { |
|---|
| 59 | __shared__ float a; |
|---|
| 60 | callee(&a); // implicit cast from parameters |
|---|
| 61 | } |
|---|
| 62 | // CHECK: define void @_Z5func1v() |
|---|
| 63 | // CHECK: call void @_Z6calleePf(float* addrspacecast (float addrspace(3)* @_ZZ5func1vE1a to float*)) |
|---|
| 64 | |
|---|
| 65 | __device__ void func2() { |
|---|
| 66 | __shared__ float a[256]; |
|---|
| 67 | float *ap = &a[128]; // implicit cast from a decayed array |
|---|
| 68 | *ap = 1.0f; |
|---|
| 69 | } |
|---|
| 70 | // CHECK: define void @_Z5func2v() |
|---|
| 71 | // CHECK: store float* getelementptr inbounds ([256 x float], [256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i{{32|64}} 0, i{{32|64}} 128), float** %{{.*}} |
|---|
| 72 | |
|---|
| 73 | __device__ void func3() { |
|---|
| 74 | __shared__ float a; |
|---|
| 75 | float *ap = reinterpret_cast<float *>(&a); // explicit cast |
|---|
| 76 | *ap = 1.0f; |
|---|
| 77 | } |
|---|
| 78 | // CHECK: define void @_Z5func3v() |
|---|
| 79 | // CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %{{.*}} |
|---|
| 80 | |
|---|
| 81 | __device__ void func4() { |
|---|
| 82 | __shared__ float a; |
|---|
| 83 | float *ap = (float *)&a; // explicit c-style cast |
|---|
| 84 | *ap = 1.0f; |
|---|
| 85 | } |
|---|
| 86 | // CHECK: define void @_Z5func4v() |
|---|
| 87 | // CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %{{.*}} |
|---|
| 88 | |
|---|
| 89 | __shared__ float b; |
|---|
| 90 | |
|---|
| 91 | __device__ float *func5() { |
|---|
| 92 | return &b; // implicit cast from a return value |
|---|
| 93 | } |
|---|
| 94 | // CHECK: define float* @_Z5func5v() |
|---|
| 95 | // CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*) |
|---|
| 96 | |
|---|