[CUDA][HIP] Fix device variables used by host

[yaxunl]

Summary

variables emitted on both host and device side with different addresses
when ODR-used by host function should not cause device side counter-part
to be force emitted.

This fixes the regression caused by https://reviews.llvm.org/D102237

[tra]

Tentative LGTM as we need it to fix the regression soon.

Summoning @rsmith for the 'big picture' opinion.
While the patch may fix this particular regression, I wonder if there's a better way to deal with this. We're growing a bit too many nuances that would be hard to explain and may cause more corner cases to appear.

clang/lib/CodeGen/CodeGenModule.cpp

2386

IIUIC, The idea here is that we do not want to emit constexpr int foo; on device, even if we happen to ODR-use it there.
And the way we detect this is by checking for implicit __constant__ we happen to add to constexpr variables.

I think this may be relying on the implementation details too much. It also makes compiler's behavior somewhat surprising -- we would potentially emit other variables that do not get any device attributes attribute, but would not emit the variables with implicit __constant__, which is a device attribute.

I'm not sure if we have any good options here. This may be an acceptable compromise, but I wonder if there's a better way to deal with this.

That said, this patch is OK to fix the regression we have now, but we may need to revisit this.

clang/test/CodeGenCUDA/host-used-device-var.cu

103–131

This definitely needs some comments. Otherwise this is nearly incomprehensible and it's impossible to tell what's going on.

[tra]

This patch does not appear to fix the second regression introduced by the D102237.

Trying to compile the following code triggers an assertion in CGExpr.cpp:

class a {
public:
  a(char *);
};
void b() {
  [](char *c) {
    static a d(c);
    d;
  };
}

With assertions disabled it eventually leads to a different error:
Module has a nontrivial global ctor, which NVPTX does not support.
https://godbolt.org/z/sYE1dKr1W

[yaxunl]

Fix the other regression

[yaxunl]

In D102801#2769936, @tra wrote:

This patch does not appear to fix the second regression introduced by the D102237.

Trying to compile the following code triggers an assertion in CGExpr.cpp:

class a {
public:
  a(char *);
};
void b() {
  [](char *c) {
    static a d(c);
    d;
  };
}

With assertions disabled it eventually leads to a different error:
Module has a nontrivial global ctor, which NVPTX does not support.
https://godbolt.org/z/sYE1dKr1W

The root cause is similar to the last regression. Basically when a variable is emitted on both sides but as different entities, we should not treat it as a device variable on host side. I have updated the patch to fix both regressions.

clang/lib/CodeGen/CodeGenModule.cpp

2386

we need to differentiate constexpr int a and __constant__ constexpr int a, since the former is emitted on both sides, and the later is only emitted on device side. It seems the only way to differentiate them is to check whether the constant attribute is explicit or not.

clang/test/CodeGenCUDA/host-used-device-var.cu

103–131

done

[yaxunl]

In D102801#2769619, @tra wrote:

Tentative LGTM as we need it to fix the regression soon.

Summoning @rsmith for the 'big picture' opinion.
While the patch may fix this particular regression, I wonder if there's a better way to deal with this. We're growing a bit too many nuances that would be hard to explain and may cause more corner cases to appear.

In the updated patch I have a simpler solution which is easier to explain to the users. Basically we classify variables by how they are emitted: device side only, host side only, both sides as different entities (e.g. default constexpr var), and both sides as unified entity (e.g. managed var). For variables emitted on both sides as separate entities, we have limited knowledge and we limit what we can do for them. I think users should understand the compiler's limitation in such cases. And they can easily workaround that by making the variable explicitly device variable.

[tra]

In D102801#2771664, @yaxunl wrote:

In the updated patch I have a simpler solution which is easier to explain to the users. Basically we classify variables by how they are emitted: device side only, host side only, both sides as different entities (e.g. default constexpr var), and both sides as unified entity (e.g. managed var). For variables emitted on both sides as separate entities, we have limited knowledge and we limit what we can do for them. I think users should understand the compiler's limitation in such cases. And they can easily workaround that by making the variable explicitly device variable.

This is really nice.

Let me test it internally and see if anything breaks.

clang/include/clang/Sema/Sema.h

12066 ↗

