[Mlir-commits] [mlir] [mlir][spirv] Fix lookup logic `spirv.target_env` for `gpu.module` (PR #147262)
Jaeho Kim
llvmlistbot at llvm.org
Wed Jul 30 20:02:42 PDT 2025
https://github.com/oojahooo updated https://github.com/llvm/llvm-project/pull/147262
>From 5d05b68214f9166dae933b6313852fa8f079e155 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Mon, 7 Jul 2025 18:15:11 +0900
Subject: [PATCH 01/16] [mlir][SPIRV] Fix lookup logic `spirv.target_env` for
`gpu.module`
The `gpu.module` operation can contain `spirv.target_env` attributes
within an array attribute named `"targets"`. So it accounts for that
case by iterating over the `"targets"` attribute, if present, and
looking up `spirv.target_env`.
---
mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp | 9 +++++++++
1 file changed, 9 insertions(+)
diff --git a/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp b/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp
index 5ecbd5d7c59d5..dbaa10e89bd42 100644
--- a/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp
+++ b/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp
@@ -184,6 +184,15 @@ spirv::TargetEnvAttr spirv::lookupTargetEnv(Operation *op) {
if (!op)
break;
+ if (auto arrAttr = op->getAttrOfType<ArrayAttr>("targets")) {
+ for (auto attr : arrAttr) {
+ if (auto spirvTargetEnvAttr =
+ llvm::dyn_cast<spirv::TargetEnvAttr>(attr)) {
+ return spirvTargetEnvAttr;
+ }
+ }
+ }
+
if (auto attr = op->getAttrOfType<spirv::TargetEnvAttr>(
spirv::getTargetEnvAttrName()))
return attr;
>From 71d3d92b442a332b20187e2c01698b062033a993 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Tue, 8 Jul 2025 18:24:35 +0900
Subject: [PATCH 02/16] Add test
---
.../GPUToSPIRV/lookup-target-env.mlir | 17 +++++++++++++++++
1 file changed, 17 insertions(+)
create mode 100644 mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
diff --git a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
new file mode 100644
index 0000000000000..7c74bbad49101
--- /dev/null
+++ b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
@@ -0,0 +1,17 @@
+// RUN: mlir-opt --spirv-attach-target='caps=Shader exts=SPV_KHR_storage_buffer_storage_class' --convert-gpu-to-spirv %s -o - | FileCheck %s
+
+module attributes {gpu.container_module} {
+ // CHECK-LABEL: spirv.module @{{.*}} GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
+ gpu.module @kernels {
+ // CHECK: spirv.func @load_kernel
+ // CHECK-SAME: %[[ARG:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>})
+ gpu.func @load_kernel(%arg0: memref<12x4xf32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
+ %c0 = arith.constant 0 : index
+ // CHECK: %[[PTR:.*]] = spirv.AccessChain %[[ARG]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
+ // CHECK-NEXT: {{%.*}} = spirv.Load "StorageBuffer" %[[PTR]] : f32
+ %0 = memref.load %arg0[%c0, %c0] : memref<12x4xf32>
+ // CHECK: spirv.Return
+ gpu.return
+ }
+ }
+}
>From ae6ba9cc1395906b28b849260fdd71024c31bab8 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Tue, 8 Jul 2025 18:33:15 +0900
Subject: [PATCH 03/16] Delete braces for complying with conding standard
It resolves
https://github.com/llvm/llvm-project/pull/147262#discussion_r2189731690
---
mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp | 6 ++----
1 file changed, 2 insertions(+), 4 deletions(-)
diff --git a/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp b/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp
index dbaa10e89bd42..c90db80b7b970 100644
--- a/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp
+++ b/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp
@@ -185,12 +185,10 @@ spirv::TargetEnvAttr spirv::lookupTargetEnv(Operation *op) {
break;
if (auto arrAttr = op->getAttrOfType<ArrayAttr>("targets")) {
- for (auto attr : arrAttr) {
+ for (auto attr : arrAttr)
if (auto spirvTargetEnvAttr =
- llvm::dyn_cast<spirv::TargetEnvAttr>(attr)) {
+ llvm::dyn_cast<spirv::TargetEnvAttr>(attr))
return spirvTargetEnvAttr;
- }
- }
}
if (auto attr = op->getAttrOfType<spirv::TargetEnvAttr>(
>From fa3b44e7027b91bffb97deb2172be0d2f574abe4 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Thu, 10 Jul 2025 11:03:37 +0900
Subject: [PATCH 04/16] fixup! [mlir][SPIRV] Fix lookup logic
`spirv.target_env` for `gpu.module`
Add lookup target env in "targets" attr logic to GPUToSPIRV pass
---
.../Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp | 27 ++++++++++++++++---
mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp | 7 -----
2 files changed, 23 insertions(+), 11 deletions(-)
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index 0b2c06a08db2d..d182198b82f47 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -48,9 +48,29 @@ struct GPUToSPIRVPass final : impl::ConvertGPUToSPIRVBase<GPUToSPIRVPass> {
void runOnOperation() override;
private:
+ spirv::TargetEnvAttr lookupTargetEnvInTargets(gpu::GPUModuleOp moduleOp);
+ spirv::TargetEnvAttr lookupTargetEnvOrDefault(gpu::GPUModuleOp moduleOp);
bool mapMemorySpace;
};
+spirv::TargetEnvAttr
+GPUToSPIRVPass::lookupTargetEnvInTargets(gpu::GPUModuleOp moduleOp) {
+ for (auto &targetAttr : moduleOp.getTargetsAttr())
+ if (auto spirvTargetEnvAttr =
+ llvm::dyn_cast<spirv::TargetEnvAttr>(targetAttr))
+ return spirvTargetEnvAttr;
+
+ return {};
+}
+
+spirv::TargetEnvAttr
+GPUToSPIRVPass::lookupTargetEnvOrDefault(gpu::GPUModuleOp moduleOp) {
+ if (auto targetEnvAttr = lookupTargetEnvInTargets(moduleOp))
+ return targetEnvAttr;
+
+ return spirv::lookupTargetEnvOrDefault(moduleOp);
+}
+
void GPUToSPIRVPass::runOnOperation() {
MLIRContext *context = &getContext();
ModuleOp module = getOperation();
@@ -58,9 +78,8 @@ void GPUToSPIRVPass::runOnOperation() {
SmallVector<Operation *, 1> gpuModules;
OpBuilder builder(context);
- auto targetEnvSupportsKernelCapability = [](gpu::GPUModuleOp moduleOp) {
- Operation *gpuModule = moduleOp.getOperation();
- auto targetAttr = spirv::lookupTargetEnvOrDefault(gpuModule);
+ auto targetEnvSupportsKernelCapability = [this](gpu::GPUModuleOp moduleOp) {
+ auto targetAttr = lookupTargetEnvOrDefault(moduleOp);
spirv::TargetEnv targetEnv(targetAttr);
return targetEnv.allows(spirv::Capability::Kernel);
};
@@ -86,7 +105,7 @@ void GPUToSPIRVPass::runOnOperation() {
// TargetEnv attributes.
for (Operation *gpuModule : gpuModules) {
spirv::TargetEnvAttr targetAttr =
- spirv::lookupTargetEnvOrDefault(gpuModule);
+ lookupTargetEnvOrDefault(llvm::cast<gpu::GPUModuleOp>(gpuModule));
// Map MemRef memory space to SPIR-V storage class first if requested.
if (mapMemorySpace) {
diff --git a/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp b/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp
index c90db80b7b970..5ecbd5d7c59d5 100644
--- a/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp
+++ b/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp
@@ -184,13 +184,6 @@ spirv::TargetEnvAttr spirv::lookupTargetEnv(Operation *op) {
if (!op)
break;
- if (auto arrAttr = op->getAttrOfType<ArrayAttr>("targets")) {
- for (auto attr : arrAttr)
- if (auto spirvTargetEnvAttr =
- llvm::dyn_cast<spirv::TargetEnvAttr>(attr))
- return spirvTargetEnvAttr;
- }
-
if (auto attr = op->getAttrOfType<spirv::TargetEnvAttr>(
spirv::getTargetEnvAttrName()))
return attr;
>From 7d43c590b32fcd45d17239ba0d4c763609fd9e4d Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Thu, 10 Jul 2025 11:05:11 +0900
Subject: [PATCH 05/16] Remove unnecessary check in test
---
mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
index 7c74bbad49101..a6a8e36a8642e 100644
--- a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
@@ -1,7 +1,7 @@
// RUN: mlir-opt --spirv-attach-target='caps=Shader exts=SPV_KHR_storage_buffer_storage_class' --convert-gpu-to-spirv %s -o - | FileCheck %s
module attributes {gpu.container_module} {
- // CHECK-LABEL: spirv.module @{{.*}} GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
+ // CHECK-LABEL: spirv.module @{{.*}} GLSL450
gpu.module @kernels {
// CHECK: spirv.func @load_kernel
// CHECK-SAME: %[[ARG:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>})
>From 40b6f07634f3a34fc0bc21bad0fc647fd52b3ab1 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Thu, 10 Jul 2025 11:09:03 +0900
Subject: [PATCH 06/16] Add same lookup target env logic in GPUModuleConversion
---
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp | 4 ++++
1 file changed, 4 insertions(+)
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
index b99ed261ecfa3..a2b86c35ea99d 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
@@ -385,6 +385,10 @@ LogicalResult GPUModuleConversion::matchAndRewrite(
if (auto attr = moduleOp->getAttrOfType<spirv::TargetEnvAttr>(
spirv::getTargetEnvAttrName()))
spvModule->setAttr(spirv::getTargetEnvAttrName(), attr);
+ for (auto targetAttr : moduleOp.getTargetsAttr())
+ if (auto spirvTargetEnvAttr =
+ llvm::dyn_cast<spirv::TargetEnvAttr>(targetAttr))
+ spvModule->setAttr(spirv::getTargetEnvAttrName(), spirvTargetEnvAttr);
rewriter.eraseOp(moduleOp);
return success();
>From 0b68019365472cac6770227499042e67f6b21cd7 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Thu, 10 Jul 2025 17:30:43 +0900
Subject: [PATCH 07/16] fixup! Add same lookup target env logic in
GPUModuleConversion
remove unnecessary namespace qualifiers
---
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp | 3 +--
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp | 5 ++---
2 files changed, 3 insertions(+), 5 deletions(-)
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
index a2b86c35ea99d..94f6d0c21e241 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
@@ -386,8 +386,7 @@ LogicalResult GPUModuleConversion::matchAndRewrite(
spirv::getTargetEnvAttrName()))
spvModule->setAttr(spirv::getTargetEnvAttrName(), attr);
for (auto targetAttr : moduleOp.getTargetsAttr())
- if (auto spirvTargetEnvAttr =
- llvm::dyn_cast<spirv::TargetEnvAttr>(targetAttr))
+ if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
spvModule->setAttr(spirv::getTargetEnvAttrName(), spirvTargetEnvAttr);
rewriter.eraseOp(moduleOp);
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index d182198b82f47..195e5295d171a 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -56,8 +56,7 @@ struct GPUToSPIRVPass final : impl::ConvertGPUToSPIRVBase<GPUToSPIRVPass> {
spirv::TargetEnvAttr
GPUToSPIRVPass::lookupTargetEnvInTargets(gpu::GPUModuleOp moduleOp) {
for (auto &targetAttr : moduleOp.getTargetsAttr())
- if (auto spirvTargetEnvAttr =
- llvm::dyn_cast<spirv::TargetEnvAttr>(targetAttr))
+ if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
return spirvTargetEnvAttr;
return {};
@@ -105,7 +104,7 @@ void GPUToSPIRVPass::runOnOperation() {
// TargetEnv attributes.
for (Operation *gpuModule : gpuModules) {
spirv::TargetEnvAttr targetAttr =
- lookupTargetEnvOrDefault(llvm::cast<gpu::GPUModuleOp>(gpuModule));
+ lookupTargetEnvOrDefault(cast<gpu::GPUModuleOp>(gpuModule));
// Map MemRef memory space to SPIR-V storage class first if requested.
if (mapMemorySpace) {
>From 0833b2a5d707482b9b7f2b53e5a7fdf1b4a5b709 Mon Sep 17 00:00:00 2001
From: Jaeho Kim <oojahooo at gmail.com>
Date: Fri, 11 Jul 2025 09:04:43 +0900
Subject: [PATCH 08/16] Update
mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
Co-authored-by: Jakub Kuderski <kubakuderski at gmail.com>
---
mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
index a6a8e36a8642e..554ace7dcbbc0 100644
--- a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt --spirv-attach-target='caps=Shader exts=SPV_KHR_storage_buffer_storage_class' --convert-gpu-to-spirv %s -o - | FileCheck %s
+// RUN: mlir-opt --spirv-attach-target='caps=Shader exts=SPV_KHR_storage_buffer_storage_class' --convert-gpu-to-spirv %s | FileCheck %s
module attributes {gpu.container_module} {
// CHECK-LABEL: spirv.module @{{.*}} GLSL450
>From a688ec24112510e33a3718ffdb78de3a469dc9e9 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Fri, 11 Jul 2025 09:13:02 +0900
Subject: [PATCH 09/16] Spell out the not obvious types
---
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp | 2 +-
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp | 4 ++--
2 files changed, 3 insertions(+), 3 deletions(-)
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
index 94f6d0c21e241..aa2528af30312 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
@@ -385,7 +385,7 @@ LogicalResult GPUModuleConversion::matchAndRewrite(
if (auto attr = moduleOp->getAttrOfType<spirv::TargetEnvAttr>(
spirv::getTargetEnvAttrName()))
spvModule->setAttr(spirv::getTargetEnvAttrName(), attr);
- for (auto targetAttr : moduleOp.getTargetsAttr())
+ for (const Attribute &targetAttr : moduleOp.getTargetsAttr())
if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
spvModule->setAttr(spirv::getTargetEnvAttrName(), spirvTargetEnvAttr);
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index 195e5295d171a..5c9b960430679 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -55,7 +55,7 @@ struct GPUToSPIRVPass final : impl::ConvertGPUToSPIRVBase<GPUToSPIRVPass> {
spirv::TargetEnvAttr
GPUToSPIRVPass::lookupTargetEnvInTargets(gpu::GPUModuleOp moduleOp) {
- for (auto &targetAttr : moduleOp.getTargetsAttr())
+ for (const Attribute &targetAttr : moduleOp.getTargetsAttr())
if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
return spirvTargetEnvAttr;
@@ -64,7 +64,7 @@ GPUToSPIRVPass::lookupTargetEnvInTargets(gpu::GPUModuleOp moduleOp) {
spirv::TargetEnvAttr
GPUToSPIRVPass::lookupTargetEnvOrDefault(gpu::GPUModuleOp moduleOp) {
- if (auto targetEnvAttr = lookupTargetEnvInTargets(moduleOp))
+ if (spirv::TargetEnvAttr targetEnvAttr = lookupTargetEnvInTargets(moduleOp))
return targetEnvAttr;
return spirv::lookupTargetEnvOrDefault(moduleOp);
>From 1f38539920a9bcda6bcad34ff5932ace578ab5c3 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Fri, 11 Jul 2025 09:13:48 +0900
Subject: [PATCH 10/16] Add documentation comments
---
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp | 6 ++++++
1 file changed, 6 insertions(+)
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index 5c9b960430679..29960791e3a4b 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -48,7 +48,13 @@ struct GPUToSPIRVPass final : impl::ConvertGPUToSPIRVBase<GPUToSPIRVPass> {
void runOnOperation() override;
private:
+ /// Queries the target environment from 'targets' attribute of the given
+ /// `moduleOp`.
spirv::TargetEnvAttr lookupTargetEnvInTargets(gpu::GPUModuleOp moduleOp);
+
+ /// Queries the target environment from 'targets' attribute of the given
+ /// `moduleOp` or returns target environment as returned by
+ /// `spirv::lookupTargetEnvOrDefault` if not provided by 'targets'.
spirv::TargetEnvAttr lookupTargetEnvOrDefault(gpu::GPUModuleOp moduleOp);
bool mapMemorySpace;
};
>From 9344c20e60ad83ec72cae0245f521fd4de26f450 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Fri, 11 Jul 2025 09:14:33 +0900
Subject: [PATCH 11/16] Simplify pass pipeline of test code
---
mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
index 554ace7dcbbc0..8efda39d9c37e 100644
--- a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
@@ -1,8 +1,8 @@
-// RUN: mlir-opt --spirv-attach-target='caps=Shader exts=SPV_KHR_storage_buffer_storage_class' --convert-gpu-to-spirv %s | FileCheck %s
+// RUN: mlir-opt --convert-gpu-to-spirv %s | FileCheck %s
module attributes {gpu.container_module} {
// CHECK-LABEL: spirv.module @{{.*}} GLSL450
- gpu.module @kernels {
+ gpu.module @kernels [#spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>] {
// CHECK: spirv.func @load_kernel
// CHECK-SAME: %[[ARG:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>})
gpu.func @load_kernel(%arg0: memref<12x4xf32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
>From 3f8a37939d3e41b98a77b422a7d6591bbf7b5207 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Wed, 16 Jul 2025 13:37:02 +0900
Subject: [PATCH 12/16] Fix failures in other tests
It had to be checked for a lack of `"targets"` attributes in a GPU
module
---
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp | 8 +++++---
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp | 8 +++++---
2 files changed, 10 insertions(+), 6 deletions(-)
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
index aa2528af30312..ebe15abaf865e 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
@@ -385,9 +385,11 @@ LogicalResult GPUModuleConversion::matchAndRewrite(
if (auto attr = moduleOp->getAttrOfType<spirv::TargetEnvAttr>(
spirv::getTargetEnvAttrName()))
spvModule->setAttr(spirv::getTargetEnvAttrName(), attr);
- for (const Attribute &targetAttr : moduleOp.getTargetsAttr())
- if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
- spvModule->setAttr(spirv::getTargetEnvAttrName(), spirvTargetEnvAttr);
+ if (const ArrayAttr &targets = moduleOp.getTargetsAttr()) {
+ for (const Attribute &targetAttr : targets)
+ if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
+ spvModule->setAttr(spirv::getTargetEnvAttrName(), spirvTargetEnvAttr);
+ }
rewriter.eraseOp(moduleOp);
return success();
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index 29960791e3a4b..1a56cf35533c9 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -61,9 +61,11 @@ struct GPUToSPIRVPass final : impl::ConvertGPUToSPIRVBase<GPUToSPIRVPass> {
spirv::TargetEnvAttr
GPUToSPIRVPass::lookupTargetEnvInTargets(gpu::GPUModuleOp moduleOp) {
- for (const Attribute &targetAttr : moduleOp.getTargetsAttr())
- if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
- return spirvTargetEnvAttr;
+ if (const ArrayAttr &targets = moduleOp.getTargetsAttr()) {
+ for (const Attribute &targetAttr : targets)
+ if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
+ return spirvTargetEnvAttr;
+ }
return {};
}
>From 3703a1e6a90a5950ea80a37951ce275bda427654 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Thu, 17 Jul 2025 09:41:35 +0900
Subject: [PATCH 13/16] Do not use reference, do pass by value when assign
attributes
---
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp | 4 ++--
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp | 4 ++--
2 files changed, 4 insertions(+), 4 deletions(-)
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
index ebe15abaf865e..2bc67e792c778 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
@@ -385,8 +385,8 @@ LogicalResult GPUModuleConversion::matchAndRewrite(
if (auto attr = moduleOp->getAttrOfType<spirv::TargetEnvAttr>(
spirv::getTargetEnvAttrName()))
spvModule->setAttr(spirv::getTargetEnvAttrName(), attr);
- if (const ArrayAttr &targets = moduleOp.getTargetsAttr()) {
- for (const Attribute &targetAttr : targets)
+ if (ArrayAttr targets = moduleOp.getTargetsAttr()) {
+ for (Attribute targetAttr : targets)
if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
spvModule->setAttr(spirv::getTargetEnvAttrName(), spirvTargetEnvAttr);
}
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index 1a56cf35533c9..25fc1fd55f9bb 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -61,8 +61,8 @@ struct GPUToSPIRVPass final : impl::ConvertGPUToSPIRVBase<GPUToSPIRVPass> {
spirv::TargetEnvAttr
GPUToSPIRVPass::lookupTargetEnvInTargets(gpu::GPUModuleOp moduleOp) {
- if (const ArrayAttr &targets = moduleOp.getTargetsAttr()) {
- for (const Attribute &targetAttr : targets)
+ if (ArrayAttr targets = moduleOp.getTargetsAttr()) {
+ for (Attribute targetAttr : targets)
if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
return spirvTargetEnvAttr;
}
>From b8b3c1898ee53164c9f11247ac90d4ad85cbdbad Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Thu, 31 Jul 2025 09:36:00 +0900
Subject: [PATCH 14/16] Select first target env when targets array has more
than one target env
---
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp | 5 ++++-
1 file changed, 4 insertions(+), 1 deletion(-)
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
index 2bc67e792c778..d5a3057f0fe2e 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
@@ -387,8 +387,11 @@ LogicalResult GPUModuleConversion::matchAndRewrite(
spvModule->setAttr(spirv::getTargetEnvAttrName(), attr);
if (ArrayAttr targets = moduleOp.getTargetsAttr()) {
for (Attribute targetAttr : targets)
- if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
+ if (auto spirvTargetEnvAttr =
+ dyn_cast<spirv::TargetEnvAttr>(targetAttr)) {
spvModule->setAttr(spirv::getTargetEnvAttrName(), spirvTargetEnvAttr);
+ break;
+ }
}
rewriter.eraseOp(moduleOp);
>From 1a7d496a1584b21bf9684c55baaf837a2fa64f78 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Thu, 31 Jul 2025 09:36:59 +0900
Subject: [PATCH 15/16] Add test case with more than one target env in targets
array
---
.../GPUToSPIRV/lookup-target-env.mlir | 24 ++++++++++++++++++-
1 file changed, 23 insertions(+), 1 deletion(-)
diff --git a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
index 8efda39d9c37e..2d934d76ac68d 100644
--- a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt --convert-gpu-to-spirv %s | FileCheck %s
+// RUN: mlir-opt --split-input-file --convert-gpu-to-spirv %s | FileCheck %s
module attributes {gpu.container_module} {
// CHECK-LABEL: spirv.module @{{.*}} GLSL450
@@ -15,3 +15,25 @@ module attributes {gpu.container_module} {
}
}
}
+
+// -----
+
+module attributes {gpu.container_module} {
+ // CHECK-LABEL: spirv.module @{{.*}} GLSL450
+ // CHECK-SAME: #spirv.target_env<#spirv.vce<v1.4, [Shader], [SPV_KHR_storage_buffer_storage_class]>
+ gpu.module @kernels [
+ #spirv.target_env<#spirv.vce<v1.4, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>,
+ #spirv.target_env<#spirv.vce<v1.0, [Kernel], []>, #spirv.resource_limits<>>,
+ #spirv.target_env<#spirv.vce<v1.0, [Shader], []>, #spirv.resource_limits<>>] {
+ // CHECK: spirv.func @load_kernel
+ // CHECK-SAME: %[[ARG:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>})
+ gpu.func @load_kernel(%arg0: memref<12x4xf32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
+ %c0 = arith.constant 0 : index
+ // CHECK: %[[PTR:.*]] = spirv.AccessChain %[[ARG]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
+ // CHECK-NEXT: {{%.*}} = spirv.Load "StorageBuffer" %[[PTR]] : f32
+ %0 = memref.load %arg0[%c0, %c0] : memref<12x4xf32>
+ // CHECK: spirv.Return
+ gpu.return
+ }
+ }
+}
>From b1e810b7c56423a65dcb56fcd0f8c7c09177c489 Mon Sep 17 00:00:00 2001
From: oojahooo <oojahooo at gmail.com>
Date: Thu, 31 Jul 2025 12:01:27 +0900
Subject: [PATCH 16/16] Add comment explaining the purpose of new test case
---
mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
index 2d934d76ac68d..983747be57995 100644
--- a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
@@ -17,7 +17,8 @@ module attributes {gpu.container_module} {
}
// -----
-
+// Checks that the `-convert-gpu-to-spirv` pass selects the first
+// `spirv.target_env` from the `targets` array attribute attached to `gpu.module`.
module attributes {gpu.container_module} {
// CHECK-LABEL: spirv.module @{{.*}} GLSL450
// CHECK-SAME: #spirv.target_env<#spirv.vce<v1.4, [Shader], [SPV_KHR_storage_buffer_storage_class]>
More information about the Mlir-commits
mailing list