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