(On Diff #346796)

Wasn't there another kind, where the variable is emitted on the host with device-side shadow? I vaguely recall it had something to do with textures.

12067 ↗

(On Diff #346796)

I think we should mention the host-side shadows, too.

clang/lib/Sema/SemaCUDA.cpp

140–141

I'm still not a fan of relying on a implicit constant.
Can we change it to more direct is-a-constexpr && !has-explicit-device-side-attr ?
We may eventually consider relaxing this to can-be-const-evaluated and allow const vars with known values.

[yaxunl]

clang/include/clang/Sema/Sema.h

12066 ↗

(On Diff #346796)

That was the first implementation, which was similar to managed var but used pinned host memory as a common memory shared by device and host.

However, that implementation was later replaced by a different implementation which is similar to nvcc. In the new implementation textures and surfaces are like usual device variables. So far I do not see the necessity to differentiate them from usual device variables.

12067 ↗

(On Diff #346796)

will do

clang/lib/Sema/SemaCUDA.cpp

140–141

will do. agree we should relax this for const var in the future

[yaxunl]

revised by Artem's comments

[tra]

LGTM.

I've verified that Tensorflow still builds with this patch and that the patch does fix the regressions we've seen.
If you could land this patch soon, that would be appreciated.

clang/lib/CodeGen/CodeGenModule.cpp

if (!DeferredVTables.empty()) {
 
     // Emitting a vtable doesn't directly cause more vtables to
     // become deferred, although it can cause functions to be
     // emitted that then need those vtables.
     assert(DeferredVTables.empty());
   }
 
   // Emit CUDA/HIP static device variables referenced by host code only.
+  // Note we should not clear CUDADeviceVarODRUsedByHost since it is still
+  // needed for further handling.
   if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice)
-    for (const auto *V : getContext().CUDADeviceVarODRUsedByHost)
+    for (const auto *V : getContext().CUDADeviceVarODRUsedByHost) {
+      // ToDo: The user of ODR-used variables may not be emitted on host side.
+      // There needs a more accurate way to determine whether a device side
+      // variable ODR-used by host only should be emitted on device side.
+      //
+      // Currently we do it conservatively. However this does not work well
+      // with implicit constant variables since use of implicit constant
+      // variable in host function cannot be easily differentiated from
+      // use of the host variable with the same name. Therefore we do not
+      // force emit implicit constant variable.
+      auto HasImplicitConstantAttr = [](const VarDecl *Var) {
+        auto *A = Var->getAttr<CUDAConstantAttr>();
+        if (!A)
+          return false;
+        return A->isImplicit();
+      };
+      if (!HasImplicitConstantAttr(V))
         DeferredDeclsToEmit.push_back(V);
+    }
 
   // Stop if we're out of both deferred vtables and deferred declarations.
   if (DeferredDeclsToEmit.empty())
     return;
 
   // Grab the list of decls to emit. If EmitGlobalDefinition schedules more
   // work, it will not interfere with this.
   std::vector<GlobalDecl> CurDeclsToEmit;

clang/lib/Sema/SemaCUDA.cpp

Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D,
 
   return CFT_Host;
 }
 
 // * CUDA Call preference table
 //
 // F - from,
 // T - to
 // Ph - preference in host mode
 // Pd - preference in device mode
 // H  - handled in (x)
 // Preferences: N:native, SS:same side, HD:host-device, WS:wrong side, --:never.
 //
 // | F  | T  | Ph  | Pd  |  H  |
 // |----+----+-----+-----+-----+
 // | d  | d  | N   | N   | (c) |
 // | d  | g  | --  | --  | (a) |
 // | d  | h  | --  | --  | (e) |
void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
   }
 
   NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
   NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
 }
 
 void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
   if (getLangOpts().CUDAIsDevice && VD->isConstexpr() &&
-      (VD->isFileVarDecl() || VD->isStaticDataMember())) {
+      (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
+      !VD->hasAttr<CUDAConstantAttr>()) {
     VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
   }
 }
 
 Sema::SemaDiagnosticBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
                                                        unsigned DiagID) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   SemaDiagnosticBuilder::Kind DiagKind = [&] {

clang/lib/Sema/SemaDeclAttr.cpp

 }
 
 static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   const auto *VD = cast<VarDecl>(D);
   if (VD->hasLocalStorage()) {
     S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
     return;
   }
