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