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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
