[PATCH] D103108: [CUDA][HIP] Promote const variables to constant

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl 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
----------------
tra wrote:
> Nit: add an empty line to separate from the preceding function.
done


================
Comment at: clang/lib/Sema/SemaCUDA.cpp:718
+                                 *this, VD, CICK_DeviceOrConstant)))) {
     VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
   }
----------------
tra wrote:
> 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.
> 
added a TODO


================
Comment at: clang/test/CodeGenCUDA/device-use-host-var.cu:44
+const A const_array[] = {0, 0, 0, 6};
+const char const_str[] = "xyz";
 
----------------
tra wrote:
> 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.
> 
added test for that


================
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}}
----------------
tra wrote:
> 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.
> 
done


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

https://reviews.llvm.org/D103108