| Justin Holewinski | 2c585b9 | 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 | 4dc34eb | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 2 | |
| 3 | #include "../SemaCUDA/cuda.h" |
| 4 | |
| Justin Holewinski | 2c585b9 | 2012-05-24 17:43:12 +0000 | [diff] [blame^] | 5 | // CHECK: @i = addrspace(1) global |
| Peter Collingbourne | 4dc34eb | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 6 | __device__ int i; |
| 7 | |
| Justin Holewinski | 2c585b9 | 2012-05-24 17:43:12 +0000 | [diff] [blame^] | 8 | // CHECK: @j = addrspace(4) global |
| Peter Collingbourne | 4dc34eb | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 9 | __constant__ int j; |
| 10 | |
| Justin Holewinski | 2c585b9 | 2012-05-24 17:43:12 +0000 | [diff] [blame^] | 11 | // CHECK: @k = addrspace(3) global |
| Peter Collingbourne | 4dc34eb | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 12 | __shared__ int k; |
| 13 | |
| 14 | __device__ void foo() { |
| Justin Holewinski | 2c585b9 | 2012-05-24 17:43:12 +0000 | [diff] [blame^] | 15 | // CHECK: load i32* bitcast (i32 addrspace(1)* @i to i32*) |
| Peter Collingbourne | 4dc34eb | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 16 | i++; |
| 17 | |
| Justin Holewinski | 2c585b9 | 2012-05-24 17:43:12 +0000 | [diff] [blame^] | 18 | // CHECK: load i32* bitcast (i32 addrspace(4)* @j to i32*) |
| Peter Collingbourne | 4dc34eb | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 19 | j++; |
| 20 | |
| Justin Holewinski | 2c585b9 | 2012-05-24 17:43:12 +0000 | [diff] [blame^] | 21 | // CHECK: load i32* bitcast (i32 addrspace(3)* @k to i32*) |
| Peter Collingbourne | 4dc34eb | 2012-05-20 21:08:35 +0000 | [diff] [blame] | 22 | k++; |
| 23 | } |
| 24 | |