[flang-commits] [flang] [llvm] [flang][cuda] share runtime type info between host and device under managed memory (PR #229213)

[via flang-commits]

github-actions[bot] wrote:

<!--LLVM CODE FORMAT COMMENT: {clang-format}-->


:warning: C/C++ code formatter, clang-format found issues in your code. :warning:

<details>
<summary>
You can test this locally with the following command:
</summary>

``````````bash
git-clang-format --diff origin/main HEAD --extensions cpp,h -- flang/lib/Optimizer/Transforms/CUDA/CUFSharedTypeInfo.cpp flang-rt/lib/cuda/registration.cpp flang/include/flang/Optimizer/Builder/CUFCommon.h flang/include/flang/Runtime/CUDA/registration.h flang/lib/Optimizer/Builder/CUFCommon.cpp flang/lib/Optimizer/CodeGen/CodeGen.cpp flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp --diff_from_common_commit
``````````

:warning:
The reproduction instructions above might return results for more than one PR
in a stack if you are using a stacked PR workflow. You can limit the results by
changing `origin/main` to the base branch/commit you want to compare against.
:warning:

</details>

<details>
<summary>
View the diff from clang-format here.
</summary>

``````````diff
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index 4c9b9e1f0..b8f7c3e28 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -257,8 +257,7 @@ loadSharedTypeDescriptor(mlir::gpu::GPUModuleOp gpuMod,
   mlir::DataLayout dataLayout(gpuMod);
   // Managed globals live in the global address space on the device.
   mlir::Value ptrAddr = replaceWithAddrOfOrASCast(
-      rewriter, loc,
-      static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Global),
+      rewriter, loc, static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Global),
       fir::factory::getProgramAddressSpace(&dataLayout), ptrSym.getValue(),
       llvmPtrTy);
   return mlir::LLVM::LoadOp::create(rewriter, loc, llvmPtrTy, ptrAddr);
diff --git a/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp b/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
index 1f57094d6..ba4f43067 100644
--- a/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
+++ b/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
@@ -444,8 +444,7 @@ struct CUFAddConstructor
       mlir::Value end = getSectionBound("__stop_");
       auto rangeFuncTy = mlir::LLVM::LLVMFunctionType::get(
           voidTy, {llvmPtrTy, llvmPtrTy}, /*isVarArg=*/false);
-      llvm::StringRef rangeFuncName =
-          RTNAME_STRING(CUFRegisterHostMemoryRange);
+      llvm::StringRef rangeFuncName = RTNAME_STRING(CUFRegisterHostMemoryRange);
       if (!mod.lookupSymbol(rangeFuncName)) {
         mlir::OpBuilder::InsertionGuard guard(builder);
         builder.setInsertionPointToEnd(mod.getBody());
@@ -595,7 +594,8 @@ struct CUFAddConstructor
                 cudaHostTypeDescAttrName);
             if (!typeDescRef)
               continue;
-            auto typeDesc = symTab.lookup<fir::GlobalOp>(typeDescRef.getValue());
+            auto typeDesc =
+                symTab.lookup<fir::GlobalOp>(typeDescRef.getValue());
             auto ptrGlobal = mod.lookupSymbol<fir::GlobalOp>(
                 (globalOp.getSymName() + managedPtrSuffix).str());
             if (!typeDesc || !ptrGlobal)
diff --git a/flang/lib/Optimizer/Transforms/CUDA/CUFSharedTypeInfo.cpp b/flang/lib/Optimizer/Transforms/CUDA/CUFSharedTypeInfo.cpp
index 124892443..0cacacecf 100644
--- a/flang/lib/Optimizer/Transforms/CUDA/CUFSharedTypeInfo.cpp
+++ b/flang/lib/Optimizer/Transforms/CUDA/CUFSharedTypeInfo.cpp
@@ -89,7 +89,8 @@ public:
     std::string unitTag = getUnitTag(mod);
     mlir::OpBuilder builder(ctx);
     auto ptrTy = fir::LLVMPointerType::get(ctx, mlir::IntegerType::get(ctx, 8));
-    auto managed = cuf::DataAttributeAttr::get(ctx, cuf::DataAttribute::Managed);
+    auto managed =
+        cuf::DataAttributeAttr::get(ctx, cuf::DataAttribute::Managed);
     llvm::SmallVector<mlir::NamedAttribute> sharedTypeDescs;
     for (fir::GlobalOp typeDesc : typeDescs) {
       if (!gpuSymTab.lookup(typeDesc.getSymName()))
@@ -108,8 +109,9 @@ public:
       fir::HasValueOp::create(builder, typeDesc.getLoc(), zero);
 
       gpuSymTab.insert(ptrGlobal->clone());
-      ptrGlobal->setAttr(cudaHostTypeDescAttrName,
-                         mlir::FlatSymbolRefAttr::get(typeDesc.getSymNameAttr()));
+      ptrGlobal->setAttr(
+          cudaHostTypeDescAttrName,
+          mlir::FlatSymbolRefAttr::get(typeDesc.getSymNameAttr()));
       sharedTypeDescs.emplace_back(
           typeDesc.getSymNameAttr(),
           mlir::FlatSymbolRefAttr::get(ptrGlobal.getSymNameAttr()));

``````````

</details>


https://github.com/llvm/llvm-project/pull/229213