[Mlir-commits] [mlir] [mlir][spirv]: Add ImageSsupport in ABI Lowering (PR #150996)
Jack Frankland
llvmlistbot at llvm.org
Tue Jul 29 02:42:31 PDT 2025
https://github.com/FranklandJack updated https://github.com/llvm/llvm-project/pull/150996
>From 2a4c34826008a40eb4ab2ba33e6e038e8e784a8e Mon Sep 17 00:00:00 2001
From: Jack Frankland <jack.frankland at arm.com>
Date: Mon, 28 Jul 2025 17:35:48 +0100
Subject: [PATCH 1/3] [mlir][spirv]: Add ImageSsupport in ABI Lowering
Add support for generating shader arguments as global variables in the
SPIR-V module when the argument in question is a SPIR-V image.
Add lit tests to execute the new logic and check global variables are
being generated.
Signed-off-by: Jack Frankland <jack.frankland at arm.com>
---
.../Transforms/LowerABIAttributesPass.cpp | 12 +++++++++-
.../SPIRV/Transforms/abi-interface.mlir | 24 +++++++++++++++++++
2 files changed, 35 insertions(+), 1 deletion(-)
diff --git a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp
index 85525a5a02fa2..e447e4bfae9dc 100644
--- a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp
+++ b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp
@@ -58,7 +58,17 @@ createGlobalVarForEntryPointArgument(OpBuilder &builder, spirv::FuncOp funcOp,
spirv::PointerType::get(spirv::StructType::get(varType), *storageClass);
}
auto varPtrType = cast<spirv::PointerType>(varType);
- auto varPointeeType = cast<spirv::StructType>(varPtrType.getPointeeType());
+ auto pointeeType = varPtrType.getPointeeType();
+
+ // Images are an opaque type and so we can just return a pointer to an image.
+ // Note that currently only sampled images are supported in the SPIR-V
+ // lowering.
+ if (isa<spirv::SampledImageType>(pointeeType))
+ return builder.create<spirv::GlobalVariableOp>(
+ funcOp.getLoc(), varType, varName, abiInfo.getDescriptorSet(),
+ abiInfo.getBinding());
+
+ auto varPointeeType = cast<spirv::StructType>(pointeeType);
// Set the offset information.
varPointeeType =
diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir
index bd51a07843652..f3a3218e5aec0 100644
--- a/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir
+++ b/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir
@@ -66,3 +66,27 @@ spirv.module Logical GLSL450 attributes {spirv.target_env = #spirv.target_env<#s
// CHECK: spirv.EntryPoint "GLCompute" [[FN]], [[VAR0]], [[VAR1]]
// CHECK: spirv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
} // end spirv.module
+
+// -----
+
+module {
+ spirv.module Logical GLSL450 attributes {spirv.target_env = #spirv.target_env<#spirv.vce<v1.6, [Shader, Sampled1D], []>, #spirv.resource_limits<>>} {
+ // CHECK-DAG: spirv.GlobalVariable @[[IMAGE_GV:.*]] bind(0, 0) : !spirv.ptr<!spirv.sampled_image<!spirv.image<f32, Dim1D, DepthUnknown, NonArrayed, SingleSampled, NeedSampler, R32f>>, UniformConstant>
+ // CHECK: spirv.func @read_image
+ spirv.func @read_image(%arg0: !spirv.ptr<!spirv.sampled_image<!spirv.image<f32, Dim1D, DepthUnknown, NonArrayed, SingleSampled, NeedSampler, R32f>>, UniformConstant> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>}, %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<1 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>}) "None" attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
+ // CHECK: %[[IMAGE_ADDR:.*]] = spirv.mlir.addressof @[[IMAGE_GV]] : !spirv.ptr<!spirv.sampled_image<!spirv.image<f32, Dim1D, DepthUnknown, NonArrayed, SingleSampled, NeedSampler, R32f>>, UniformConstant>
+ %cst0_i32 = spirv.Constant 0 : i32
+ // CHECK: spirv.Load "UniformConstant" %[[IMAGE_ADDR]]
+ %0 = spirv.Load "UniformConstant" %arg0 : !spirv.sampled_image<!spirv.image<f32, Dim1D, DepthUnknown, NonArrayed, SingleSampled, NeedSampler, R32f>>
+ %1 = spirv.Image %0 : !spirv.sampled_image<!spirv.image<f32, Dim1D, DepthUnknown, NonArrayed, SingleSampled, NeedSampler, R32f>>
+ %2 = spirv.ImageFetch %1, %cst0_i32 : !spirv.image<f32, Dim1D, DepthUnknown, NonArrayed, SingleSampled, NeedSampler, R32f>, i32 -> vector<4xf32>
+ %3 = spirv.CompositeExtract %2[0 : i32] : vector<4xf32>
+ %cst0_i32_0 = spirv.Constant 0 : i32
+ %cst0_i32_1 = spirv.Constant 0 : i32
+ %cst1_i32 = spirv.Constant 1 : i32
+ %4 = spirv.AccessChain %arg1[%cst0_i32_0, %cst0_i32] : !spirv.ptr<!spirv.struct<(!spirv.array<1 x f32, stride=4> [0])>, StorageBuffer>, i32, i32 -> !spirv.ptr<f32, StorageBuffer>
+ spirv.Store "StorageBuffer" %4, %3 : f32
+ spirv.Return
+ }
+ }
+}
>From fd5c33dfb67110551a06b5a259167f2d7f98f67c Mon Sep 17 00:00:00 2001
From: Jack Frankland <jack.frankland at arm.com>
Date: Tue, 29 Jul 2025 07:38:53 +0100
Subject: [PATCH 2/3] [mlir][spirv]: Address Feedback
Make variable type explicit in definition.
Signed-off-by: Jack Frankland <jack.frankland at arm.com>
---
mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp
index e447e4bfae9dc..6d215f47e2d93 100644
--- a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp
+++ b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp
@@ -58,7 +58,7 @@ createGlobalVarForEntryPointArgument(OpBuilder &builder, spirv::FuncOp funcOp,
spirv::PointerType::get(spirv::StructType::get(varType), *storageClass);
}
auto varPtrType = cast<spirv::PointerType>(varType);
- auto pointeeType = varPtrType.getPointeeType();
+ Type pointeeType = varPtrType.getPointeeType();
// Images are an opaque type and so we can just return a pointer to an image.
// Note that currently only sampled images are supported in the SPIR-V
>From 5d57b0b2412f76303044b235e48f49f5bf8b5257 Mon Sep 17 00:00:00 2001
From: Jack Frankland <jack.frankland at arm.com>
Date: Tue, 29 Jul 2025 10:40:21 +0100
Subject: [PATCH 3/3] [mlir][spirv]: Address Feedback
Use new builder APIs over builder methods.
Signed-off-by: Jack Frankland <jack.frankland at arm.com>
---
.../lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp
index 6d215f47e2d93..316b6bb88a567 100644
--- a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp
+++ b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp
@@ -64,9 +64,9 @@ createGlobalVarForEntryPointArgument(OpBuilder &builder, spirv::FuncOp funcOp,
// Note that currently only sampled images are supported in the SPIR-V
// lowering.
if (isa<spirv::SampledImageType>(pointeeType))
- return builder.create<spirv::GlobalVariableOp>(
- funcOp.getLoc(), varType, varName, abiInfo.getDescriptorSet(),
- abiInfo.getBinding());
+ return spirv::GlobalVariableOp::create(builder, funcOp.getLoc(), varType,
+ varName, abiInfo.getDescriptorSet(),
+ abiInfo.getBinding());
auto varPointeeType = cast<spirv::StructType>(pointeeType);
More information about the Mlir-commits
mailing list