[PATCH] D80858: [CUDA][HIP] Support accessing static device variable in host code

[John McCall via Phabricator via cfe-commits]

rjmccall added inline comments.


================
Comment at: clang/lib/AST/ASTContext.cpp:10068
+        isa<VarDecl>(D) && cast<VarDecl>(D)->isFileVarDecl() &&
+        cast<VarDecl>(D)->getStorageClass() == SC_Static) {
+      return GVA_StrongExternal;
----------------
JonChesterfield wrote:
> yaxunl wrote:
> > rjmccall wrote:
> > > Are you sure this doesn't apply to e.g. local statics?  Can't you have kernel lambdas, or am I confusing HIP with another language?
> > function-scope static var in a device function is only visible to the device function. Host code cannot access it, therefore no need to externalize it.
> This doesn't sound right. An inline function can return a pointer to a function scope static variable, e.g. to implement a singleton in a header file.  I think host code can then access said variable.
Right, and IIRC you can declare __host__ __device__ functions as well, which ought to agree on the variable if they agree on globals.


================
Comment at: clang/lib/Driver/Action.cpp:170
+  if (!Id)
+    Id = llvm::sys::Process::GetRandomNumber();
+}
----------------
yaxunl wrote:
> rjmccall wrote:
> > I'm sure GetRandomNumber can return 0, so this logic is faulty.
> > 
> > This also seems like an unfortunate intrusion of HIP-specific semantics on the rest of the driver.  Maybe HIP can generate the shared id when it's setting up and cloning the job.
> Changed type of ID to string. Now empty ID means disabled. 0 is now allowed as a usual ID.
> 
> Moved initialization of ID to HIP action builder.
Thanks, that works.


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

https://reviews.llvm.org/D80858