diff options
| author | Artem Belevich <tra@google.com> | 2018-12-22 01:11:09 +0000 |
|---|---|---|
| committer | Artem Belevich <tra@google.com> | 2018-12-22 01:11:09 +0000 |
| commit | 5ce73e9a7953d8b153974c3c5a1ed1330558a73f (patch) | |
| tree | 74de8f26960271103db27115354e1ac41b47da9a /test/CodeGenCUDA | |
| parent | bd282df269cc1b61ea433838d2d3d5b4316c0ace (diff) | |
| download | clang-5ce73e9a7953d8b153974c3c5a1ed1330558a73f.tar.gz | |
[CUDA] Treat extern global variable shadows same as regular extern vars.
This fixes compiler crash when we attempted to compile this code:
extern __device__ int data;
__device__ int data = 1;
Differential Revision: https://reviews.llvm.org/D56033
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@349981 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenCUDA')
| -rw-r--r-- | test/CodeGenCUDA/device-stub.cu | 19 |
1 files changed, 13 insertions, 6 deletions
diff --git a/test/CodeGenCUDA/device-stub.cu b/test/CodeGenCUDA/device-stub.cu index 7abb7ae3a0..ea45c391d2 100644 --- a/test/CodeGenCUDA/device-stub.cu +++ b/test/CodeGenCUDA/device-stub.cu @@ -42,13 +42,20 @@ int host_var; // ALL-DAG: @ext_host_var = external global i32 extern int ext_host_var; -// Shadows for external device-side variables are *definitions* of -// those variables. -// ALL-DAG: @ext_device_var = internal global i32 +// external device-side variables -> extern references to their shadows. +// ALL-DAG: @ext_device_var = external global i32 extern __device__ int ext_device_var; -// ALL-DAG: @ext_device_var = internal global i32 +// ALL-DAG: @ext_device_var = external global i32 extern __constant__ int ext_constant_var; +// external device-side variables with definitions should generate +// definitions for the shadows. +// ALL-DAG: @ext_device_var_def = internal global i32 undef, +extern __device__ int ext_device_var_def; +__device__ int ext_device_var_def = 1; +// ALL-DAG: @ext_device_var_def = internal global i32 undef, +__constant__ int ext_constant_var_def = 2; + void use_pointers() { int *p; p = &device_var; @@ -114,8 +121,8 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); } // ALL: call{{.*}}[[PREFIX]]RegisterFunction(i8** %0, {{.*}}kernelfunc // ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}device_var{{.*}}i32 0, i32 4, i32 0, i32 0 // ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}constant_var{{.*}}i32 0, i32 4, i32 1, i32 0 -// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_device_var{{.*}}i32 1, i32 4, i32 0, i32 0 -// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_constant_var{{.*}}i32 1, i32 4, i32 1, i32 0 +// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_device_var_def{{.*}}i32 0, i32 4, i32 0, i32 0 +// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_constant_var_def{{.*}}i32 0, i32 4, i32 1, i32 0 // ALL: ret void // Test that we've built a constructor. |