From 0372db05bb1552c2b39fc735f949977e0a863a25 Mon Sep 17 00:00:00 2001 From: Frederik Gossen Date: Tue, 21 Apr 2020 10:16:41 +0000 Subject: [PATCH] [MLIR] Use nested symbol to identify kernel in `LaunchFuncOp`. Summary: Use a nested symbol to identify the kernel to be invoked by a `LaunchFuncOp` in the GPU dialect. This replaces the two attributes that were used to identify the kernel module and the kernel within seperately. Differential Revision: https://reviews.llvm.org/D78551 --- mlir/include/mlir/Dialect/GPU/GPUOps.td | 14 +++-- mlir/include/mlir/IR/SymbolTable.h | 16 ++++++ .../GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp | 14 ++--- .../ConvertGPULaunchFuncToVulkanLaunchFunc.cpp | 2 +- mlir/lib/Dialect/GPU/IR/GPUDialect.cpp | 40 +++++++-------- .../GPUToCUDA/lower-launch-func-to-cuda.mlir | 2 +- mlir/test/Conversion/GPUToSPIRV/builtins.mlir | 16 +++--- mlir/test/Conversion/GPUToSPIRV/if.mlir | 2 +- mlir/test/Conversion/GPUToSPIRV/load-store.mlir | 2 +- mlir/test/Conversion/GPUToSPIRV/loop.mlir | 2 +- mlir/test/Conversion/GPUToSPIRV/simple.mlir | 4 +- .../lower-gpu-launch-vulkan-launch.mlir | 2 +- mlir/test/Dialect/GPU/invalid.mlir | 59 ++++++++++------------ mlir/test/Dialect/GPU/ops.mlir | 8 +-- mlir/test/Dialect/GPU/outlining.mlir | 8 +-- mlir/test/mlir-vulkan-runner/addf.mlir | 2 +- mlir/test/mlir-vulkan-runner/mulf.mlir | 2 +- mlir/test/mlir-vulkan-runner/subf.mlir | 2 +- mlir/test/mlir-vulkan-runner/time.mlir | 2 +- 19 files changed, 98 insertions(+), 101 deletions(-) diff --git a/mlir/include/mlir/Dialect/GPU/GPUOps.td b/mlir/include/mlir/Dialect/GPU/GPUOps.td index 5d91ff6..342b36b 100644 --- a/mlir/include/mlir/Dialect/GPU/GPUOps.td +++ b/mlir/include/mlir/Dialect/GPU/GPUOps.td @@ -334,15 +334,17 @@ def GPU_LaunchFuncOp : GPU_Op<"launch_func">, let extraClassDeclaration = [{ /// The kernel function specified by the operation's `kernel` attribute. - StringRef kernel(); + SymbolRefAttr kernel(); /// The number of operands passed to the kernel function. unsigned getNumKernelOperands(); - /// The name of the kernel module specified by the operation's - /// `kernel_module` attribute. + /// The name of the kernel's containing module. StringRef getKernelModuleName(); + /// The name of the kernel. + StringRef getKernelName(); + /// The i-th operand passed to the kernel function. Value getKernelOperand(unsigned i); @@ -361,12 +363,8 @@ def GPU_LaunchFuncOp : GPU_Op<"launch_func">, friend LogicalResult GPUDialect::verifyOperationAttribute(Operation *, NamedAttribute); - /// The name of the symbolRef attribute specifying the kernel to launch. + /// The name of the symbol reference attribute specifying the kernel to launch. static StringRef getKernelAttrName() { return "kernel"; } - - /// The name of the symbolRef attribute specifying the name of the module - /// containing the kernel to launch. - static StringRef getKernelModuleAttrName() { return "kernel_module"; } }]; let verifier = [{ return ::verify(*this); }]; diff --git a/mlir/include/mlir/IR/SymbolTable.h b/mlir/include/mlir/IR/SymbolTable.h index 6f5c07e..c61efb0 100644 --- a/mlir/include/mlir/IR/SymbolTable.h +++ b/mlir/include/mlir/IR/SymbolTable.h @@ -9,6 +9,7 @@ #ifndef MLIR_IR_SYMBOLTABLE_H #define MLIR_IR_SYMBOLTABLE_H +#include "mlir/IR/Attributes.h" #include "mlir/IR/OpDefinition.h" #include "llvm/ADT/StringMap.h" @@ -106,6 +107,14 @@ public: static Operation *lookupNearestSymbolFrom(Operation *from, StringRef symbol); static Operation *lookupNearestSymbolFrom(Operation *from, SymbolRefAttr symbol); + template + static T lookupNearestSymbolFrom(Operation *from, StringRef symbol) { + return dyn_cast_or_null(lookupNearestSymbolFrom(from, symbol)); + } + template + static T lookupNearestSymbolFrom(Operation *from, SymbolRefAttr symbol) { + return dyn_cast_or_null(lookupNearestSymbolFrom(from, symbol)); + } /// This class represents a specific symbol use. class SymbolUse { @@ -227,6 +236,13 @@ public: template T lookupSymbol(StringRef name) { return dyn_cast_or_null(lookupSymbol(name)); } + Operation *lookupSymbol(SymbolRefAttr symbol) { + return mlir::SymbolTable::lookupSymbolIn(this->getOperation(), symbol); + } + template + T lookupSymbol(SymbolRefAttr symbol) { + return dyn_cast_or_null(lookupSymbol(symbol)); + } }; /// A trait used to define a symbol that can be used on operations within a diff --git a/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp b/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp index bdd9bb6..cfdcb0f 100644 --- a/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp +++ b/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp @@ -273,14 +273,8 @@ Value GpuLaunchFuncToCudaCallsPass::setupParamsArray(gpu::LaunchFuncOp launchOp, OpBuilder &builder) { // Get the launch target. - auto containingModule = launchOp.getParentOfType(); - if (!containingModule) - return {}; - auto gpuModule = containingModule.lookupSymbol( - launchOp.getKernelModuleName()); - if (!gpuModule) - return {}; - auto gpuFunc = gpuModule.lookupSymbol(launchOp.kernel()); + auto gpuFunc = SymbolTable::lookupNearestSymbolFrom( + launchOp, launchOp.kernel()); if (!gpuFunc) return {}; @@ -416,8 +410,8 @@ void GpuLaunchFuncToCudaCallsPass::translateGpuLaunchCalls( // the kernel function. auto cuOwningModuleRef = builder.create(loc, getPointerType(), cuModule); - auto kernelName = generateKernelNameConstant(launchOp.getKernelModuleName(), - launchOp.kernel(), loc, builder); + auto kernelName = generateKernelNameConstant( + launchOp.getKernelModuleName(), launchOp.getKernelName(), loc, builder); auto cuFunction = allocatePointer(builder, loc); auto cuModuleGetFunction = getOperation().lookupSymbol(cuModuleGetFunctionName); diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp index b33edb9..2658804 100644 --- a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp +++ b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp @@ -182,7 +182,7 @@ void ConvertGpuLaunchFuncToVulkanLaunchFunc::convertGpuLaunchFunc( // Set entry point name as an attribute. vulkanLaunchCallOp.setAttr( kSPIRVEntryPointAttrName, - StringAttr::get(launchOp.kernel(), loc->getContext())); + StringAttr::get(launchOp.getKernelName(), loc->getContext())); launchOp.erase(); } diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp index 1474be7..e751107 100644 --- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp +++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp @@ -11,8 +11,10 @@ //===----------------------------------------------------------------------===// #include "mlir/Dialect/GPU/GPUDialect.h" + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" +#include "mlir/IR/Attributes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/Function.h" #include "mlir/IR/FunctionImplementation.h" @@ -62,10 +64,8 @@ LogicalResult GPUDialect::verifyOperationAttribute(Operation *op, // Ignore launch ops with missing attributes here. The errors will be // reported by the verifiers of those ops. - if (!launchOp.getAttrOfType( - LaunchFuncOp::getKernelAttrName()) || - !launchOp.getAttrOfType( - LaunchFuncOp::getKernelModuleAttrName())) + if (!launchOp.getAttrOfType( + LaunchFuncOp::getKernelAttrName())) return success(); // Check that `launch_func` refers to a well-formed GPU kernel module. @@ -76,13 +76,12 @@ LogicalResult GPUDialect::verifyOperationAttribute(Operation *op, << "kernel module '" << kernelModuleName << "' is undefined"; // Check that `launch_func` refers to a well-formed kernel function. - StringRef kernelName = launchOp.kernel(); - Operation *kernelFunc = kernelModule.lookupSymbol(kernelName); + Operation *kernelFunc = module.lookupSymbol(launchOp.kernel()); auto kernelGPUFunction = dyn_cast_or_null(kernelFunc); auto kernelLLVMFunction = dyn_cast_or_null(kernelFunc); if (!kernelGPUFunction && !kernelLLVMFunction) return launchOp.emitOpError("kernel function '") - << kernelName << "' is undefined"; + << launchOp.kernel() << "' is undefined"; if (!kernelFunc->getAttrOfType( GPUDialect::getKernelFuncAttrName())) return launchOp.emitOpError("kernel function is missing the '") @@ -397,11 +396,11 @@ void LaunchFuncOp::build(Builder *builder, OperationState &result, result.addOperands( {gridSizeX, gridSizeY, gridSizeZ, blockSizeX, blockSizeY, blockSizeZ}); result.addOperands(kernelOperands); - result.addAttribute(getKernelAttrName(), - builder->getStringAttr(kernelFunc.getName())); auto kernelModule = kernelFunc.getParentOfType(); - result.addAttribute(getKernelModuleAttrName(), - builder->getSymbolRefAttr(kernelModule.getName())); + auto kernelSymbol = builder->getSymbolRefAttr( + kernelModule.getName(), + {builder->getSymbolRefAttr(kernelFunc.getName())}); + result.addAttribute(getKernelAttrName(), kernelSymbol); } void LaunchFuncOp::build(Builder *builder, OperationState &result, @@ -411,8 +410,8 @@ void LaunchFuncOp::build(Builder *builder, OperationState &result, blockSize.x, blockSize.y, blockSize.z, kernelOperands); } -StringRef LaunchFuncOp::kernel() { - return getAttrOfType(getKernelAttrName()).getValue(); +SymbolRefAttr LaunchFuncOp::kernel() { + return getAttrOfType(getKernelAttrName()); } unsigned LaunchFuncOp::getNumKernelOperands() { @@ -420,10 +419,11 @@ unsigned LaunchFuncOp::getNumKernelOperands() { } StringRef LaunchFuncOp::getKernelModuleName() { - return getAttrOfType(getKernelModuleAttrName()) - .getRootReference(); + return kernel().getRootReference(); } +StringRef LaunchFuncOp::getKernelName() { return kernel().getLeafReference(); } + Value LaunchFuncOp::getKernelOperand(unsigned i) { return getOperation()->getOperand(i + kNumConfigOperands); } @@ -446,16 +446,10 @@ static LogicalResult verify(LaunchFuncOp op) { "expected the closest surrounding module to have the '" + GPUDialect::getContainerModuleAttrName() + "' attribute"); - auto kernelAttr = op.getAttrOfType(op.getKernelAttrName()); + auto kernelAttr = op.getAttrOfType(op.getKernelAttrName()); if (!kernelAttr) - return op.emitOpError("string attribute '" + op.getKernelAttrName() + - "' must be specified"); - - auto kernelModuleAttr = - op.getAttrOfType(op.getKernelModuleAttrName()); - if (!kernelModuleAttr) return op.emitOpError("symbol reference attribute '" + - op.getKernelModuleAttrName() + "' must be specified"); + op.getKernelAttrName() + "' must be specified"); return success(); } diff --git a/mlir/test/Conversion/GPUToCUDA/lower-launch-func-to-cuda.mlir b/mlir/test/Conversion/GPUToCUDA/lower-launch-func-to-cuda.mlir index 17244ce..20b76a2 100644 --- a/mlir/test/Conversion/GPUToCUDA/lower-launch-func-to-cuda.mlir +++ b/mlir/test/Conversion/GPUToCUDA/lower-launch-func-to-cuda.mlir @@ -27,7 +27,7 @@ module attributes {gpu.container_module} { // CHECK: llvm.call @mcuGetStreamHelper // CHECK: llvm.call @mcuLaunchKernel // CHECK: llvm.call @mcuStreamSynchronize - "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = "kernel", kernel_module = @kernel_module } + "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = @kernel_module::@kernel } : (!llvm.i64, !llvm.i64, !llvm.i64, !llvm.i64, !llvm.i64, !llvm.i64, !llvm.float, !llvm<"float*">) -> () llvm.return diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir index 2a73884..84afa22 100644 --- a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir @@ -3,7 +3,7 @@ module attributes {gpu.container_module} { func @builtin() { %c0 = constant 1 : index - "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_id_x", kernel_module = @kernels} : (index, index, index, index, index, index) -> () + "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_id_x} : (index, index, index, index, index, index) -> () return } @@ -26,7 +26,7 @@ module attributes {gpu.container_module} { module attributes {gpu.container_module} { func @builtin() { %c0 = constant 1 : index - "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_id_y", kernel_module = @kernels} : (index, index, index, index, index, index) -> () + "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_id_y} : (index, index, index, index, index, index) -> () return } @@ -49,7 +49,7 @@ module attributes {gpu.container_module} { module attributes {gpu.container_module} { func @builtin() { %c0 = constant 1 : index - "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_id_z", kernel_module = @kernels} : (index, index, index, index, index, index) -> () + "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_id_z} : (index, index, index, index, index, index) -> () return } @@ -72,7 +72,7 @@ module attributes {gpu.container_module} { module attributes {gpu.container_module} { func @builtin() { %c0 = constant 1 : index - "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_size_x", kernel_module = @kernels} : (index, index, index, index, index, index) -> () + "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_size_x} : (index, index, index, index, index, index) -> () return } @@ -96,7 +96,7 @@ module attributes {gpu.container_module} { module attributes {gpu.container_module} { func @builtin() { %c0 = constant 1 : index - "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_size_y", kernel_module = @kernels} : (index, index, index, index, index, index) -> () + "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_size_y} : (index, index, index, index, index, index) -> () return } @@ -117,7 +117,7 @@ module attributes {gpu.container_module} { module attributes {gpu.container_module} { func @builtin() { %c0 = constant 1 : index - "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_size_z", kernel_module = @kernels} : (index, index, index, index, index, index) -> () + "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_workgroup_size_z} : (index, index, index, index, index, index) -> () return } @@ -138,7 +138,7 @@ module attributes {gpu.container_module} { module attributes {gpu.container_module} { func @builtin() { %c0 = constant 1 : index - "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_local_id_x", kernel_module = @kernels} : (index, index, index, index, index, index) -> () + "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_local_id_x} : (index, index, index, index, index, index) -> () return } @@ -161,7 +161,7 @@ module attributes {gpu.container_module} { module attributes {gpu.container_module} { func @builtin() { %c0 = constant 1 : index - "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_num_workgroups_x", kernel_module = @kernels} : (index, index, index, index, index, index) -> () + "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = @kernels::@builtin_num_workgroups_x} : (index, index, index, index, index, index) -> () return } diff --git a/mlir/test/Conversion/GPUToSPIRV/if.mlir b/mlir/test/Conversion/GPUToSPIRV/if.mlir index 3fefc04..8a63745 100644 --- a/mlir/test/Conversion/GPUToSPIRV/if.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/if.mlir @@ -9,7 +9,7 @@ module attributes { } { func @main(%arg0 : memref<10xf32>, %arg1 : i1) { %c0 = constant 1 : index - "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0, %arg0, %arg1) { kernel = "kernel_simple_selection", kernel_module = @kernels} : (index, index, index, index, index, index, memref<10xf32>, i1) -> () + "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0, %arg0, %arg1) { kernel = @kernels::@kernel_simple_selection} : (index, index, index, index, index, index, memref<10xf32>, i1) -> () return } diff --git a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir index acb18e7..077a1c0 100644 --- a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir @@ -17,7 +17,7 @@ module attributes { %1 = subi %c4, %c0_0 : index %c1_1 = constant 1 : index %c1_2 = constant 1 : index - "gpu.launch_func"(%0, %c1_2, %c1_2, %1, %c1_2, %c1_2, %arg0, %arg1, %arg2, %c0, %c0_0, %c1, %c1_1) {kernel = "load_store_kernel", kernel_module = @kernels} : (index, index, index, index, index, index, memref<12x4xf32>, memref<12x4xf32>, memref<12x4xf32>, index, index, index, index) -> () + "gpu.launch_func"(%0, %c1_2, %c1_2, %1, %c1_2, %c1_2, %arg0, %arg1, %arg2, %c0, %c0_0, %c1, %c1_1) {kernel = @kernels::@load_store_kernel} : (index, index, index, index, index, index, memref<12x4xf32>, memref<12x4xf32>, memref<12x4xf32>, index, index, index, index) -> () return } diff --git a/mlir/test/Conversion/GPUToSPIRV/loop.mlir b/mlir/test/Conversion/GPUToSPIRV/loop.mlir index 6f0b209..56bff8a 100644 --- a/mlir/test/Conversion/GPUToSPIRV/loop.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/loop.mlir @@ -9,7 +9,7 @@ module attributes { } { func @loop(%arg0 : memref<10xf32>, %arg1 : memref<10xf32>) { %c0 = constant 1 : index - "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0, %arg0, %arg1) { kernel = "loop_kernel", kernel_module = @kernels} : (index, index, index, index, index, index, memref<10xf32>, memref<10xf32>) -> () + "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0, %arg0, %arg1) { kernel = @kernels::@loop_kernel} : (index, index, index, index, index, index, memref<10xf32>, memref<10xf32>) -> () return } diff --git a/mlir/test/Conversion/GPUToSPIRV/simple.mlir b/mlir/test/Conversion/GPUToSPIRV/simple.mlir index c657d5f..f688233 100644 --- a/mlir/test/Conversion/GPUToSPIRV/simple.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/simple.mlir @@ -18,7 +18,7 @@ module attributes {gpu.container_module} { %0 = "op"() : () -> (f32) %1 = "op"() : () -> (memref<12xf32>) %cst = constant 1 : index - "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = "basic_module_structure", kernel_module = @kernels } + "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = @kernels::@basic_module_structure } : (index, index, index, index, index, index, f32, memref<12xf32>) -> () return } @@ -39,7 +39,7 @@ module attributes {gpu.container_module} { %0 = "op"() : () -> (f32) %1 = "op"() : () -> (memref<12xf32>) %cst = constant 1 : index - "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = "missing_entry_point_abi", kernel_module = @kernels } + "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = @kernels::@missing_entry_point_abi } : (index, index, index, index, index, index, f32, memref<12xf32>) -> () return } diff --git a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir index 30dc3a4..726b276 100644 --- a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir +++ b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir @@ -26,7 +26,7 @@ module attributes {gpu.container_module} { func @foo() { %0 = alloc() : memref<12xf32> %c1 = constant 1 : index - "gpu.launch_func"(%c1, %c1, %c1, %c1, %c1, %c1, %0) {kernel = "kernel", kernel_module = @kernels} : (index, index, index, index, index, index, memref<12xf32>) -> () + "gpu.launch_func"(%c1, %c1, %c1, %c1, %c1, %c1, %0) {kernel = @kernels::@kernel} : (index, index, index, index, index, index, memref<12xf32>) -> () return } } diff --git a/mlir/test/Dialect/GPU/invalid.mlir b/mlir/test/Dialect/GPU/invalid.mlir index 885ad32..be02dec 100644 --- a/mlir/test/Dialect/GPU/invalid.mlir +++ b/mlir/test/Dialect/GPU/invalid.mlir @@ -54,7 +54,7 @@ func @launch_func_missing_parent_module_attribute(%sz : index) { module attributes {gpu.container_module} { func @launch_func_missing_callee_attribute(%sz : index) { - // expected-error@+1 {{string attribute 'kernel' must be specified}} + // expected-error@+1 {{symbol reference attribute 'kernel' must be specified}} "gpu.launch_func"(%sz, %sz, %sz, %sz, %sz, %sz) {foo = "bar"} : (index, index, index, index, index, index) -> () return @@ -64,19 +64,8 @@ module attributes {gpu.container_module} { // ----- module attributes {gpu.container_module} { - func @launch_func_missing_module_attribute(%sz : index) { - // expected-error@+1 {{attribute 'kernel_module' must be specified}} - "gpu.launch_func"(%sz, %sz, %sz, %sz, %sz, %sz) {kernel = "launch_func_missing_kernel_attr"} - : (index, index, index, index, index, index) -> () - return - } -} - -// ----- - -module attributes {gpu.container_module} { func @launch_func_no_function_attribute(%sz : index) { - // expected-error@+1 {{string attribute 'kernel' must be specified}} + // expected-error@+1 {{symbol reference attribute 'kernel' must be specified}} "gpu.launch_func"(%sz, %sz, %sz, %sz, %sz, %sz) {kernel = 10} : (index, index, index, index, index, index) -> () return @@ -86,22 +75,10 @@ module attributes {gpu.container_module} { // ----- module attributes {gpu.container_module} { - func @launch_func_module_attribute_wrong_type(%sz : index) { - // expected-error@+1 {{symbol reference attribute 'kernel_module' must be specified}} - "gpu.launch_func"(%sz, %sz, %sz, %sz, %sz, %sz) - {kernel = "launch_func_module_attribute_wrong_type", kernel_module = 10} - : (index, index, index, index, index, index) -> () - return - } -} - -// ----- - -module attributes {gpu.container_module} { func @launch_func_undefined_module(%sz : index) { // expected-error@+1 {{kernel module 'kernels' is undefined}} "gpu.launch_func"(%sz, %sz, %sz, %sz, %sz, %sz) - { kernel = "kernel_1", kernel_module = @kernels } + { kernel = @kernels::@kernel_1 } : (index, index, index, index, index, index) -> () return } @@ -116,7 +93,7 @@ module attributes {gpu.container_module} { func @launch_func_missing_module_attribute(%sz : index) { // expected-error@+1 {{kernel module 'kernels' is undefined}} "gpu.launch_func"(%sz, %sz, %sz, %sz, %sz, %sz) - { kernel = "kernel_1", kernel_module = @kernels } + { kernel = @kernels::@kernel_1 } : (index, index, index, index, index, index) -> () return } @@ -128,9 +105,9 @@ module attributes {gpu.container_module} { gpu.module @kernels { } func @launch_func_undefined_function(%sz : index) { - // expected-error@+1 {{kernel function 'kernel_1' is undefined}} + // expected-error@+1 {{kernel function '@kernels::@kernel_1' is undefined}} "gpu.launch_func"(%sz, %sz, %sz, %sz, %sz, %sz) - { kernel = "kernel_1", kernel_module = @kernels } + { kernel = @kernels::@kernel_1 } : (index, index, index, index, index, index) -> () return } @@ -139,6 +116,24 @@ module attributes {gpu.container_module} { // ----- module attributes {gpu.container_module} { + module @kernels { + gpu.func @kernel_1(%arg1 : !llvm<"float*">) kernel { + gpu.return + } + } + + func @launch_func_missing_kernel_attr(%sz : index, %arg : !llvm<"float*">) { + // expected-error@+1 {{kernel module 'kernels' is undefined}} + "gpu.launch_func"(%sz, %sz, %sz, %sz, %sz, %sz, %arg) + {kernel = @kernels::@kernel_1} + : (index, index, index, index, index, index, !llvm<"float*">) -> () + return + } +} + +// ----- + +module attributes {gpu.container_module} { gpu.module @kernels { gpu.func @kernel_1(%arg1 : !llvm<"float*">) { gpu.return @@ -148,7 +143,7 @@ module attributes {gpu.container_module} { func @launch_func_missing_kernel_attr(%sz : index, %arg : !llvm<"float*">) { // expected-error@+1 {{kernel function is missing the 'gpu.kernel' attribute}} "gpu.launch_func"(%sz, %sz, %sz, %sz, %sz, %sz, %arg) - {kernel = "kernel_1", kernel_module = @kernels} + {kernel = @kernels::@kernel_1} : (index, index, index, index, index, index, !llvm<"float*">) -> () return } @@ -166,7 +161,7 @@ module attributes {gpu.container_module} { func @launch_func_kernel_operand_size(%sz : index, %arg : !llvm<"float*">) { // expected-error@+1 {{got 2 kernel operands but expected 1}} "gpu.launch_func"(%sz, %sz, %sz, %sz, %sz, %sz, %arg, %arg) - {kernel = "kernel_1", kernel_module = @kernels} + {kernel = @kernels::@kernel_1} : (index, index, index, index, index, index, !llvm<"float*">, !llvm<"float*">) -> () return @@ -185,7 +180,7 @@ module attributes {gpu.container_module} { func @launch_func_kernel_operand_types(%sz : index, %arg : f32) { // expected-err@+1 {{type of function argument 0 does not match}} "gpu.launch_func"(%sz, %sz, %sz, %sz, %sz, %sz, %arg) - {kernel = "kernel_1", kernel_module = @kernels} + {kernel = @kernels::@kernel_1} : (index, index, index, index, index, index, f32) -> () return } diff --git a/mlir/test/Dialect/GPU/ops.mlir b/mlir/test/Dialect/GPU/ops.mlir index 1cb1b53..f500d71 100644 --- a/mlir/test/Dialect/GPU/ops.mlir +++ b/mlir/test/Dialect/GPU/ops.mlir @@ -70,14 +70,14 @@ module attributes {gpu.container_module} { // CHECK: %{{.*}} = constant 8 %cst = constant 8 : index - // CHECK: "gpu.launch_func"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) {kernel = "kernel_1", kernel_module = @kernels} : (index, index, index, index, index, index, f32, memref) -> () + // CHECK: "gpu.launch_func"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) {kernel = @kernels::@kernel_1} : (index, index, index, index, index, index, f32, memref) -> () "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) - { kernel = "kernel_1", kernel_module = @kernels } + { kernel = @kernels::@kernel_1} : (index, index, index, index, index, index, f32, memref) -> () - // CHECK: "gpu.launch_func"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) {kernel = "kernel_2", kernel_module = @kernels} : (index, index, index, index, index, index, f32, memref) -> () + // CHECK: "gpu.launch_func"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) {kernel = @kernels::@kernel_2} : (index, index, index, index, index, index, f32, memref) -> () "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) - { kernel = "kernel_2", kernel_module = @kernels } + { kernel = @kernels::@kernel_2} : (index, index, index, index, index, index, f32, memref) -> () return diff --git a/mlir/test/Dialect/GPU/outlining.mlir b/mlir/test/Dialect/GPU/outlining.mlir index 0d87347..d15f10f 100644 --- a/mlir/test/Dialect/GPU/outlining.mlir +++ b/mlir/test/Dialect/GPU/outlining.mlir @@ -21,7 +21,7 @@ func @launch() { // CHECK: %[[BDIMZ:.*]] = constant 28 %bDimZ = constant 28 : index - // CHECK: "gpu.launch_func"(%[[GDIMX]], %[[GDIMY]], %[[GDIMZ]], %[[BDIMX]], %[[BDIMY]], %[[BDIMZ]], %[[ARG0]], %[[ARG1]]) {kernel = "launch_kernel", kernel_module = @launch_kernel} : (index, index, index, index, index, index, f32, memref) -> () + // CHECK: "gpu.launch_func"(%[[GDIMX]], %[[GDIMY]], %[[GDIMZ]], %[[BDIMX]], %[[BDIMY]], %[[BDIMZ]], %[[ARG0]], %[[ARG1]]) {kernel = @launch_kernel::@launch_kernel} : (index, index, index, index, index, index, f32, memref) -> () // CHECK-NOT: gpu.launch blocks gpu.launch blocks(%bx, %by, %bz) in (%grid_x = %gDimX, %grid_y = %gDimY, %grid_z = %gDimZ) @@ -64,14 +64,14 @@ func @launch() { func @multiple_launches() { // CHECK: %[[CST:.*]] = constant 8 : index %cst = constant 8 : index - // CHECK: "gpu.launch_func"(%[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]]) {kernel = "multiple_launches_kernel", kernel_module = @multiple_launches_kernel} : (index, index, index, index, index, index) -> () + // CHECK: "gpu.launch_func"(%[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]]) {kernel = @multiple_launches_kernel::@multiple_launches_kernel} : (index, index, index, index, index, index) -> () gpu.launch blocks(%bx, %by, %bz) in (%grid_x = %cst, %grid_y = %cst, %grid_z = %cst) threads(%tx, %ty, %tz) in (%block_x = %cst, %block_y = %cst, %block_z = %cst) { gpu.terminator } - // CHECK: "gpu.launch_func"(%[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]]) {kernel = "multiple_launches_kernel", kernel_module = @multiple_launches_kernel_0} : (index, index, index, index, index, index) -> () + // CHECK: "gpu.launch_func"(%[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]]) {kernel = @multiple_launches_kernel_0::@multiple_launches_kernel} : (index, index, index, index, index, index) -> () gpu.launch blocks(%bx2, %by2, %bz2) in (%grid_x2 = %cst, %grid_y2 = %cst, %grid_z2 = %cst) threads(%tx2, %ty2, %tz2) in (%block_x2 = %cst, %block_y2 = %cst, @@ -93,7 +93,7 @@ func @extra_constants(%arg0 : memref) { %cst = constant 8 : index %cst2 = constant 2 : index %cst3 = dim %arg0, 0 : memref - // CHECK: "gpu.launch_func"(%[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]], %{{.*}}) {kernel = "extra_constants_kernel", kernel_module = @extra_constants_kernel} : (index, index, index, index, index, index, memref) -> () + // CHECK: "gpu.launch_func"(%[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]], %{{.*}}) {kernel = @extra_constants_kernel::@extra_constants_kernel} : (index, index, index, index, index, index, memref) -> () gpu.launch blocks(%bx, %by, %bz) in (%grid_x = %cst, %grid_y = %cst, %grid_z = %cst) threads(%tx, %ty, %tz) in (%block_x = %cst, %block_y = %cst, diff --git a/mlir/test/mlir-vulkan-runner/addf.mlir b/mlir/test/mlir-vulkan-runner/addf.mlir index 2fb3a94..3ba86ef 100644 --- a/mlir/test/mlir-vulkan-runner/addf.mlir +++ b/mlir/test/mlir-vulkan-runner/addf.mlir @@ -39,7 +39,7 @@ module attributes { %cst1 = constant 1 : index %cst8 = constant 8 : index - "gpu.launch_func"(%cst8, %cst1, %cst1, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = "kernel_add", kernel_module = @kernels } + gpu.launch_func"(%cst8, %cst1, %cst1, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_add } : (index, index, index, index, index, index, memref<8xf32>, memref<8xf32>, memref<8xf32>) -> () %arg6 = memref_cast %arg5 : memref to memref<*xf32> call @print_memref_f32(%arg6) : (memref<*xf32>) -> () diff --git a/mlir/test/mlir-vulkan-runner/mulf.mlir b/mlir/test/mlir-vulkan-runner/mulf.mlir index 0da888b..89175e8 100644 --- a/mlir/test/mlir-vulkan-runner/mulf.mlir +++ b/mlir/test/mlir-vulkan-runner/mulf.mlir @@ -40,7 +40,7 @@ module attributes { %cst1 = constant 1 : index %cst4 = constant 4 : index - "gpu.launch_func"(%cst4, %cst4, %cst1, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = "kernel_mul", kernel_module = @kernels } + "gpu.launch_func"(%cst4, %cst4, %cst1, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_mul } : (index, index, index, index, index, index, memref<4x4xf32>, memref<4x4xf32>, memref<4x4xf32>) -> () %arg6 = memref_cast %arg5 : memref to memref<*xf32> call @print_memref_f32(%arg6) : (memref<*xf32>) -> () diff --git a/mlir/test/mlir-vulkan-runner/subf.mlir b/mlir/test/mlir-vulkan-runner/subf.mlir index c77a14b..b410946 100644 --- a/mlir/test/mlir-vulkan-runner/subf.mlir +++ b/mlir/test/mlir-vulkan-runner/subf.mlir @@ -42,7 +42,7 @@ module attributes { %cst1 = constant 1 : index %cst4 = constant 4 : index %cst8 = constant 8 : index - "gpu.launch_func"(%cst8, %cst4, %cst4, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = "kernel_sub", kernel_module = @kernels } + "gpu.launch_func"(%cst8, %cst4, %cst4, %cst1, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_sub } : (index, index, index, index, index, index, memref<8x4x4xf32>, memref<4x4xf32>, memref<8x4x4xf32>) -> () %arg6 = memref_cast %arg5 : memref to memref<*xf32> call @print_memref_f32(%arg6) : (memref<*xf32>) -> () diff --git a/mlir/test/mlir-vulkan-runner/time.mlir b/mlir/test/mlir-vulkan-runner/time.mlir index b95452e..ffa8985 100644 --- a/mlir/test/mlir-vulkan-runner/time.mlir +++ b/mlir/test/mlir-vulkan-runner/time.mlir @@ -46,7 +46,7 @@ module attributes { %cst1 = constant 1 : index %cst128 = constant 128 : index - "gpu.launch_func"(%cst128, %cst1, %cst1, %cst128, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = "kernel_add", kernel_module = @kernels } + "gpu.launch_func"(%cst128, %cst1, %cst1, %cst128, %cst1, %cst1, %arg0, %arg1, %arg2) { kernel = @kernels::@kernel_add } : (index, index, index, index, index, index, memref<16384xf32>, memref<16384xf32>, memref<16384xf32>) -> () %arg6 = memref_cast %arg5 : memref to memref<*xf32> return -- 2.7.4