blob: 9df7e3f4d262ac04d77ad349204987c8893e034c [file] [log] [blame]
Justin Holewinski2c585b92012-05-24 17:43:12 +00001// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
Peter Collingbourne4dc34eb2012-05-20 21:08:35 +00002
3#include "../SemaCUDA/cuda.h"
4
Justin Holewinski2c585b92012-05-24 17:43:12 +00005// CHECK: @i = addrspace(1) global
Peter Collingbourne4dc34eb2012-05-20 21:08:35 +00006__device__ int i;
7
Justin Holewinski2c585b92012-05-24 17:43:12 +00008// CHECK: @j = addrspace(4) global
Peter Collingbourne4dc34eb2012-05-20 21:08:35 +00009__constant__ int j;
10
Justin Holewinski2c585b92012-05-24 17:43:12 +000011// CHECK: @k = addrspace(3) global
Peter Collingbourne4dc34eb2012-05-20 21:08:35 +000012__shared__ int k;
13
14__device__ void foo() {
Justin Holewinski2c585b92012-05-24 17:43:12 +000015 // CHECK: load i32* bitcast (i32 addrspace(1)* @i to i32*)
Peter Collingbourne4dc34eb2012-05-20 21:08:35 +000016 i++;
17
Justin Holewinski2c585b92012-05-24 17:43:12 +000018 // CHECK: load i32* bitcast (i32 addrspace(4)* @j to i32*)
Peter Collingbourne4dc34eb2012-05-20 21:08:35 +000019 j++;
20
Justin Holewinski2c585b92012-05-24 17:43:12 +000021 // CHECK: load i32* bitcast (i32 addrspace(3)* @k to i32*)
Peter Collingbourne4dc34eb2012-05-20 21:08:35 +000022 k++;
Peter Collingbourne1aba7782012-08-28 20:37:10 +000023
24 static int li;
25 // CHECK: load i32 addrspace(1)* @_ZZ3foovE2li
26 li++;
Peter Collingbournec0c00662012-08-28 20:37:50 +000027
28 __constant__ int lj;
29 // CHECK: load i32 addrspace(4)* @_ZZ3foovE2lj
30 lj++;
31
32 __shared__ int lk;
33 // CHECK: load i32 addrspace(3)* @_ZZ3foovE2lk
34 lk++;
Peter Collingbourne4dc34eb2012-05-20 21:08:35 +000035}
36