Skip to content

Commit 9487980

Browse files
committed
Looks at LLVMFuncOp
1 parent b2d4f9b commit 9487980

File tree

2 files changed

+40
-5
lines changed

2 files changed

+40
-5
lines changed

flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp

Lines changed: 10 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@
1616
#include "flang/Optimizer/Dialect/FIRAttr.h"
1717
#include "flang/Optimizer/Dialect/FIRType.h"
1818
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
19+
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
1920
#include "mlir/IR/Attributes.h"
2021
#include "mlir/IR/BuiltinAttributes.h"
2122
#include "mlir/IR/BuiltinOps.h"
@@ -280,14 +281,18 @@ mlir::LogicalResult cuf::RegisterKernelOp::verify() {
280281
return emitOpError("gpu module not found");
281282

282283
mlir::SymbolTable gpuSymTab(gpuMod);
283-
auto func = gpuSymTab.lookup<mlir::gpu::GPUFuncOp>(getKernelName());
284-
if (func) {
285-
// Only check if the gpu.func is found. It might be converted to LLVMFuncOp
286-
// already.
284+
if (auto func = gpuSymTab.lookup<mlir::gpu::GPUFuncOp>(getKernelName())) {
287285
if (!func.isKernel())
288286
return emitOpError("only kernel gpu.func can be registered");
287+
return mlir::success();
288+
} else if (auto func =
289+
gpuSymTab.lookup<mlir::LLVM::LLVMFuncOp>(getKernelName())) {
290+
if (!func->getAttrOfType<mlir::UnitAttr>(
291+
mlir::gpu::GPUDialect::getKernelFuncAttrName()))
292+
return emitOpError("only gpu.kernel llvm.func can be registered");
293+
return mlir::success();
289294
}
290-
return mlir::success();
295+
return emitOpError("device function not found");
291296
}
292297

293298
// Tablegen operators

flang/test/Fir/cuf-invalid.fir

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -143,6 +143,21 @@ module attributes {gpu.container_module} {
143143

144144
// -----
145145

146+
module attributes {gpu.container_module} {
147+
gpu.module @cuda_device_mod {
148+
gpu.func @_QPsub_device1() {
149+
gpu.return
150+
}
151+
}
152+
llvm.func internal @__cudaFortranConstructor() {
153+
// expected-error@+1{{'cuf.register_kernel' op device function not found}}
154+
cuf.register_kernel @cuda_device_mod::@_QPsub_device2
155+
llvm.return
156+
}
157+
}
158+
159+
// -----
160+
146161
module attributes {gpu.container_module} {
147162
llvm.func internal @__cudaFortranConstructor() {
148163
// expected-error@+1{{'cuf.register_kernel' op gpu module not found}}
@@ -160,3 +175,18 @@ module attributes {gpu.container_module} {
160175
llvm.return
161176
}
162177
}
178+
179+
// -----
180+
181+
module attributes {gpu.container_module} {
182+
gpu.module @cuda_device_mod {
183+
llvm.func @_QPsub_device1() {
184+
llvm.return
185+
}
186+
}
187+
llvm.func internal @__cudaFortranConstructor() {
188+
// expected-error@+1{{'cuf.register_kernel' op only gpu.kernel llvm.func can be registered}}
189+
cuf.register_kernel @cuda_device_mod::@_QPsub_device1
190+
llvm.return
191+
}
192+
}

0 commit comments

Comments
 (0)