[HIP][Clang][Sema][RFC] Add Sema support for C++ Parallel Algorithm Offload

[AlexVlx]

Summary

This patch adds the Sema changes needed by the standard algorithm offload feature being proposed here: https://discourse.llvm.org/t/rfc-adding-c-parallel-algorithm-offload-support-to-clang-llvm/72159/1. The verbose documentation is included in the head of the patch series. This change impacts the CUDA / HIP language specific checks, and only manifests if compiling in stdpar mode. In this case, we essentially do two things:

1. Allow device side callers to call host side callees - since the user visible HLL would be standard C++, with no annotations / restriction mechanisms, we cannot unambiguously establish that such a call is an error, so we conservatively allow all such calls, deferring actual cleanup to a subsequent pass over IR;
2. Allow host formed lambdas to capture by reference;
3. Allow device functions to use host global variables.

Please note that host and device here are used to match existing nomenclature, they would not be present in user code.

[AlexVlx]

Remove noise / unintended file. Add support for dealing with unsupported ASM.

[AlexVlx]

Fix typo.

[yaxunl]

clang/test/SemaStdPar/device-can-call-host.cpp

1 ↗

(On Diff #549117)

can we use %clang_cc1 and omit the stdpar path options?

[AlexVlx]

Update test.

[yaxunl]

LGTM. Thanks.

[AlexVlx]

Updating this to reflect the outcome of the RFC, which is that it will be added as a HIP extension exclusively.

clang/lib/Sema/SemaCUDA.cpp

if (CalleeTarget == CFT_HostDevice)
     return CFP_HostDevice;
 
   // (c) Best case scenarios
   if (CalleeTarget == CallerTarget ||
       (CallerTarget == CFT_Host && CalleeTarget == CFT_Global) ||
       (CallerTarget == CFT_Global && CalleeTarget == CFT_Device))
     return CFP_Native;
 
+  // HipStdPar mode is special, in that assessing whether a device side call to
+  // a host target is deferred to a subsequent pass, and cannot unambiguously be
+  // adjudicated in the AST, hence we optimistically allow them to pass here.
+  if (getLangOpts().HIPStdPar &&
+      (CallerTarget == CFT_Global || CallerTarget == CFT_Device ||
+       CallerTarget == CFT_HostDevice) &&
+      CalleeTarget == CFT_Host)
+    return CFP_HostDevice;
+
   // (d) HostDevice behavior depends on compilation mode.
   if (CallerTarget == CFT_HostDevice) {
     // It's OK to call a compilation-mode matching function from an HD one.
     if ((getLangOpts().CUDAIsDevice && CalleeTarget == CFT_Device) ||
         (!getLangOpts().CUDAIsDevice &&
          (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)))
       return CFP_SameSide;
 
void Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
   // to and called in a device function or kernel.
   bool CalleeIsDevice = Callee->hasAttr<CUDADeviceAttr>();
   bool CallerIsHost =
       !Caller->hasAttr<CUDAGlobalAttr>() && !Caller->hasAttr<CUDADeviceAttr>();
   bool ShouldCheck = CalleeIsDevice && CallerIsHost;
   if (!ShouldCheck || !Capture.isReferenceCapture())
     return;
   auto DiagKind = SemaDiagnosticBuilder::K_Deferred;
-  if (Capture.isVariableCapture()) {
+  if (Capture.isVariableCapture() && !getLangOpts().HIPStdPar) {
     SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
                           diag::err_capture_bad_target, Callee, *this)
         << Capture.getVariable();
   } else if (Capture.isThisCapture()) {
     // Capture of this pointer is allowed since this pointer may be pointing to
     // managed memory which is accessible on both device and host sides. It only
     // results in invalid memory access if this pointer points to memory not
     // accessible on device side.

clang/lib/Sema/SemaExpr.cpp

if (SemaRef.LangOpts.CUDA && Var->hasGlobalStorage()) {
     auto VarTarget = SemaRef.IdentifyCUDATarget(Var);
     auto UserTarget = SemaRef.IdentifyCUDATarget(FD);
     if (VarTarget == Sema::CVT_Host &&
         (UserTarget == Sema::CFT_Device || UserTarget == Sema::CFT_HostDevice ||
          UserTarget == Sema::CFT_Global)) {
       // Diagnose ODR-use of host global variables in device functions.
       // Reference of device global variables in host functions is allowed
       // through shadow variables therefore it is not diagnosed.
-      if (SemaRef.LangOpts.CUDAIsDevice) {
+      if (SemaRef.LangOpts.CUDAIsDevice && !SemaRef.LangOpts.HIPStdPar) {
         SemaRef.targetDiag(Loc, diag::err_ref_bad_target)
             << /*host*/ 2 << /*variable*/ 1 << Var << UserTarget;
         SemaRef.targetDiag(Var->getLocation(),
                            Var->getType().isConstQualified()
                                ? diag::note_cuda_const_var_unpromoted
                                : diag::note_cuda_host_var);
       }
     } else if (VarTarget == Sema::CVT_Device &&

clang/lib/Sema/SemaStmtAsm.cpp

for (unsigned i = 0; i != NumOutputs; i++) {
     StringLiteral *Literal = Constraints[i];
     assert(Literal->isOrdinary());
 
     StringRef OutputName;
     if (Names[i])
       OutputName = Names[i]->getName();
 
     TargetInfo::ConstraintInfo Info(Literal->getString(), OutputName);
-    if (!Context.getTargetInfo().validateOutputConstraint(Info)) {
+    if (!Context.getTargetInfo().validateOutputConstraint(Info) &&
+        !(LangOpts.HIPStdPar && LangOpts.CUDAIsDevice)) {
       targetDiag(Literal->getBeginLoc(),
                  diag::err_asm_invalid_output_constraint)
           << Info.getConstraintStr();
       return new (Context)
           GCCAsmStmt(Context, AsmLoc, IsSimple, IsVolatile, NumOutputs,
                      NumInputs, Names, Constraints, Exprs.data(), AsmString,
                      NumClobbers, Clobbers, NumLabels, RParenLoc);
     }

clang/test/SemaHipStdPar/device-can-call-host.cpp

+// RUN: %clang_cc1 -x hip %s --hipstdpar -triple amdgcn-amd-amdhsa --std=c++17 \
+// RUN:   -fcuda-is-device -emit-llvm -o /dev/null -verify
+
+// Note: These would happen implicitly, within the implementation of the
+//       accelerator specific algorithm library, and not from user code.
+
+// Calls from the accelerator side to implicitly host (i.e. unannotated)
+// functions are fine.
+
+// expected-no-diagnostics
+
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+
+extern "C" void host_fn() {}
+
+struct Dummy {};
+
+struct S {
+  S() {}
+  ~S() { host_fn(); }
+
+  int x;
+};
+
+struct T {
+  __device__ void hd() { host_fn(); }
+
+  __device__ void hd3();
+
+  void h() {}
+
+  void operator+();
+  void operator-(const T&) {}
+
+  operator Dummy() { return Dummy(); }
+};
+
+__device__ void T::hd3() { host_fn(); }
+
+template <typename T> __device__ void hd2() { host_fn(); }
+
+__global__ void kernel() { hd2<int>(); }
+
+__device__ void hd() { host_fn(); }
+
+template <typename T> __device__ void hd3() { host_fn(); }
+__device__ void device_fn() { hd3<int>(); }
+
+__device__ void local_var() {
+  S s;
+}
+
+__device__ void explicit_destructor(S *s) {
+  s->~S();
+}
+
+__device__ void hd_member_fn() {
+  T t;
+
+  t.hd();
+}
+
+__device__ void h_member_fn() {
+  T t;
+  t.h();
+}
+
+__device__ void unaryOp() {
+  T t;
+  (void) +t;
+}
+
+__device__ void binaryOp() {
+  T t;
+  (void) (t - t);
+}
+
+__device__ void implicitConversion() {
+  T t;
+  Dummy d = t;
+}
+
+template <typename T>
+struct TmplStruct {
+  template <typename U> __device__ void fn() {}
+};
+
+template <>
+template <>
+__device__ void TmplStruct<int>::fn<int>() { host_fn(); }
+
+__device__ void double_specialization() { TmplStruct<int>().fn<int>(); }