[PATCH] D95560: [CUDA][HIP] Fix function scope static variable

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl marked 6 inline comments as done.
yaxunl added inline comments.


================
Comment at: clang/lib/CodeGen/CodeGenModule.cpp:101
+// does that.
+class CUDAStaticDeviceVarEmitter
+    : public StmtVisitor<CUDAStaticDeviceVarEmitter> {
----------------
tra wrote:
> Nit. "This class does that" could be dropped. I'd generally follow a `"<this thing> does <that> for <this reason>"` structure.
> E.g something along these lines:
> ```
> Helper class for emitting device-side static variables created in host-side functions. While we do not emit host-side functions on device, we still need to emit the static variables the host code will expect to see on the device.
> ```
done


================
Comment at: clang/lib/Sema/SemaCUDA.cpp:533-540
+      // isConstantInitializer cannot be called with dependent value, therefore
+      // we skip checking dependent value here. This is OK since
+      // checkAllowedCUDAInitializer is called again when the template is
+      // instantiated.
       AllowedInit =
-          ((VD->getType()->isDependentType() || Init->isValueDependent()) &&
-           VD->isConstexpr()) ||
+          (VD->getType()->isDependentType() || Init->isValueDependent()) ||
           Init->isConstantInitializer(Context,
----------------
tra wrote:
> This does not seem to be directly relevant for this patch. Perhaps move it into a separate patch?
separated to another patch


================
Comment at: clang/lib/Sema/SemaDecl.cpp:7247-7250
+  // CUDA/HIP: Function-scope static variables in device or global functions
+  // have implicit device or constant attribute. Function-scope static variables
+  // in host device functions have implicit device or constant attribute in
+  // device compilation only.
----------------
tra wrote:
> This is somewhat confusing. I guess the issue is that we're conflating all the functionality implied by the `__device__` attribute and the `accessible on device` which is a subset of it. For the static vars in D functions you only need for it to be accessible on device, IMO. For HD functions, you do need the full `__device__` functionality, with host shadow and runtime registration.
> 
> While adding implicit `__device__` works for statics in the device-only functions, it's a bit of an overkill. It also gives us a somewhat different AST between host/device compilations.
> 
> Perhaps we can handle statics in device-only functions w/o adding implicit `__device__`. Can we check the parent of the variable instead when we check whether we're allowed to reference the variable? 
Before we consider a function scope static variable without explicit device attribute, let's consider the difference between a static variable with explicit device attribute and a global device variable. They are both emitted in device compilation and have shadow variables in host compilation. The only difference is the linkage. A global device variable is supposed to be visible to other compilation units, whereas a static device variable is supposed to be visible to the same compilation unit only. A function scope static variable with device attribute has similar traits: It needs to be emitted in device compilation, and it needs a shadow variable in host compilation in case it needs to be accessed in host code. The only difference is that it is only visible inside the function.

Now let's consider a static var without device attribute in a device function. From sema and codegen point of view, it should have difference from a function scope static var with device attribute. Adding an implicit device attribute would simplify its handling.

Now let's consider a static var without device attribute in a host device function. The following code is valid for both nvcc and cuda-clang:

```
int __device__ __host__ func(int x) {
  static int a = 1;
  return a + x;
}
```
This requires the static variable is directly accessible in both device and host compilation. This requires that in device compilation, the static var behaves like a static var with explicit device attribute, whereas in host compilation, the static var behaves like a normal host static var. By adding implicit device attribute, we can clearly distinguish these situations and reuse the sema and codegen logic of device attribute.


================
Comment at: clang/test/CodeGenCUDA/func-scope-static-var.cu:54
+// NORDC: @_ZZ4fun2vE1b = dso_local addrspace(1) global i32 2
+// RDC: @_ZZ4fun2vE1b = internal addrspace(1) global i32 2
+// HOST: @_ZZ4fun2vE1b = internal global i32 2
----------------
tra wrote:
> What's the reason for externalizing the variables for no-rdc only?
> If we do not externalize them, then we'll potentially have a problem with the host code attempting to get variable's device-side address and fail at runtime, because it's not visible on device.
> 
> I think the right thing to do here is to always externalize them, but add unique suffix for RDC.
Yes this will be fixed by the patch for externalizing static var for -fgpu-rdc


================
Comment at: clang/test/CodeGenCUDA/func-scope-static-var.cu:87
+// In host device function, explicit static device variables are externalized
+// if used and registered. Default static variables are implicit device
+// variables in device compilation and host variables in host compilation,
----------------
tra wrote:
> Nit: `static variables w/o attributes are implicitly __device__`. Or `By default, static variables are implicitly __device__`.
> 
> It's also not clear what you mean by `which are independent`.  It may be better to add more details in a separate sentence.
revised


================
Comment at: clang/test/CodeGenCUDA/func-scope-static-var.cu:126-127
+
+// In kernels, static device variables are not externalized nor shadowed.
+// Static managed variable behaves like a normal static device variable.
+
----------------
tra wrote:
> We could use an explanation why we're not externalizing or shadowing them.
> 
added explanation


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

https://reviews.llvm.org/D95560