+  // constexpr variable may already get an implicit constant attr, which should
+  // be replaced by the explicit constant attr.
+  if (auto *A = D->getAttr<CUDAConstantAttr>()) {
+    if (!A->isImplicit())
+      return;
+    D->dropAttr<CUDAConstantAttr>();
+  }
   D->addAttr(::new (S.Context) CUDAConstantAttr(S.Context, AL));
 }
 
 static void handleSharedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   const auto *VD = cast<VarDecl>(D);
   // extern __shared__ is only allowed on arrays with no length (e.g.
   // "int x[]").
   if (!S.getLangOpts().GPURelocatableDeviceCode && VD->hasExternalStorage() &&

clang/test/AST/ast-dump-constant-var.cu

+// RUN: %clang_cc1 -std=c++14 -ast-dump -x hip %s | FileCheck -check-prefixes=CHECK,HOST %s
+// RUN: %clang_cc1 -std=c++14 -ast-dump -fcuda-is-device -x hip %s | FileCheck -check-prefixes=CHECK,DEV %s
+
+#include "Inputs/cuda.h"
+
+// CHECK-LABEL: VarDecl {{.*}} m1 'int'
+// CHECK-NEXT: CUDAConstantAttr {{.*}}cuda.h
+__constant__ int m1;
+
+// CHECK-LABEL: VarDecl {{.*}} m2 'int'
+// CHECK-NEXT: CUDAConstantAttr {{.*}}cuda.h
+// CHECK-NOT: CUDAConstantAttr
+__constant__ __constant__ int m2;
+
+// CHECK-LABEL: VarDecl {{.*}} m3 'const int'
+// HOST-NOT: CUDAConstantAttr
+// DEV-NOT: CUDAConstantAttr {{.*}}cuda.h
+// DEV: CUDAConstantAttr {{.*}}Implicit
+// DEV-NOT: CUDAConstantAttr {{.*}}cuda.h
+constexpr int m3 = 1;
+
+// CHECK-LABEL: VarDecl {{.*}} m3a 'const int'
+// CHECK-NOT: CUDAConstantAttr {{.*}}Implicit
+// CHECK: CUDAConstantAttr {{.*}}cuda.h
+// CHECK-NOT: CUDAConstantAttr {{.*}}Implicit
+constexpr __constant__ int m3a = 2;
+
+// CHECK-LABEL: VarDecl {{.*}} m3b 'const int'
+// CHECK-NOT: CUDAConstantAttr {{.*}}Implicit
+// CHECK: CUDAConstantAttr {{.*}}cuda.h
+// CHECK-NOT: CUDAConstantAttr {{.*}}Implicit
+__constant__ constexpr int m3b = 3;

clang/test/CodeGenCUDA/host-used-device-var.cu

 {
   return x + y;
 }
 
 // DEV-DAG: @_Z10p_add_funcIiE = linkonce_odr addrspace(1) externally_initialized global i32 (i32, i32)* @_Z8add_funcIiET_S0_S0_
 template <typename T>
 __device__ func_t<T> p_add_func = add_func<T>;
 
+// Check non-constant constexpr variables ODR-used by host code only is not emitted.
+// DEV-NEG-NOT: constexpr_var1a
+// DEV-NEG-NOT: constexpr_var1b
+constexpr int constexpr_var1a = 1;
+inline constexpr int constexpr_var1b = 1;
+
+// Check constant constexpr variables ODR-used by host code only.
+// Non-inline constexpr variable has internal linkage, therefore it is not accessible by host and not kept.
+// Inline constexpr variable has linkonce_ord linkage, therefore it can be accessed by host and kept.
+// DEV-NEG-NOT: constexpr_var2a
+// DEV-DAG: @constexpr_var2b = linkonce_odr addrspace(4) externally_initialized constant i32 2
+__constant__ constexpr int constexpr_var2a = 2;
+inline __constant__ constexpr int constexpr_var2b = 2;
+
 void use(func_t<int> p);
-void use(int *p);
+void use(const int *p);
 
 void fun1() {
   use(&u1);
   use(&u2);
   use(&u3);
   use(&ext_var);
   use(&inline_var);
   use(p_add_func<int>);
+  use(&constexpr_var1a);
+  use(&constexpr_var1b);
+  use(&constexpr_var2a);
+  use(&constexpr_var2b);
 }
 
 __global__ void kern1(int **x) {
   *x = &u4;
 }
 
+// Check implicit constant variable ODR-used by host code is not emitted.
+// DEV-NEG-NOT: _ZN16TestConstexprVar1oE
+namespace TestConstexprVar {
+char o;
+class ou {
+public:
+  ou(char) { __builtin_strlen(&o); }
+};
+template < typename ao > struct aw { static constexpr ao c; };
+class x {
+protected:
+  typedef ou (*y)(const x *);
+  constexpr x(y ag) : ah(ag) {}
+  template < bool * > struct ak;
+  template < typename > struct al {
+    static bool am;
+    static ak< &am > an;
+  };
+  template < typename ao > static x ap() { (void)aw< ao >::c; return x(nullptr); }
+  y ah;
+};
+template < typename ao > bool x::al< ao >::am(&ap< ao >);
+class ar : x {
+public:
+  constexpr ar() : x(as) {}
+  static ou as(const x *) { return 0; }
+  al< ar > av;
+};
+}
+
 // Check the exact list of variables to ensure @_ZL2u4 is not among them.
-// DEV: @llvm.compiler.used = {{[^@]*}} @_Z10p_add_funcIiE {{[^@]*}} @_ZL2u3 {{[^@]*}} @inline_var {{[^@]*}} @u1 {{[^@]*}} @u2 {{[^@]*}} @u5
+// DEV: @llvm.compiler.used = {{[^@]*}} @_Z10p_add_funcIiE {{[^@]*}} @_ZL2u3 {{[^@]*}} @constexpr_var2b {{[^@]*}} @inline_var {{[^@]*}} @u1 {{[^@]*}} @u2 {{[^@]*}} @u5
 
 // HOST-DAG: hipRegisterVar{{.*}}@u1
 // HOST-DAG: hipRegisterVar{{.*}}@u2
 // HOST-DAG: hipRegisterVar{{.*}}@_ZL2u3
+// HOST-DAG: hipRegisterVar{{.*}}@constexpr_var2b
 // HOST-DAG: hipRegisterVar{{.*}}@u5
 // HOST-DAG: hipRegisterVar{{.*}}@inline_var
 // HOST-DAG: hipRegisterVar{{.*}}@_Z10p_add_funcIiE
 // HOST-NEG-NOT: hipRegisterVar{{.*}}@ext_var
 // HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZL2u4
+// HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var1a
+// HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var1b
+// HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var2a