maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[MLIR][GPUToSPIRV] Passing gpu module name to SPIR-V module 

This patch allows to pass the gpu module name to SPIR-V
module during conversion. This has many benefits as we can lookup
converted to SPIR-V kernel in the symbol table.

In order to avoid symbol conflicts, `"__spv__"` is added to the
gpu module name to form the new one.

Reviewed By: mravishankar

Differential Revision: https://reviews.llvm.org/D86384
This commit is contained in:
George Mitenkov 2020-08-27 08:50:57 +03:00
parent 888c02deee
commit d48b84eb8a
5 changed files with 22 additions and 17 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp
View File

@ -19,6 +19,8 @@
using namespace mlir;
static constexpr const char kSPIRVModule[] = "__spv__";
namespace {
/// Pattern lowering GPU block/thread size/id to loading SPIR-V invocation
/// builtin variables.
@ -285,8 +287,11 @@ LogicalResult GPUModuleConversion::matchAndRewrite(
return moduleOp.emitRemark("match failure: could not selected memory model "
"based on 'spv.target_env'");
auto spvModule = rewriter.create<spirv::ModuleOp>(
moduleOp.getLoc(), addressingModel, memoryModel.getValue());
// Add a keyword to the module name to avoid symbolic conflict.
auto spvModuleName = StringRef(kSPIRVModule + moduleOp.getName().str());
auto spvModule =
rewriter.create<spirv::ModuleOp>(moduleOp.getLoc(), addressingModel,
memoryModel.getValue(), spvModuleName);
// Move the region from the module op into the SPIR-V module.
Region &spvModuleRegion = spvModule.body();

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

@ -7,7 +7,7 @@ module attributes {gpu.container_module} {
return
}
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
// CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
gpu.module @kernels {
gpu.func @builtin_workgroup_id_x() kernel
@ -30,7 +30,7 @@ module attributes {gpu.container_module} {
return
}
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
// CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
gpu.module @kernels {
gpu.func @builtin_workgroup_id_y() kernel
@ -53,7 +53,7 @@ module attributes {gpu.container_module} {
return
}
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
// CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
gpu.module @kernels {
gpu.func @builtin_workgroup_id_z() kernel
@ -76,7 +76,7 @@ module attributes {gpu.container_module} {
return
}
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
gpu.module @kernels {
gpu.func @builtin_workgroup_size_x() kernel
attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]>: vector<3xi32>}} {
@ -100,7 +100,7 @@ module attributes {gpu.container_module} {
return
}
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
gpu.module @kernels {
gpu.func @builtin_workgroup_size_y() kernel
attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
@ -121,7 +121,7 @@ module attributes {gpu.container_module} {
return
}
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
gpu.module @kernels {
gpu.func @builtin_workgroup_size_z() kernel
attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
@ -142,7 +142,7 @@ module attributes {gpu.container_module} {
return
}
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
// CHECK: spv.globalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
gpu.module @kernels {
gpu.func @builtin_local_id_x() kernel
@ -165,7 +165,7 @@ module attributes {gpu.container_module} {
return
}
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
// CHECK: spv.globalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
gpu.module @kernels {
gpu.func @builtin_num_workgroups_x() kernel
@ -182,7 +182,7 @@ module attributes {gpu.container_module} {
// -----
module attributes {gpu.container_module} {
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
// CHECK: spv.globalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId")
gpu.module @kernels {
gpu.func @builtin_subgroup_id() kernel
@ -198,7 +198,7 @@ module attributes {gpu.container_module} {
// -----
module attributes {gpu.container_module} {
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
// CHECK: spv.globalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups")
gpu.module @kernels {
gpu.func @builtin_num_subgroups() kernel
@ -214,7 +214,7 @@ module attributes {gpu.container_module} {
// -----
module attributes {gpu.container_module} {
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
// CHECK: spv.globalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize")
gpu.module @kernels {
gpu.func @builtin_subgroup_size() kernel

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

@ -21,7 +21,7 @@ module attributes {
return
}
// CHECK-LABEL: spv.module Logical GLSL450
// CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
gpu.module @kernels {
// CHECK-DAG: spv.globalVariable @[[NUMWORKGROUPSVAR:.*]] built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
// CHECK-DAG: spv.globalVariable @[[$LOCALINVOCATIONIDVAR:.*]] built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>

2
mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir
View File

@ -8,7 +8,7 @@ module attributes {
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
gpu.module @kernels {
// CHECK-LABEL: spv.module Physical64 OpenCL
// CHECK-LABEL: spv.module @{{.*}} Physical64 OpenCL
// CHECK: spv.func
// CHECK-SAME: {{%.*}}: f32
// CHECK-NOT: spv.interface_var_abi

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

@ -2,7 +2,7 @@
module attributes {gpu.container_module} {
gpu.module @kernels {
// CHECK: spv.module Logical GLSL450 {
// CHECK: spv.module @{{.*}} Logical GLSL450 {
// CHECK-LABEL: spv.func @basic_module_structure
// CHECK-SAME: {{%.*}}: f32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>}
// CHECK-SAME: {{%.*}}: !spv.ptr<!spv.struct<!spv.array<12 x f32, stride=4> [0]>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
@ -28,7 +28,7 @@ module attributes {gpu.container_module} {
module attributes {gpu.container_module} {
gpu.module @kernels {
// CHECK: spv.module Logical GLSL450 {
// CHECK: spv.module @{{.*}} Logical GLSL450 {
// CHECK-LABEL: spv.func @basic_module_structure_preset_ABI
// CHECK-SAME: {{%[a-zA-Z0-9_]*}}: f32
// CHECK-SAME: spv.interface_var_abi = #spv.interface_var_abi<(1, 2), StorageBuffer>