tra added inline comments.
================
Comment at: clang/lib/Sema/SemaCUDA.cpp:568
+}
+// Check whether a variable has an allowed initializer for a CUDA device side
+// variable with global storage. \p VD may be a host variable to be checked for
----------------
Nit: add an empty line to separate from the preceding function.
================
Comment at: clang/lib/Sema/SemaCUDA.cpp:718
+ *this, VD, CICK_DeviceOrConstant)))) {
VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
}
----------------
One of the issues with automatic promotion I've ran into in the past is that
`__constant__` is a very limited resource.
It would be nearly impossible for the end user to predict how much
`__constant__` space is available for their own use and if we use too much, the
failure would only manifest at runtime when we fail to load the GPU code. That
will be hard to troubleshoot.
We may need to add some sort of safeguard and report an error when we see too
much __constant__ data for a given GPU.
That would have to be done somewhere at the end of the compilation pipeline.
It's not likely to affect any/many users yet so it can be done later. It's
probably worth a comment and a TODO here in case someone does run into this.
================
Comment at: clang/test/CodeGenCUDA/device-use-host-var.cu:44
+const A const_array[] = {0, 0, 0, 6};
+const char const_str[] = "xyz";
----------------
It would be great to have a test which verifies that we only emit const
variables that we need.
It would be bad if we'd end up with *all* host-side consts in the GPU binary.
================
Comment at: clang/test/SemaCUDA/device-use-host-var.cu:82
*out = global_constexpr_var;
+ *out = b1.x; // dev-error {{reference to __host__ variable 'b1' in
__device__ function}}
+ *out = b2.x; // dev-error {{reference to __host__ variable 'b2' in
__device__ function}}
----------------
This diag is a bit misleading now. While it's technically true that we're not
allowed to refer to a host variables from device in general, the real reason
it's not allowed in *this* case is that the variable has a non-empty ctor.
I wonder if we could add a note pointing that out. Otherwise it would be rather
confusing why I can access other host vars few lines above, but not here.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D103108/new/
https://reviews.llvm.org/D103108
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
