[flang-commits] [flang] [flang][cuda] duplicate host_device procedures early and outline the copies under their original name (PR #224400)
via flang-commits
flang-commits at lists.llvm.org
Thu Sep 17 14:03:33 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-flang-fir-hlfir
Author: Zhen Wang (wangzpgi)
<details>
<summary>Changes</summary>
An `attributes(host,device)` procedure is lowered as a single `func.func`, which `cuf-transform-device-func` clones into the GPU module at the end of the FIR pipeline. By then the optimizer has already placed the one body under the host policy, so with `-fstack-arrays` an automatic array stays a `fir.alloca` and overflows the device stack at run time:
```fortran
attributes(host,device) integer function sumauto(n)
integer :: auto(n)
...
end function
attributes(global) subroutine kern(a, n)
a(1) = sumauto(n) ! device call
end subroutine
```
Procedures reached from device code only through such a procedure, and device bodies inlined into a kernel, have the same problem.
This adds `cuf-duplicate-device-func`, meant to run before the FIR optimizer. Every `host_device` procedure, and every procedure without a device attribute that device code reaches through calls or procedure references, gets a device copy in the host module, and device code is redirected to it while host code keeps the original:
```mlir
func.func @<!-- -->sumauto(...) attributes {cuf.proc_attr = #cuf.cuda_proc<host_device>} {
%auto = fir.alloca !fir.array<?xi32>, %n // host copy, unchanged
}
func.func @<!-- -->sumauto.device(...) attributes {cuf.proc_attr = #cuf.cuda_proc<device>,
cuf.device_copy_of = @<!-- -->sumauto,
fir.allocation_policy = #fir.allocation_policy<stack_arrays = false, ...>} {
%auto = fir.alloca !fir.array<?xi32>, %n // placed to the heap by the optimizer
}
func.func @<!-- -->kern(...) attributes {cuf.proc_attr = #cuf.cuda_proc<global>} {
fir.call @<!-- -->sumauto.device(...) // was @<!-- -->sumauto
}
```
The optimizer then treats the two as ordinary functions with their own policies, and inlining a copy into a kernel carries the heap allocation along. Declarations are not copied, so runtime calls keep their real names and the descriptor I/O check in `cuf-transform-device-func` still fires.
`cuf-transform-device-func` recognizes the `cuf.device_copy_of` marker, outlines the copy under the original name, since device symbol names are the cross-unit ABI, restores the original names inside outlined code and `cuf.kernel` regions, leaves the `host_device` original in the host module and erases the copy:
```mlir
func.func @<!-- -->sumauto(...) { %auto = fir.alloca ... } // host module
gpu.module @<!-- -->cuda_device_mod {
gpu.func @<!-- -->sumauto(...) { %auto = fir.allocmem ... } // the copy, renamed back
gpu.func @<!-- -->kern(...) kernel { fir.call @<!-- -->sumauto(...) }
}
```
Without markers the pass behaves as before, so it is correct in either pipeline order.
The device code gathering is factored out of `cuf-transform-device-func` into `cuf::collectDeviceCode`, shared by both passes. `fir::getPresentedFunction` and `fir::getPresentedCallee` let diagnostics and remarks report a copy as the procedure it copies, and `fir::getPresentableFunctionName` uses them, so a remark inside `@<!-- -->sumauto.device` names `sumauto`. Lowering and the copies share `cuf::setDeviceAllocationPolicy`.
---
Patch is 41.94 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/224400.diff
15 Files Affected:
- (modified) flang/include/flang/Optimizer/Builder/CUFCommon.h (+44)
- (modified) flang/include/flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.h (+18)
- (modified) flang/include/flang/Optimizer/Support/Utils.h (+12-1)
- (modified) flang/include/flang/Optimizer/Transforms/Passes.td (+16)
- (modified) flang/lib/Lower/CallInterface.cpp (+2-5)
- (modified) flang/lib/Optimizer/Builder/CUFCommon.cpp (+121)
- (modified) flang/lib/Optimizer/Support/Utils.cpp (+24)
- (modified) flang/lib/Optimizer/Transforms/CMakeLists.txt (+1)
- (modified) flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp (+46-86)
- (modified) flang/lib/Optimizer/Transforms/CUDA/CUFDeviceGlobal.cpp (+1-3)
- (added) flang/lib/Optimizer/Transforms/CUDA/CUFDuplicateDeviceFunc.cpp (+107)
- (modified) flang/lib/Optimizer/Transforms/CUDA/CUFPredefinedVarToGPU.cpp (+4-10)
- (modified) flang/test/Fir/CUDA/cuda-device-func-transform-aio.mlir (+5)
- (added) flang/test/Fir/CUDA/cuda-device-func-transform-copies.mlir (+81)
- (added) flang/test/Fir/CUDA/cuda-duplicate-device-func.mlir (+111)
``````````diff
diff --git a/flang/include/flang/Optimizer/Builder/CUFCommon.h b/flang/include/flang/Optimizer/Builder/CUFCommon.h
index 736f90123969c..fc297a8919126 100644
--- a/flang/include/flang/Optimizer/Builder/CUFCommon.h
+++ b/flang/include/flang/Optimizer/Builder/CUFCommon.h
@@ -10,11 +10,18 @@
#define FORTRAN_OPTIMIZER_TRANSFORMS_CUFCOMMON_H_
#include "flang/Optimizer/Dialect/FIROps.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/IR/BuiltinOps.h"
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/SetVector.h"
static constexpr llvm::StringRef cudaDeviceModuleName = "cuda_device_mod";
static constexpr llvm::StringRef cudaSharedMemSuffix = "__shared_mem__";
+/// Appended to the name of the device copy of a procedure while it lives in
+/// the host module, where it cannot share the symbol of the original. The dot
+/// keeps it clear of any Fortran or C identifier.
+static constexpr llvm::StringRef cudaDeviceCopySuffix = ".device";
namespace fir {
class FirOpBuilder;
@@ -33,6 +40,43 @@ bool isCUDADeviceContext(mlir::Region &,
bool isRegisteredDeviceGlobal(fir::GlobalOp op);
bool isRegisteredDeviceAttr(std::optional<cuf::DataAttribute> attr);
+/// True for procedures that have a device side: attributes(device), (global),
+/// (grid_global) and (host,device). Unlike isCUDADeviceContext, host_device
+/// counts, since its body is compiled for the device as well.
+bool isDeviceProcedure(mlir::func::FuncOp funcOp);
+
+/// The device code of a module, as gathered by collectDeviceCode.
+struct DeviceCodeSet {
+ /// Procedures with a device proc attribute (see isDeviceProcedure).
+ llvm::SetVector<mlir::func::FuncOp> deviceFuncs;
+ /// Procedures without one that device code reaches, through calls or
+ /// procedure references, directly or through other such procedures. OpenACC
+ /// routines are excluded, the OpenACC pipeline moves those to the device
+ /// itself. Declarations are included: device code needs them too.
+ llvm::SetVector<mlir::func::FuncOp> calledFromDevice;
+ /// Device procedures referenced from a derived-type binding table, whose
+ /// host symbol must be kept for the table to verify.
+ llvm::SetVector<mlir::func::FuncOp> keepInModule;
+};
+
+/// Gather the device code of \p mod. A host_device procedure that already has
+/// a device copy is host code and is not walked for callees. With
+/// \p rejectDynamicDispatch, a type-bound call with dynamic dispatch in device
+/// code is reported as not yet implemented.
+DeviceCodeSet collectDeviceCode(mlir::ModuleOp mod, mlir::SymbolTable &symTab,
+ bool rejectDynamicDispatch = false);
+
+/// Point every fir.call and fir.address_of under \p root whose symbol is a key
+/// of \p map at the mapped symbol instead.
+void remapProcedureSymbols(
+ mlir::Operation *root,
+ const llvm::DenseMap<mlir::StringAttr, mlir::FlatSymbolRefAttr> &map);
+
+/// Record on \p func the allocation policy of device code: the policy in
+/// effect for it with stack arrays disabled, since the device stack is far
+/// smaller than the host one.
+void setDeviceAllocationPolicy(mlir::Operation *func);
+
void genPointerSync(const mlir::Value box, fir::FirOpBuilder &builder);
int computeElementByteSize(mlir::Location loc, mlir::Type type,
diff --git a/flang/include/flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.h b/flang/include/flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.h
index 88a47749b8808..6ce3c5ef71636 100644
--- a/flang/include/flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.h
+++ b/flang/include/flang/Optimizer/Dialect/CUF/Attributes/CUFAttr.h
@@ -15,6 +15,8 @@
#include "flang/Support/Fortran.h"
#include "mlir/IR/BuiltinAttributes.h"
+#include "mlir/IR/Operation.h"
+#include <optional>
namespace llvm {
class StringRef;
@@ -35,6 +37,22 @@ namespace cuf {
static constexpr llvm::StringRef dataAttrName = "data_attr";
static constexpr llvm::StringRef getDataAttrName() { return "cuf.data_attr"; }
static constexpr llvm::StringRef getProcAttrName() { return "cuf.proc_attr"; }
+/// On the device copy of a procedure, the symbol of the procedure it copies.
+static constexpr llvm::StringRef getDeviceCopyOfAttrName() {
+ return "cuf.device_copy_of";
+}
+/// Mark \p copy as the device copy of the procedure named \p original.
+inline void setDeviceCopyOf(mlir::Operation *copy, llvm::StringRef original) {
+ copy->setAttr(getDeviceCopyOfAttrName(),
+ mlir::FlatSymbolRefAttr::get(copy->getContext(), original));
+}
+/// The name of the procedure \p op is the device copy of, if it is one.
+inline std::optional<llvm::StringRef> getDeviceCopyOf(mlir::Operation *op) {
+ if (auto ref =
+ op->getAttrOfType<mlir::FlatSymbolRefAttr>(getDeviceCopyOfAttrName()))
+ return ref.getValue();
+ return std::nullopt;
+}
/// Attribute to carry CUDA launch_bounds values.
static constexpr llvm::StringRef getLaunchBoundsAttrName() {
diff --git a/flang/include/flang/Optimizer/Support/Utils.h b/flang/include/flang/Optimizer/Support/Utils.h
index d2f0be15d7dd9..ef7bff4e719db 100644
--- a/flang/include/flang/Optimizer/Support/Utils.h
+++ b/flang/include/flang/Optimizer/Support/Utils.h
@@ -255,8 +255,19 @@ mlir::Value integerCast(const fir::LLVMTypeConverter &converter,
/// otherwise it returns std::nullopt.
std::optional<bool> isNewAllocationResult(mlir::OpResult result);
+/// The procedure \p func stands for in diagnostics and remarks: itself, or the
+/// procedure it is a compiler-made copy of (the device copy of a CUDA Fortran
+/// host_device procedure) when that one can be found.
+mlir::FunctionOpInterface getPresentedFunction(mlir::FunctionOpInterface func);
+
+/// Same for the callee of \p call named by \p callee; null if it does not
+/// resolve to a function.
+mlir::FunctionOpInterface getPresentedCallee(mlir::Operation *call,
+ mlir::SymbolRefAttr callee);
+
/// Used to obtain user-facing function name that can be used in
-/// diagnostics and remarks without mangling or underscores.
+/// diagnostics and remarks without mangling or underscores. Compiler-made
+/// copies are reported under the name of the procedure they copy.
std::string getPresentableFunctionName(mlir::FunctionOpInterface func);
} // namespace fir
diff --git a/flang/include/flang/Optimizer/Transforms/Passes.td b/flang/include/flang/Optimizer/Transforms/Passes.td
index be2cbe31811e8..653f490715707 100644
--- a/flang/include/flang/Optimizer/Transforms/Passes.td
+++ b/flang/include/flang/Optimizer/Transforms/Passes.td
@@ -661,6 +661,22 @@ def CUFComputeSharedMemoryOffsetsAndSize
];
}
+def CUFDuplicateDeviceFunc
+ : Pass<"cuf-duplicate-device-func", "::mlir::ModuleOp"> {
+ let summary = "Give host_device procedures a device copy before optimization";
+ let description = [{
+ Creates the device copy of every attributes(host,device) procedure, and of
+ every procedure without a device attribute that device code calls, as a
+ func.func in the host module. The copy carries the device proc attribute,
+ the device allocation policy and a cuf.device_copy_of marker naming the
+ original, and device code is redirected to it while host code keeps the
+ original. Running before the FIR optimizer lets each copy be optimized for
+ its own target; cuf-transform-device-func later moves the copies into the
+ GPU module under their original name.
+ }];
+ let dependentDialects = ["cuf::CUFDialect", "fir::FIROpsDialect"];
+}
+
def CUFDeviceFuncTransform
: Pass<"cuf-transform-device-func", "::mlir::ModuleOp"> {
let summary = "Transform device function to GPU func";
diff --git a/flang/lib/Lower/CallInterface.cpp b/flang/lib/Lower/CallInterface.cpp
index 9885345bed3a5..2bfe94559cd88 100644
--- a/flang/lib/Lower/CallInterface.cpp
+++ b/flang/lib/Lower/CallInterface.cpp
@@ -16,11 +16,11 @@
#include "flang/Lower/PFTBuilder.h"
#include "flang/Lower/StatementContext.h"
#include "flang/Lower/Support/Utils.h"
+#include "flang/Optimizer/Builder/CUFCommon.h"
#include "flang/Optimizer/Builder/Character.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
#include "flang/Optimizer/Builder/Todo.h"
#include "flang/Optimizer/Dialect/FIROpsSupport.h"
-#include "flang/Optimizer/Support/AllocationPolicy.h"
#include "flang/Optimizer/Support/InternalNames.h"
#include "flang/Optimizer/Support/Utils.h"
#include "flang/Semantics/symbol.h"
@@ -706,10 +706,7 @@ setCUDAAttributes(mlir::func::FuncOp func,
cuf::ProcAttribute proc = procAttr.getValue();
if (proc != cuf::ProcAttribute::Host &&
proc != cuf::ProcAttribute::HostDevice) {
- fir::AllocationPolicy policy =
- fir::getAllocationPolicy(func.getOperation());
- policy.stackArrays = false;
- fir::setAllocationPolicy(func.getOperation(), policy);
+ cuf::setDeviceAllocationPolicy(func.getOperation());
}
}
diff --git a/flang/lib/Optimizer/Builder/CUFCommon.cpp b/flang/lib/Optimizer/Builder/CUFCommon.cpp
index 601f9fd6e2b31..2ba5e445634cf 100644
--- a/flang/lib/Optimizer/Builder/CUFCommon.cpp
+++ b/flang/lib/Optimizer/Builder/CUFCommon.cpp
@@ -8,11 +8,16 @@
#include "flang/Optimizer/Builder/CUFCommon.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
+#include "flang/Optimizer/Builder/Todo.h"
#include "flang/Optimizer/Dialect/CUF/CUFOps.h"
#include "flang/Optimizer/Dialect/Support/KindMapping.h"
#include "flang/Optimizer/HLFIR/HLFIROps.h"
+#include "flang/Optimizer/Support/AllocationPolicy.h"
+#include "flang/Optimizer/Support/InternalNames.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/SymbolTable.h"
+#include "llvm/ADT/StringSet.h"
/// Retrieve or create the CUDA Fortran GPU module in the give in \p mod.
mlir::gpu::GPUModuleOp cuf::getOrCreateGPUModule(mlir::ModuleOp mod,
@@ -156,3 +161,119 @@ mlir::Value cuf::computeElementCount(mlir::PatternRewriter &rewriter,
}
return mlir::Value();
}
+
+// Factored out of CUFDeviceFuncTransform.cpp (was isDeviceFunc) so that every
+// CUF pass tests "has a device side" the same way.
+bool cuf::isDeviceProcedure(mlir::func::FuncOp funcOp) {
+ auto procAttr =
+ funcOp->getAttrOfType<cuf::ProcAttributeAttr>(cuf::getProcAttrName());
+ if (!procAttr)
+ return false;
+ switch (procAttr.getValue()) {
+ case cuf::ProcAttribute::Device:
+ case cuf::ProcAttribute::Global:
+ case cuf::ProcAttribute::GridGlobal:
+ case cuf::ProcAttribute::HostDevice:
+ return true;
+ case cuf::ProcAttribute::Host:
+ return false;
+ }
+ return false;
+}
+
+// Factored out of CUFDeviceFuncTransform.cpp so that cuf-duplicate-device-func
+// and cuf-transform-device-func agree on what device code is. Compared to the
+// original: transitive, follows any symbol reference, skips copied originals.
+cuf::DeviceCodeSet cuf::collectDeviceCode(mlir::ModuleOp mod,
+ mlir::SymbolTable &symTab,
+ bool rejectDynamicDispatch) {
+ cuf::DeviceCodeSet code;
+ // Originals that already have a device copy are host code from now on.
+ llvm::StringSet<> hasDeviceCopy;
+ mod.walk([&](mlir::func::FuncOp funcOp) {
+ if (cuf::isDeviceProcedure(funcOp)) {
+ code.deviceFuncs.insert(funcOp);
+ if (std::optional<llvm::StringRef> original =
+ cuf::getDeviceCopyOf(funcOp))
+ hasDeviceCopy.insert(*original);
+ }
+ });
+
+ // Everything device code reaches without a device attribute of its own. A
+ // worklist makes this transitive: a plain procedure two calls below a kernel
+ // is device code too, and must be optimized as such.
+ llvm::SmallVector<mlir::Operation *> worklist;
+ auto found = [&](mlir::StringAttr name) {
+ auto callee = symTab.lookup<mlir::func::FuncOp>(name);
+ if (!callee || mlir::acc::isAccRoutine(callee) ||
+ code.deviceFuncs.count(callee))
+ return;
+ if (code.calledFromDevice.insert(callee) && !callee.isDeclaration())
+ worklist.push_back(callee);
+ };
+ auto scan = [&](mlir::Operation *root) {
+ // Every symbol referenced in the body, whatever op carries it: fir.call,
+ // fir.address_of, and anything added later. The root's own attributes
+ // (such as cuf.device_copy_of) are not references to device code.
+ for (mlir::Region ®ion : root->getRegions())
+ if (std::optional<mlir::SymbolTable::UseRange> uses =
+ mlir::SymbolTable::getSymbolUses(®ion))
+ for (const mlir::SymbolTable::SymbolUse &use : *uses)
+ found(use.getSymbolRef().getLeafReference());
+ // Only the outliner asks for this; before the optimizer fir.dispatch is
+ // still present in ordinary host code and must not be rejected.
+ if (rejectDynamicDispatch)
+ root->walk([](fir::DispatchOp op) {
+ TODO(op.getLoc(),
+ "type-bound procedure call with dynamic dispatch in device code");
+ });
+ };
+ for (mlir::func::FuncOp funcOp : code.deviceFuncs)
+ if (!hasDeviceCopy.contains(funcOp.getSymName()))
+ worklist.push_back(funcOp);
+ mod.walk([&](cuf::KernelOp kernelOp) { worklist.push_back(kernelOp); });
+ while (!worklist.empty())
+ scan(worklist.pop_back_val());
+
+ // A device procedure referenced from a binding table keeps a host symbol,
+ // or the table fails to verify once lowered to the LLVM dialect.
+ for (fir::GlobalOp globalOp : mod.getOps<fir::GlobalOp>()) {
+ if (!globalOp.getName().contains(fir::kBindingTableSeparator))
+ continue;
+ globalOp.walk([&](fir::AddrOfOp addrOfOp) {
+ auto funcOp = symTab.lookup<mlir::func::FuncOp>(
+ addrOfOp.getSymbol().getLeafReference());
+ if (funcOp && code.deviceFuncs.count(funcOp))
+ code.keepInModule.insert(funcOp);
+ });
+ }
+ return code;
+}
+
+// Shared by both passes: the duplicator points device code at the copies and
+// the outliner points it back, and the two must rewrite the same ops.
+void cuf::remapProcedureSymbols(
+ mlir::Operation *root,
+ const llvm::DenseMap<mlir::StringAttr, mlir::FlatSymbolRefAttr> &map) {
+ if (map.empty())
+ return;
+ root->walk([&](mlir::Operation *op) {
+ if (auto call = mlir::dyn_cast<fir::CallOp>(op)) {
+ if (mlir::SymbolRefAttr callee = call.getCalleeAttr())
+ if (auto it = map.find(callee.getLeafReference()); it != map.end())
+ call.setCalleeAttr(it->second);
+ } else if (auto addrOf = mlir::dyn_cast<fir::AddrOfOp>(op)) {
+ if (auto it = map.find(addrOf.getSymbol().getLeafReference());
+ it != map.end())
+ addrOf.setSymbolAttr(it->second);
+ }
+ });
+}
+
+// Shared by lowering (device and global procedures) and the duplicator (device
+// copies), so both record the same device policy.
+void cuf::setDeviceAllocationPolicy(mlir::Operation *func) {
+ fir::AllocationPolicy policy = fir::getAllocationPolicy(func);
+ policy.stackArrays = false;
+ fir::setAllocationPolicy(func, policy);
+}
diff --git a/flang/lib/Optimizer/Support/Utils.cpp b/flang/lib/Optimizer/Support/Utils.cpp
index a9590d3073ff2..93dc0c569a97d 100644
--- a/flang/lib/Optimizer/Support/Utils.cpp
+++ b/flang/lib/Optimizer/Support/Utils.cpp
@@ -14,6 +14,7 @@
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
#include "flang/Optimizer/Support/InternalNames.h"
+#include "mlir/IR/SymbolTable.h"
fir::TypeInfoOp fir::lookupTypeInfoOp(fir::RecordType recordType,
mlir::ModuleOp module,
@@ -149,7 +150,30 @@ std::optional<bool> fir::isNewAllocationResult(mlir::OpResult result) {
return false;
}
+mlir::FunctionOpInterface
+fir::getPresentedFunction(mlir::FunctionOpInterface func) {
+ if (std::optional<llvm::StringRef> original =
+ cuf::getDeviceCopyOf(func.getOperation()))
+ if (auto originalFunc = mlir::SymbolTable::lookupNearestSymbolFrom<
+ mlir::FunctionOpInterface>(
+ func.getOperation(),
+ mlir::StringAttr::get(func.getContext(), *original)))
+ return originalFunc;
+ return func;
+}
+
+mlir::FunctionOpInterface fir::getPresentedCallee(mlir::Operation *call,
+ mlir::SymbolRefAttr callee) {
+ if (!callee)
+ return nullptr;
+ auto func =
+ mlir::SymbolTable::lookupNearestSymbolFrom<mlir::FunctionOpInterface>(
+ call, callee);
+ return func ? getPresentedFunction(func) : nullptr;
+}
+
std::string fir::getPresentableFunctionName(mlir::FunctionOpInterface func) {
+ func = getPresentedFunction(func);
if (func.getName() == fir::NameUniquer::doProgramEntry()) {
// Main program entry is all uppercase - to avoid name conflicts. But
// from a reporting perspective, keep it lowercase for consistency with
diff --git a/flang/lib/Optimizer/Transforms/CMakeLists.txt b/flang/lib/Optimizer/Transforms/CMakeLists.txt
index 9b25e590764ab..7fb77ff1ed224 100644
--- a/flang/lib/Optimizer/Transforms/CMakeLists.txt
+++ b/flang/lib/Optimizer/Transforms/CMakeLists.txt
@@ -15,6 +15,7 @@ add_flang_library(FIRTransforms
CUDA/CUFComputeSharedMemoryOffsetsAndSize.cpp
CUDA/CUFDeviceFuncTransform.cpp
CUDA/CUFDeviceGlobal.cpp
+ CUDA/CUFDuplicateDeviceFunc.cpp
CUDA/CUFFunctionRewrite.cpp
CUDA/CUFGPUToLLVMConversion.cpp
CUDA/CUFLaunchAttachAttr.cpp
diff --git a/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp b/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp
index a0303a181775e..f5d605fbd4094 100644
--- a/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp
+++ b/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp
@@ -7,20 +7,18 @@
//===----------------------------------------------------------------------===//
#include "flang/Optimizer/Builder/CUFCommon.h"
-#include "flang/Optimizer/Builder/Todo.h"
#include "flang/Optimizer/Dialect/CUF/CUFOps.h"
#include "flang/Optimizer/Dialect/FIRAttr.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIRType.h"
-#include "flang/Optimizer/Support/InternalNames.h"
#include "flang/Optimizer/Transforms/Passes.h"
#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
-#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/IR/IRMapping.h"
#include "mlir/Pass/Pass.h"
+#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/SetVector.h"
#include "llvm/ADT/StringSet.h"
@@ -85,7 +83,8 @@ class CUFDeviceFuncTransform
}
static gpu::GPUFuncOp createGPUFuncOp(mlir::func::FuncOp funcOp,
- bool isGlobal, int computeCap) {
+ llvm::StringRef name, bool isGlobal,
+ int computeCap) {
mlir::OpBuilder builder(funcOp.getContext());
mlir::Region &funcOpBody = funcOp.getBody();
@@ -107,9 +106,8 @@ class CUFDeviceFuncTransform
mlir::FunctionType type = mlir::FunctionType::get(
funcOp.getContext(), funcOperandTypes, funcResultTypes);
- auto deviceFuncOp =
- gpu::GPUFuncOp::create(builder, loc, funcOp.getName(), type,
- mlir::TypeRange{}, mlir::TypeRange{});
+ auto deviceFuncOp = gpu::GPUFuncOp::create(
+ builder, loc, name, type, mlir::TypeRange{}, mlir::TypeRange{});
if (mlir::ArrayAttr argAttrs = funcOp.getAllArgAttrs())
deviceFuncOp.setAllArgAttrs(argAttrs);
setIntentInKernelArgAttrs(funcOp, deviceFuncOp);
@@ -183,18 +181,6 @@ class CUFDeviceFuncTransform
symTab.insert(emptyStub);
}
- static bool isDeviceFunc(mlir::func::FuncOp funcOp) {
- if (auto cudaProcAttr =
- funcOp.getOperation()->getAttrOfType<cuf::ProcAttributeAttr>(
- cuf::getProcAttrName()))
- if (cudaProcAttr.getValue() == cuf::ProcAttribute::Device ||
- cudaProcAttr.getValue() == cuf::ProcAttribute::Global ||
- cudaProcAttr.getValue() == cuf::ProcAttribute::GridGlobal ||
- cudaProcAttr.getValue() == cuf::ProcAttribute::HostDevice)
- return true;
- return false;
- }
-
void runOnOperation() override {
// Working on Module operation because inserting/removing function from the
// module is not thread-safe.
@@ -207,64 +...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/224400
More information about the flang-commits
mailing list