Skip to content

[mlir][spirv] Fix lookup logic spirv.target_env for gpu.module #147262

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 17 commits into from
Aug 1, 2025
Merged

[mlir][spirv] Fix lookup logic spirv.target_env for gpu.module #147262

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
17 commits
Select commit Hold shift + click to select a range
5d05b68
[mlir][SPIRV] Fix lookup logic `spirv.target_env` for `gpu.module`
oojahooo Jul 7, 2025
71d3d92
Add test
oojahooo Jul 8, 2025
ae6ba9c
Delete braces for complying with conding standard
oojahooo Jul 8, 2025
fa3b44e
fixup! [mlir][SPIRV] Fix lookup logic `spirv.target_env` for `gpu.mod…
oojahooo Jul 10, 2025
7d43c59
Remove unnecessary check in test
oojahooo Jul 10, 2025
40b6f07
Add same lookup target env logic in GPUModuleConversion
oojahooo Jul 10, 2025
0b68019
fixup! Add same lookup target env logic in GPUModuleConversion
oojahooo Jul 10, 2025
0833b2a
Update mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
oojahooo Jul 11, 2025
a688ec2
Spell out the not obvious types
oojahooo Jul 11, 2025
1f38539
Add documentation comments
oojahooo Jul 11, 2025
9344c20
Simplify pass pipeline of test code
oojahooo Jul 11, 2025
3f8a379
Fix failures in other tests
oojahooo Jul 16, 2025
3703a1e
Do not use reference, do pass by value when assign attributes
oojahooo Jul 17, 2025
b8b3c18
Select first target env when targets array has more than one target env
oojahooo Jul 31, 2025
1a7d496
Add test case with more than one target env in targets array
oojahooo Jul 31, 2025
b1e810b
Add comment explaining the purpose of new test case
oojahooo Jul 31, 2025
6a56521
Merge branch 'main' into spirv_lookup_target
oojahooo Jul 31, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 4 additions & 1 deletion mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
24 changes: 23 additions & 1 deletion mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand All @@ -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
}
}
}