[PATCH] D85223: [CUDA][HIP] Support accessing static device variable in host code for -fgpu-rdc

[Artem Belevich via Phabricator via cfe-commits]

tra accepted this revision.
tra added a comment.
This revision is now accepted and ready to land.

LGTM with new test nits.

@JonChesterfield -- are you OK with the patch?



================
Comment at: clang/test/CodeGenCUDA/device-var-linkage.cu:40
 // NORDC-DAG: @_ZL3sv1 = dso_local addrspace(1) externally_initialized global i32 0
-// RDC-DAG: @_ZL3sv1 = internal addrspace(1) global i32 0
+// RDC-DAG: @_ZL3sv1.static.[[HASH:b04fd23c98500190]] = dso_local addrspace(1) externally_initialized global i32 0
 // HOST-DAG: @_ZL3sv1 = internal global i32 undef
----------------
It should probably be a regex after `HASH:`, not the hash value itself.


================
Comment at: clang/test/CodeGenCUDA/managed-var.cu:42
+// NORDC-D-DAG: @_ZL2sx = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// RDC-D-DAG: @_ZL2sx.static.[[HASH:b04fd23c98500190]] = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
 // HOST-DAG: @_ZL2sx.init = internal global i32 1
----------------
Same here.


================
Comment at: clang/test/CodeGenCUDA/static-device-var-rdc.cu:34
+// Test externalized static device variables
+// EXT-DEV-DAG: @_ZL1x.static.[[HASH:b04fd23c98500190]] = dso_local addrspace(1) externally_initialized global i32 0
+// EXT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH:b04fd23c98500190]]\00"
----------------
ditto.


================
Comment at: clang/test/SemaCUDA/static-device-var.cu:10
+
+// expected-no-diagnostics
+
----------------
A comment explaining what we're testing would be helpful. `no-diagnostics` gives no clues about what is it we're looking for here.




================
Comment at: clang/test/SemaCUDA/static-device-var.cu:14-22
+__device__ void f1() {
+  const static int b = 123;
+  static int a;
+}
+
+__global__ void k1() {
+  const static int b = 123;
----------------
So, this verifies that we're allowed to use static local vars in device code. A comment would be useful.


================
Comment at: clang/test/SemaCUDA/static-device-var.cu:23-37
+
+static __device__ int x;
+static __constant__ int y;
+
+__global__ void kernel(int *a) {
+  a[0] = x;
+  a[1] = y;
----------------
And this verifies that global static vars can be referenced from both host and device. 
I'd also add a negative test with `static int host_only;` and would verify that we still don't allow accessing it from the device.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D85223/new/

https://reviews.llvm.org/D85223