| Justin Holewinski | 83e9668 | 2012-05-24 17:43:12 +0000 | [diff] [blame] | 1 | // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s |
| Peter Collingbourne | f44bdf9 | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 2 | |
| 3 | #include "../SemaCUDA/cuda.h" |
| 4 | |
| Justin Holewinski | 83e9668 | 2012-05-24 17:43:12 +0000 | [diff] [blame] | 5 | // CHECK: @i = addrspace(1) global |
| Peter Collingbourne | f44bdf9 | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 6 | __device__ int i; |
| 7 | |
| Justin Holewinski | 83e9668 | 2012-05-24 17:43:12 +0000 | [diff] [blame] | 8 | // CHECK: @j = addrspace(4) global |
| Peter Collingbourne | f44bdf9 | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 9 | __constant__ int j; |
| 10 | |
| Justin Holewinski | 83e9668 | 2012-05-24 17:43:12 +0000 | [diff] [blame] | 11 | // CHECK: @k = addrspace(3) global |
| Peter Collingbourne | f44bdf9 | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 12 | __shared__ int k; |
| 13 | |
| 14 | __device__ void foo() { |
| Matt Arsenault | 00e65b2 | 2013-11-15 02:19:52 +0000 | [diff] [blame] | 15 | // CHECK: load i32* addrspacecast (i32 addrspace(1)* @i to i32*) |
| Peter Collingbourne | f44bdf9 | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 16 | i++; |
| 17 | |
| Matt Arsenault | 00e65b2 | 2013-11-15 02:19:52 +0000 | [diff] [blame] | 18 | // CHECK: load i32* addrspacecast (i32 addrspace(4)* @j to i32*) |
| Peter Collingbourne | f44bdf9 | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 19 | j++; |
| 20 | |
| Matt Arsenault | 00e65b2 | 2013-11-15 02:19:52 +0000 | [diff] [blame] | 21 | // CHECK: load i32* addrspacecast (i32 addrspace(3)* @k to i32*) |
| Peter Collingbourne | f44bdf9 | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 22 | k++; |
| Peter Collingbourne | ee0502d | 2012-08-28 20:37:10 +0000 | [diff] [blame] | 23 | |
| 24 | static int li; |
| 25 | // CHECK: load i32 addrspace(1)* @_ZZ3foovE2li |
| 26 | li++; |
| Peter Collingbourne | c6b0857 | 2012-08-28 20:37:50 +0000 | [diff] [blame] | 27 | |
| 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 Collingbourne | f44bdf9 | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 35 | } |
| 36 | |