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