Merged main:e3f5a1bfc58b into origin/amd-gfx:6b6e30e6b6dc

Local branch origin/amd-gfx 6b6e30e Merged main:74f69c49fed8 into origin/amd-gfx:c4e1c89c4f41
Remote branch main e3f5a1b [LLDB][NFC]Also includes the error in log msg. (llvm#134922)

Author: SC llvm team

clang/include/clang/AST/OpenACCClause.h

@@ -38,6 +38,7 @@ class OpenACCClause {
   OpenACCClauseKind getClauseKind() const { return Kind; }
   SourceLocation getBeginLoc() const { return Location.getBegin(); }
   SourceLocation getEndLoc() const { return Location.getEnd(); }
+  SourceRange getSourceRange() const { return Location; }
 
   static bool classof(const OpenACCClause *) { return true; }
 

clang/lib/AST/ByteCode/Compiler.cpp

@@ -6823,7 +6823,7 @@ bool Compiler<Emitter>::emitDestruction(const Descriptor *Desc,
         return true;
     }
 
-    if (size_t N = Desc->getNumElems()) {
+    if (unsigned N = Desc->getNumElems()) {
       for (ssize_t I = N - 1; I >= 0; --I) {
         if (!this->emitConstUint64(I, Loc))
           return false;

clang/lib/CIR/CodeGen/CIRGenFunction.h

@@ -570,6 +570,16 @@ class CIRGenFunction : public CIRGenTypeCache {
   //===--------------------------------------------------------------------===//
   //                         OpenACC Emission
   //===--------------------------------------------------------------------===//
+private:
+  // Function to do the basic implementation of a 'compute' operation, including
+  // the clauses/etc. This might be generalizable in the future to work for
+  // other constructs, or at least be the base for construct emission.
+  template <typename Op, typename TermOp>
+  mlir::LogicalResult
+  emitOpenACCComputeOp(mlir::Location start, mlir::Location end,
+                       llvm::ArrayRef<const OpenACCClause *> clauses,
+                       const Stmt *structuredBlock);
+
 public:
   mlir::LogicalResult
   emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s);

clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp

@@ -12,16 +12,79 @@
 
 #include "CIRGenBuilder.h"
 #include "CIRGenFunction.h"
+#include "clang/AST/OpenACCClause.h"
 #include "clang/AST/StmtOpenACC.h"
 
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+
 using namespace clang;
 using namespace clang::CIRGen;
 using namespace cir;
+using namespace mlir::acc;
+
+namespace {
+class OpenACCClauseCIREmitter final
+    : public OpenACCClauseVisitor<OpenACCClauseCIREmitter> {
+  CIRGenModule &cgm;
+
+  void clauseNotImplemented(const OpenACCClause &c) {
+    cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind());
+  }
+
+public:
+  OpenACCClauseCIREmitter(CIRGenModule &cgm) : cgm(cgm) {}
+
+#define VISIT_CLAUSE(CN)                                                       \
+  void Visit##CN##Clause(const OpenACC##CN##Clause &clause) {                  \
+    clauseNotImplemented(clause);                                              \
+  }
+#include "clang/Basic/OpenACCClauses.def"
+};
+} // namespace
+
+template <typename Op, typename TermOp>
+mlir::LogicalResult CIRGenFunction::emitOpenACCComputeOp(
+    mlir::Location start, mlir::Location end,
+    llvm::ArrayRef<const OpenACCClause *> clauses,
+    const Stmt *structuredBlock) {
+  mlir::LogicalResult res = mlir::success();
+
+  OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule());
+  clauseEmitter.VisitClauseList(clauses);
+
+  llvm::SmallVector<mlir::Type> retTy;
+  llvm::SmallVector<mlir::Value> operands;
+  auto op = builder.create<Op>(start, retTy, operands);
+
+  mlir::Block &block = op.getRegion().emplaceBlock();
+  mlir::OpBuilder::InsertionGuard guardCase(builder);
+  builder.setInsertionPointToEnd(&block);
+
+  LexicalScope ls{*this, start, builder.getInsertionBlock()};
+  res = emitStmt(structuredBlock, /*useCurrentScope=*/true);
+
+  builder.create<TermOp>(end);
+  return res;
+}
 
 mlir::LogicalResult
 CIRGenFunction::emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s) {
-  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Compute Construct");
-  return mlir::failure();
+  mlir::Location start = getLoc(s.getSourceRange().getEnd());
+  mlir::Location end = getLoc(s.getSourceRange().getEnd());
+
+  switch (s.getDirectiveKind()) {
+  case OpenACCDirectiveKind::Parallel:
+    return emitOpenACCComputeOp<ParallelOp, mlir::acc::YieldOp>(
+        start, end, s.clauses(), s.getStructuredBlock());
+  case OpenACCDirectiveKind::Serial:
+    return emitOpenACCComputeOp<SerialOp, mlir::acc::YieldOp>(
+        start, end, s.clauses(), s.getStructuredBlock());
+  case OpenACCDirectiveKind::Kernels:
+    return emitOpenACCComputeOp<KernelsOp, mlir::acc::TerminatorOp>(
+        start, end, s.clauses(), s.getStructuredBlock());
+  default:
+    llvm_unreachable("invalid compute construct kind");
+  }
 }
 
 mlir::LogicalResult

clang/lib/CIR/CodeGen/CIRGenerator.cpp

@@ -12,6 +12,7 @@
 
 #include "CIRGenModule.h"
 
+#include "mlir/Dialect/OpenACC/OpenACC.h"
 #include "mlir/IR/MLIRContext.h"
 
 #include "clang/AST/DeclGroup.h"
@@ -36,6 +37,7 @@ void CIRGenerator::Initialize(ASTContext &astContext) {
 
   mlirContext = std::make_unique<mlir::MLIRContext>();
   mlirContext->loadDialect<cir::CIRDialect>();
+  mlirContext->getOrLoadDialect<mlir::acc::OpenACCDialect>();
   cgm = std::make_unique<clang::CIRGen::CIRGenModule>(
       *mlirContext.get(), astContext, codeGenOpts, diags);
 }

clang/lib/CodeGen/CGCUDANV.cpp

@@ -1281,7 +1281,8 @@ llvm::Function *CGNVCUDARuntime::finalizeModule() {
     return nullptr;
   }
   if (CGM.getLangOpts().OffloadViaLLVM ||
-      (CGM.getLangOpts().OffloadingNewDriver && RelocatableDeviceCode))
+      (CGM.getLangOpts().OffloadingNewDriver &&
+       (CGM.getLangOpts().HIP || RelocatableDeviceCode)))
     createOffloadingEntries();
   else
     return makeModuleCtorFunction();

