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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl marked 6 inline comments as done.
yaxunl added a comment.

In D85223#2551894 <https://reviews.llvm.org/D85223#2551894>, @JonChesterfield wrote:

> This works around the limitations of the binary format nvptx and amdgpu are using in the compiler. It's the wrong place in the stack to fix it - we could introduce another symbol table in the binary to capture the per-tu-between-arch scoping.
>
> However, if we later reach consensus on what to do in the elf instead, we can still do that. In particular, embedding an elf for one arch in a named section of an elf for a host arch is crude. This workaround seems acceptable in the meantime.

Yes we should revisit this if there is a better solution.



================
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
----------------
tra wrote:
> It should probably be a regex after `HASH:`, not the hash value itself.
will do


================
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
----------------
tra wrote:
> Same here.
will do


================
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"
----------------
tra wrote:
> ditto.
will do


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


================
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;
----------------
tra wrote:
> So, this verifies that we're allowed to use static local vars in device code. A comment would be useful.
will do


================
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;
----------------
tra wrote:
> 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.
will do


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

https://reviews.llvm.org/D85223