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);
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); }];
#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"
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 {
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
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 {};
// 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);
// Set entry point name as an attribute.
vulkanLaunchCallOp.setAttr(
kSPIRVEntryPointAttrName,
- StringAttr::get(launchOp.kernel(), loc->getContext()));
+ StringAttr::get(launchOp.getKernelName(), loc->getContext()));
launchOp.erase();
}
//===----------------------------------------------------------------------===//
#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"
// 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.
<< "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 '")
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,
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() {
}
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);
}
"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();
}
// 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
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
}
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
}
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
}
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
}
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
}
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
}
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
}
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
}
} {
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
}
%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
}
} {
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
}
%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
}
%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
}
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
}
}
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
// -----
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
// -----
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
}
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
}
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
}
// -----
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
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
}
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
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
}
// 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
// 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)
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,
%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,
%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>) -> ()
%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>) -> ()
%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>) -> ()
%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