[Mlir-commits] [mlir] d48b84e - [MLIR][GPUToSPIRV] Passing gpu module name to SPIR-V module
George Mitenkov
llvmlistbot at llvm.org
Wed Aug 26 23:19:57 PDT 2020
Author: George Mitenkov
Date: 2020-08-27T09:19:24+03:00
New Revision: d48b84eb8a902a00866ef5c26dc05a7df830dbee
URL: https://github.com/llvm/llvm-project/commit/d48b84eb8a902a00866ef5c26dc05a7df830dbee
DIFF: https://github.com/llvm/llvm-project/commit/d48b84eb8a902a00866ef5c26dc05a7df830dbee.diff
LOG: [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
Added:
Modified:
mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp
mlir/test/Conversion/GPUToSPIRV/builtins.mlir
mlir/test/Conversion/GPUToSPIRV/load-store.mlir
mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir
mlir/test/Conversion/GPUToSPIRV/simple.mlir
Removed:
################################################################################
diff --git a/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp
index af44b59ba309..9fe35693bfe2 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp
@@ -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();
diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
index 173ebad6ffbb..b4f05fa27127 100644
--- a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
@@ -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
diff --git a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir
index f33ee9783a51..da57db15bedc 100644
--- a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir
@@ -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>
diff --git a/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir b/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir
index 027bfa80a34c..1b5b4d52d8b8 100644
--- a/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir
@@ -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
diff --git a/mlir/test/Conversion/GPUToSPIRV/simple.mlir b/mlir/test/Conversion/GPUToSPIRV/simple.mlir
index 0ecb83ebdcee..0c25c296efa2 100644
--- a/mlir/test/Conversion/GPUToSPIRV/simple.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/simple.mlir
@@ -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>
More information about the Mlir-commits
mailing list