Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2025-01-14 12:12:07 +00:00
Code Issues Actions 4 Packages Projects Releases Wiki Activity

[MLIR] Use nested symbol to identify kernel in LaunchFuncOp.

Browse Source 
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
This commit is contained in:
Frederik Gossen 2020-04-21 10:16:41 +00:00
parent 648fc95083
commit 0372db05bb
19 changed files with 98 additions and 101 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
mlir/include/mlir/Dialect/GPU/GPUOps.td
View File

@ -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); }];

16
mlir/include/mlir/IR/SymbolTable.h
View File

@ -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 <typename T>
static T lookupNearestSymbolFrom(Operation *from, StringRef symbol) {
return dyn_cast_or_null<T>(lookupNearestSymbolFrom(from, symbol));
}
template <typename T>
static T lookupNearestSymbolFrom(Operation *from, SymbolRefAttr symbol) {
return dyn_cast_or_null<T>(lookupNearestSymbolFrom(from, symbol));
}
/// This class represents a specific symbol use.
class SymbolUse {
@ -227,6 +236,13 @@ public:
template <typename T> T lookupSymbol(StringRef name) {
return dyn_cast_or_null<T>(lookupSymbol(name));
}
Operation *lookupSymbol(SymbolRefAttr symbol) {
return mlir::SymbolTable::lookupSymbolIn(this->getOperation(), symbol);
}
template <typename T>
T lookupSymbol(SymbolRefAttr symbol) {
return dyn_cast_or_null<T>(lookupSymbol(symbol));
}
};
/// A trait used to define a symbol that can be used on operations within a

14
mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp
View File

@ -273,14 +273,8 @@ Value GpuLaunchFuncToCudaCallsPass::setupParamsArray(gpu::LaunchFuncOp launchOp,
OpBuilder &builder) {
// Get the launch target.
auto containingModule = launchOp.getParentOfType<ModuleOp>();
if (!containingModule)
return {};
auto gpuModule = containingModule.lookupSymbol<gpu::GPUModuleOp>(
launchOp.getKernelModuleName());
if (!gpuModule)
return {};
auto gpuFunc = gpuModule.lookupSymbol<LLVM::LLVMFuncOp>(launchOp.kernel());
auto gpuFunc = SymbolTable::lookupNearestSymbolFrom<LLVM::LLVMFuncOp>(
launchOp, launchOp.kernel());
if (!gpuFunc)
return {};
@ -416,8 +410,8 @@ void GpuLaunchFuncToCudaCallsPass::translateGpuLaunchCalls(
// the kernel function.
auto cuOwningModuleRef =
builder.create<LLVM::LoadOp>(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<LLVM::LLVMFuncOp>(cuModuleGetFunctionName);

2
mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
View File

@ -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();
}

40
mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
View File

@ -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<StringAttr>(
LaunchFuncOp::getKernelAttrName()) ||
!launchOp.getAttrOfType<SymbolRefAttr>(
LaunchFuncOp::getKernelModuleAttrName()))
if (!launchOp.getAttrOfType<SymbolRefAttr>(
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<gpu::GPUFuncOp>(kernelFunc);
auto kernelLLVMFunction = dyn_cast_or_null<LLVM::LLVMFuncOp>(kernelFunc);
if (!kernelGPUFunction && !kernelLLVMFunction)
return launchOp.emitOpError("kernel function '")
<< kernelName << "' is undefined";
<< launchOp.kernel() << "' is undefined";
if (!kernelFunc->getAttrOfType<mlir::UnitAttr>(
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<GPUModuleOp>();
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<StringAttr>(getKernelAttrName()).getValue();
SymbolRefAttr LaunchFuncOp::kernel() {
return getAttrOfType<SymbolRefAttr>(getKernelAttrName());
}
unsigned LaunchFuncOp::getNumKernelOperands() {
@ -420,10 +419,11 @@ unsigned LaunchFuncOp::getNumKernelOperands() {
}
StringRef LaunchFuncOp::getKernelModuleName() {
return getAttrOfType<SymbolRefAttr>(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<StringAttr>(op.getKernelAttrName());
auto kernelAttr = op.getAttrOfType<SymbolRefAttr>(op.getKernelAttrName());
if (!kernelAttr)
return op.emitOpError("string attribute '" + op.getKernelAttrName() +
"' must be specified");
auto kernelModuleAttr =
op.getAttrOfType<SymbolRefAttr>(op.getKernelModuleAttrName());
if (!kernelModuleAttr)
return op.emitOpError("symbol reference attribute '" +
op.getKernelModuleAttrName() + "' must be specified");
op.getKernelAttrName() + "' must be specified");
return success();
}

2
mlir/test/Conversion/GPUToCUDA/lower-launch-func-to-cuda.mlir
View File

@ -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

16
mlir/test/Conversion/GPUToSPIRV/builtins.mlir
View File

@ -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
}

2
mlir/test/Conversion/GPUToSPIRV/if.mlir
View File

@ -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
}

2
mlir/test/Conversion/GPUToSPIRV/load-store.mlir
View File

@ -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
}

2
mlir/test/Conversion/GPUToSPIRV/loop.mlir
View File

@ -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
}

4
mlir/test/Conversion/GPUToSPIRV/simple.mlir
View File

@ -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
}

2
mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
View File

@ -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
}
}

59
mlir/test/Dialect/GPU/invalid.mlir
View File

@ -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
@ -63,20 +63,9 @@ 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
@ -85,23 +74,11 @@ 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
}
@ -138,6 +115,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*">) {
@ -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
}

8
mlir/test/Dialect/GPU/ops.mlir
View File

@ -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<?xf32, 1>) -> ()
// CHECK: "gpu.launch_func"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) {kernel = @kernels::@kernel_1} : (index, index, index, index, index, index, f32, memref<?xf32, 1>) -> ()
"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<?xf32, 1>) -> ()
// CHECK: "gpu.launch_func"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) {kernel = "kernel_2", kernel_module = @kernels} : (index, index, index, index, index, index, f32, memref<?xf32, 1>) -> ()
// CHECK: "gpu.launch_func"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) {kernel = @kernels::@kernel_2} : (index, index, index, index, index, index, f32, memref<?xf32, 1>) -> ()
"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<?xf32, 1>) -> ()
return

8
mlir/test/Dialect/GPU/outlining.mlir
View File

@ -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<?xf32, 1>) -> ()
// 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<?xf32, 1>) -> ()
// 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<?xf32>) {
%cst = constant 8 : index
%cst2 = constant 2 : index
%cst3 = dim %arg0, 0 : memref<?xf32>
// 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<?xf32>) -> ()
// CHECK: "gpu.launch_func"(%[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]], %[[CST]], %{{.*}}) {kernel = @extra_constants_kernel::@extra_constants_kernel} : (index, index, index, index, index, index, memref<?xf32>) -> ()
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,

2
mlir/test/mlir-vulkan-runner/addf.mlir
View File

@ -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<?xf32> to memref<*xf32>
call @print_memref_f32(%arg6) : (memref<*xf32>) -> ()

2
mlir/test/mlir-vulkan-runner/mulf.mlir
View File

@ -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<?x?xf32> to memref<*xf32>
call @print_memref_f32(%arg6) : (memref<*xf32>) -> ()

2
mlir/test/mlir-vulkan-runner/subf.mlir
View File

@ -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<?x?x?xf32> to memref<*xf32>
call @print_memref_f32(%arg6) : (memref<*xf32>) -> ()

2
mlir/test/mlir-vulkan-runner/time.mlir
View File

@ -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<?xf32> to memref<*xf32>
return