clang/lib/Driver/Driver.cpp

@@ -4405,6 +4405,10 @@ void Driver::BuildActions(Compilation &C, DerivedArgList &Args,
                    options::OPT_no_offload_new_driver,
                    C.isOffloadingHostKind(Action::OFK_Cuda));
 
+  bool HIPNoRDC =
+      C.isOffloadingHostKind(Action::OFK_HIP) &&
+      !Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false);
+
   // Builder to be used to build offloading actions.
   std::unique_ptr<OffloadingActionBuilder> OffloadBuilder =
       !UseNewOffloadingDriver
@@ -4538,7 +4542,7 @@ void Driver::BuildActions(Compilation &C, DerivedArgList &Args,
     // Check if this Linker Job should emit a static library.
     if (ShouldEmitStaticLibrary(Args)) {
       LA = C.MakeAction<StaticLibJobAction>(LinkerInputs, types::TY_Image);
-    } else if (UseNewOffloadingDriver ||
+    } else if ((UseNewOffloadingDriver && !HIPNoRDC) ||
                Args.hasArg(options::OPT_offload_link)) {
       LA = C.MakeAction<LinkerWrapperJobAction>(LinkerInputs, types::TY_Image);
       LA->propagateHostOffloadInfo(C.getActiveOffloadKinds(),
@@ -4849,10 +4853,28 @@ Action *Driver::BuildOffloadingActions(Compilation &C,
                                        const InputTy &Input, StringRef CUID,
                                        Action *HostAction) const {
   // Don't build offloading actions if explicitly disabled or we do not have a
-  // valid source input and compile action to embed it in. If preprocessing only
-  // ignore embedding.
-  if (offloadHostOnly() || !types::isSrcFile(Input.first) ||
-      !(isa<CompileJobAction>(HostAction) ||
+  // valid source input.
+  if (offloadHostOnly() || !types::isSrcFile(Input.first))
+    return HostAction;
+
+  bool HIPNoRDC =
+      C.isOffloadingHostKind(Action::OFK_HIP) &&
+      !Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false);
+
+  // For HIP non-rdc non-device-only compilation, create a linker wrapper
+  // action for each host object to link, bundle and wrap device files in
+  // it.
+  if (isa<AssembleJobAction>(HostAction) && HIPNoRDC && !offloadDeviceOnly()) {
+    ActionList AL{HostAction};
+    HostAction = C.MakeAction<LinkerWrapperJobAction>(AL, types::TY_Object);
+    HostAction->propagateHostOffloadInfo(C.getActiveOffloadKinds(),
+                                         /*BoundArch=*/nullptr);
+    return HostAction;
+  }
+
+  // Don't build offloading actions if we do not have a compile action. If
+  // preprocessing only ignore embedding.
+  if (!(isa<CompileJobAction>(HostAction) ||
         getFinalPhase(Args) == phases::Preprocess))
     return HostAction;
 
@@ -4948,12 +4970,12 @@ Action *Driver::BuildOffloadingActions(Compilation &C,
       }
     }
 
-    // Compiling HIP in non-RDC mode requires linking each action individually.
+    // Compiling HIP in device-only non-RDC mode requires linking each action
+    // individually.
     for (Action *&A : DeviceActions) {
       if ((A->getType() != types::TY_Object &&
            A->getType() != types::TY_LTO_BC) ||
-          Kind != Action::OFK_HIP ||
-          Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false))
+          !HIPNoRDC || !offloadDeviceOnly())
         continue;
       ActionList LinkerInput = {A};
       A = C.MakeAction<LinkJobAction>(LinkerInput, types::TY_Image);
@@ -4977,12 +4999,12 @@ Action *Driver::BuildOffloadingActions(Compilation &C,
     }
   }
 
-  // HIP code in non-RDC mode will bundle the output if it invoked the linker.
+  // HIP code in device-only non-RDC mode will bundle the output if it invoked
+  // the linker.
   bool ShouldBundleHIP =
-      C.isOffloadingHostKind(Action::OFK_HIP) &&
+      HIPNoRDC && offloadDeviceOnly() &&
       Args.hasFlag(options::OPT_gpu_bundle_output,
                    options::OPT_no_gpu_bundle_output, true) &&
-      !Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false) &&
       !llvm::any_of(OffloadActions,
                     [](Action *A) { return A->getType() != types::TY_Image; });
 
@@ -5002,11 +5024,9 @@ Action *Driver::BuildOffloadingActions(Compilation &C,
         C.MakeAction<LinkJobAction>(OffloadActions, types::TY_CUDA_FATBIN);
     DDep.add(*FatbinAction, *C.getSingleOffloadToolChain<Action::OFK_Cuda>(),
              nullptr, Action::OFK_Cuda);
-  } else if (C.isOffloadingHostKind(Action::OFK_HIP) &&
-             !Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
-                           false)) {
-    // If we are not in RDC-mode we just emit the final HIP fatbinary for each
-    // translation unit, linking each input individually.
+  } else if (HIPNoRDC && offloadDeviceOnly()) {
+    // If we are in device-only non-RDC-mode we just emit the final HIP
+    // fatbinary for each translation unit, linking each input individually.
     Action *FatbinAction =
         C.MakeAction<LinkJobAction>(OffloadActions, types::TY_HIP_FATBIN);
     DDep.add(*FatbinAction, *C.getSingleOffloadToolChain<Action::OFK_HIP>(),
@@ -5159,8 +5179,11 @@ Action *Driver::ConstructPhaseAction(
         (((Input->getOffloadingToolChain() &&
            Input->getOffloadingToolChain()->getTriple().isAMDGPU()) ||
           TargetDeviceOffloadKind == Action::OFK_HIP) &&
-         (Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
-                       false) ||
+         ((Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
+                        false) ||
+           (Args.hasFlag(options::OPT_offload_new_driver,
+                         options::OPT_no_offload_new_driver, false) &&
+            !offloadDeviceOnly())) ||
           TargetDeviceOffloadKind == Action::OFK_OpenMP))) {
       types::ID Output =
           Args.hasArg(options::OPT_S) &&

clang/lib/Driver/ToolChains/Clang.cpp

@@ -7832,7 +7832,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     CmdArgs.push_back("-fcuda-include-gpubinary");
     CmdArgs.push_back(CudaDeviceInput->getFilename());
   } else if (!HostOffloadingInputs.empty()) {
-    if ((IsCuda || IsHIP) && !IsRDCMode) {
+    if (IsCuda && !IsRDCMode) {
       assert(HostOffloadingInputs.size() == 1 && "Only one input expected");
       CmdArgs.push_back("-fcuda-include-gpubinary");
       CmdArgs.push_back(HostOffloadingInputs.front().getFilename());
@@ -9344,11 +9344,22 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
   // Add the linker arguments to be forwarded by the wrapper.
   CmdArgs.push_back(Args.MakeArgString(Twine("--linker-path=") +
                                        LinkCommand->getExecutable()));
-  for (const char *LinkArg : LinkCommand->getArguments())
-    CmdArgs.push_back(LinkArg);
 
-  addOffloadCompressArgs(Args, CmdArgs);
+  // We use action type to differentiate two use cases of the linker wrapper.
+  // TY_Image for normal linker wrapper work.
+  // TY_Object for HIP fno-gpu-rdc embedding device binary in a relocatable
+  // object.
+  assert(JA.getType() == types::TY_Object || JA.getType() == types::TY_Image);
+  if (JA.getType() == types::TY_Object) {
+    CmdArgs.append({"-o", Output.getFilename()});
+    for (auto Input : Inputs)
+      CmdArgs.push_back(Input.getFilename());
+    CmdArgs.push_back("-r");
+  } else
+    for (const char *LinkArg : LinkCommand->getArguments())
+      CmdArgs.push_back(LinkArg);
 
+  addOffloadCompressArgs(Args, CmdArgs);
   const char *Exec =
       Args.MakeArgString(getToolChain().GetProgramPath("clang-linker-wrapper"));
 

clang/test/CIR/CodeGenOpenACC/kernels.c

@@ -0,0 +1,30 @@
+// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_kernels(void) {
+  // CHECK: cir.func @acc_kernels() {
+#pragma acc kernels
+  {}
+
+  // CHECK-NEXT: acc.kernels {
+  // CHECK-NEXT:acc.terminator
+  // CHECK-NEXT:}
+
+#pragma acc kernels
+  while(1){}
+  // CHECK-NEXT: acc.kernels {
+  // CHECK-NEXT: cir.scope {
+  // CHECK-NEXT: cir.while {
+  // CHECK-NEXT: %[[INT:.*]] = cir.const #cir.int<1>
+  // CHECK-NEXT: %[[CAST:.*]] = cir.cast(int_to_bool, %[[INT]] :
+  // CHECK-NEXT: cir.condition(%[[CAST]])
+  // CHECK-NEXT: } do {
+  // CHECK-NEXT: cir.yield
+  // cir.while do end:
+  // CHECK-NEXT: }
+  // cir.scope end:
+  // CHECK-NEXT: }
+  // CHECK-NEXT:acc.terminator
+  // CHECK-NEXT:}
+
+  // CHECK-NEXT: cir.return
+}

clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp

@@ -3,9 +3,9 @@
 
 void HelloWorld(int *A, int *B, int *C, int N) {
 
-// expected-error@+2{{ClangIR code gen Not Yet Implemented: OpenACC Compute Construct}}
+// expected-error@+2{{ClangIR code gen Not Yet Implemented: OpenACC Combined Construct}}
 // expected-error@+1{{ClangIR code gen Not Yet Implemented: statement}}
-#pragma acc parallel
+#pragma acc parallel loop
   for (unsigned I = 0; I < N; ++I)
     A[I] = B[I] + C[I];