diff options
| author | Jingyue Wu <jingyue@google.com> | 2015-08-22 05:49:28 +0000 |
|---|---|---|
| committer | Jingyue Wu <jingyue@google.com> | 2015-08-22 05:49:28 +0000 |
| commit | e0848b6353721eb1b278a5bbea257bbf6316251e (patch) | |
| tree | defb9a45a6dbafd05e200832027d4f2baf7ec912 /test/CodeGenCUDA | |
| parent | 6ab2ad5a9b02f13441dcf8533ca33e8f7213a5e1 (diff) | |
[CUDA] Change initializer for CUDA device code based on CUDA documentation.
Summary:
According to CUDA documentation, global variables declared with __device__,
__constant__ can be initialized from host code, so mark them as
externally initialized. Because __shared__ variables cannot have an
initialization as part of their declaration and since the value maybe kept
across different kernel invocation, the value of __shared__ is effectively
undefined instead of zero initialized.
Wrongly using zero initializer may cause illegitimate optimization, e.g.
removing unused __constant__ variable because it's not updated in the device
code and the value is initialized with zero.
Test Plan: test/CodeGenCUDA/address-spaces.cu
Patch by Xuetian Weng
Reviewers: jholewinski, eliben, tra, jingyue
Subscribers: llvm-commits
Differential Revision: http://reviews.llvm.org/D12241
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@245786 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenCUDA')
| -rw-r--r-- | test/CodeGenCUDA/address-spaces.cu | 6 | ||||
| -rw-r--r-- | test/CodeGenCUDA/filter-decl.cu | 8 |
2 files changed, 7 insertions, 7 deletions
diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu index b0ef3558e2..5e7ff66b69 100644 --- a/test/CodeGenCUDA/address-spaces.cu +++ b/test/CodeGenCUDA/address-spaces.cu @@ -5,10 +5,10 @@ #include "Inputs/cuda.h" -// CHECK: @i = addrspace(1) global +// CHECK: @i = addrspace(1) externally_initialized global __device__ int i; -// CHECK: @j = addrspace(4) global +// CHECK: @j = addrspace(4) externally_initialized global __constant__ int j; // CHECK: @k = addrspace(3) global @@ -24,7 +24,7 @@ struct MyStruct { // CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] zeroinitializer // CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00 // CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00 -// CHECK: @b = addrspace(3) global float 0.000000e+00 +// CHECK: @b = addrspace(3) global float undef __device__ void foo() { // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*) diff --git a/test/CodeGenCUDA/filter-decl.cu b/test/CodeGenCUDA/filter-decl.cu index e69473f3e8..023ae61f3a 100644 --- a/test/CodeGenCUDA/filter-decl.cu +++ b/test/CodeGenCUDA/filter-decl.cu @@ -9,12 +9,12 @@ // CHECK-DEVICE-NOT: module asm "file scope asm is host only" __asm__("file scope asm is host only"); -// CHECK-HOST-NOT: constantdata = global -// CHECK-DEVICE: constantdata = global +// CHECK-HOST-NOT: constantdata = externally_initialized global +// CHECK-DEVICE: constantdata = externally_initialized global __constant__ char constantdata[256]; -// CHECK-HOST-NOT: devicedata = global -// CHECK-DEVICE: devicedata = global +// CHECK-HOST-NOT: devicedata = externally_initialized global +// CHECK-DEVICE: devicedata = externally_initialized global __device__ char devicedata[256]; // CHECK-HOST-NOT: shareddata = global |