96 lines · plain
1// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s2// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple amdgcn | FileCheck %s3 4// Verifies Clang emits correct address spaces and addrspacecast instructions5// for CUDA code.6 7#include "Inputs/cuda.h"8 9// CHECK: @i ={{.*}} addrspace(1) externally_initialized global10__device__ int i;11 12// CHECK: @j ={{.*}} addrspace(4) externally_initialized constant13__constant__ int j;14 15// CHECK: @k ={{.*}} addrspace(3) global16__shared__ int k;17 18struct MyStruct {19 int data1;20 int data2;21};22 23// CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct undef24// CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float undef25// CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] undef26// CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float undef27// CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float undef28// CHECK: @b ={{.*}} addrspace(3) global float undef29 30__device__ void foo() {31 // CHECK: load i32, ptr addrspacecast (ptr addrspace(1) @i to ptr)32 i++;33 34 // CHECK: load i32, ptr addrspacecast (ptr addrspace(4) @j to ptr)35 j++;36 37 // CHECK: load i32, ptr addrspacecast (ptr addrspace(3) @k to ptr)38 k++;39 40 __shared__ int lk;41 // CHECK: load i32, ptr addrspacecast (ptr addrspace(3) @_ZZ3foovE2lk to ptr)42 lk++;43}44 45__device__ void func0() {46 __shared__ MyStruct a;47 MyStruct *ap = &a; // composite type48 ap->data1 = 1;49 ap->data2 = 2;50}51// CHECK: define{{.*}} void @_Z5func0v()52// CHECK: store ptr addrspacecast (ptr addrspace(3) @_ZZ5func0vE1a to ptr), ptr %{{.*}}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 parameters61}62// CHECK: define{{.*}} void @_Z5func1v()63// CHECK: call void @_Z6calleePf(ptr noundef addrspacecast (ptr addrspace(3) @_ZZ5func1vE1a to ptr))64 65__device__ void func2() {66 __shared__ float a[256];67 float *ap = &a[128]; // implicit cast from a decayed array68 *ap = 1.0f;69}70// CHECK: define{{.*}} void @_Z5func2v()71// CHECK: store ptr getelementptr inbounds ([256 x float], ptr addrspacecast (ptr addrspace(3) @_ZZ5func2vE1a to ptr), i{{32|64}} 0, i{{32|64}} 128), ptr %{{.*}}72 73__device__ void func3() {74 __shared__ float a;75 float *ap = reinterpret_cast<float *>(&a); // explicit cast76 *ap = 1.0f;77}78// CHECK: define{{.*}} void @_Z5func3v()79// CHECK: store ptr addrspacecast (ptr addrspace(3) @_ZZ5func3vE1a to ptr), ptr %{{.*}}80 81__device__ void func4() {82 __shared__ float a;83 float *ap = (float *)&a; // explicit c-style cast84 *ap = 1.0f;85}86// CHECK: define{{.*}} void @_Z5func4v()87// CHECK: store ptr addrspacecast (ptr addrspace(3) @_ZZ5func4vE1a to ptr), ptr %{{.*}}88 89__shared__ float b;90 91__device__ float *func5() {92 return &b; // implicit cast from a return value93}94// CHECK: define{{.*}} ptr @_Z5func5v()95// CHECK: ret ptr addrspacecast (ptr addrspace(3) @b to ptr)96