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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl marked 9 inline comments as done.
yaxunl added inline comments.
Herald added a subscriber: dang.


================
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.
As long as we are not accessing the static variable by symbol we do not need externalize it.

If a device function returns a pointer to its static variable and somehow passes that pointer to host code, the host code can use it directly by hipMemCpy.


================
Comment at: clang/lib/AST/ASTContext.cpp:10068
+        isa<VarDecl>(D) && cast<VarDecl>(D)->isFileVarDecl() &&
+        cast<VarDecl>(D)->getStorageClass() == SC_Static) {
+      return GVA_StrongExternal;
----------------
rjmccall wrote:
> yaxunl wrote:
> > 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.
> > As long as we are not accessing the static variable by symbol we do not need externalize it.
> > 
> > If a device function returns a pointer to its static variable and somehow passes that pointer to host code, the host code can use it directly by hipMemCpy.
> Right, and IIRC you can declare __host__ __device__ functions as well, which ought to agree on the variable if they agree on globals.
If we have a static variable in a device function, it is only visible in the function and not visible by any host code. We only need externalize it if it needs to be accessed `by symbol` in the host code, however, that is impossible, therefore we do not need externalize it.

For static variables in a host device function, the static variables should be different instances on host side and device side. The rationale is that a static variable is per function, whereas a host device function is actually two functions: a host instance and a device instance, which could be totally different by using conditional macros. If it is requested that the static variable in a host device function is one instance, it requires special handling in runtime so that the same variable can be accessed on both device side and host side by common load/store instructions, but that is not the case. Therefore the device side instance of a static variable in a host device function is still only visible to device codes, not visible to host codes. Since it cannot be accessed `by symbol` by host code, it does not needs to be externalized.
 


================
Comment at: clang/lib/CodeGen/CodeGenModule.cpp:6069
+    llvm::raw_ostream &OS) const {
+  OS << ".static." << getLangOpts().CUID;
+}
----------------
tra wrote:
> I suspect that will have interesting issues if CUID is an arbitrary user-supplied string. We may want to impose some sort of sanity check or filtering on the cuid value. Considering that it's a CC1 flag, it's not a critical problem, but some safeguards would be useful there, too. Should we limit allowed character set?
will only allow alphanumeric and underscore in CUID for simplicity.


================
Comment at: clang/test/Driver/hip-cuid.hip:35
+// RUN:   --offload-arch=gfx906 \
+// RUN:   -c -nogpulib -cuid=abcd \
+// RUN:   %S/Inputs/hip_multiple_inputs/a.cu \
----------------
tra wrote:
> Nit: `abcd` could potentially match the value generated by hash. I'd change it to contain characters other than hex.
done


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

https://reviews.llvm.org/D80858