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

Open
wants to merge 7 commits into
base: main
Choose a base branch
from
+42 −4
Open

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 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
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
3 changes: 3 additions & 0 deletions mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -385,6 +385,9 @@ LogicalResult GPUModuleConversion::matchAndRewrite(
if (auto attr = moduleOp->getAttrOfType<spirv::TargetEnvAttr>(
spirv::getTargetEnvAttrName()))
spvModule->setAttr(spirv::getTargetEnvAttrName(), attr);
for (auto targetAttr : moduleOp.getTargetsAttr())
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Use the actual type here, it's not obvious outside of an IDE: https://llvm.org/docs/CodingStandards.html#use-auto-type-deduction-to-make-code-more-readable

if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
spvModule->setAttr(spirv::getTargetEnvAttrName(), spirvTargetEnvAttr);

rewriter.eraseOp(moduleOp);
return success();
Expand Down
26 changes: 22 additions & 4 deletions mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,19 +48,37 @@ struct GPUToSPIRVPass final : impl::ConvertGPUToSPIRVBase<GPUToSPIRVPass> {
void runOnOperation() override;

private:
spirv::TargetEnvAttr lookupTargetEnvInTargets(gpu::GPUModuleOp moduleOp);
spirv::TargetEnvAttr lookupTargetEnvOrDefault(gpu::GPUModuleOp moduleOp);
Comment on lines +51 to +52
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can you add documentation comments explaining what these do?

bool mapMemorySpace;
};

spirv::TargetEnvAttr
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is there a reason to have those two functions? Why not to call spirv::lookupTargetEnvOrDefault directly in GPUToSPIRVPass::lookupTargetEnvInTargets?

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you for your comment.

Since I intentionally did not include OrDefault in the function name GPUToSPIRVPass::lookupTargetEnvInTargets, I wanted the function to only attempt to retrieve the target environment from the "targets" attribute, reflecting that limited scope directly in the name.

Then, I defined a separate GPUToSPIRVPass::lookupTargetEnvOrDefault function that calls it and falls back to spirv::lookupTargetEnvOrDefault if needed, making the fallback behavior more explicit.

However, if adding multiple new functions is considered undesirable, I'm happy to simplify the implementation by having only GPUToSPIRVPass::lookupTargetEnvOrDefault and letting it handle whole condition.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

However, if adding multiple new functions is considered undesirable, I'm happy to simplify the implementation by having only GPUToSPIRVPass::lookupTargetEnvOrDefault and letting it handle whole condition.

I don't think it's undesirable. I just felt like removing the extra indirection may offer a nicer solution, but you made some good argument about why it makes sense the way it is. Initially I thought that you could add OrDefault in the name of the first function, but now I think it may be more confusing. Let's keep it the way it is, unless someone else has a strong opinion about it.

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you for the thoughtful suggestion and for being open to my perspective!

GPUToSPIRVPass::lookupTargetEnvInTargets(gpu::GPUModuleOp moduleOp) {
for (auto &targetAttr : moduleOp.getTargetsAttr())
Copy link
Member

@kuhar kuhar Jul 10, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Spell out the type here since it's not obvious without an IDE/lsp: https://llvm.org/docs/CodingStandards.html#use-auto-type-deduction-to-make-code-more-readable

if (auto spirvTargetEnvAttr = dyn_cast<spirv::TargetEnvAttr>(targetAttr))
return spirvTargetEnvAttr;

return {};
}

spirv::TargetEnvAttr
GPUToSPIRVPass::lookupTargetEnvOrDefault(gpu::GPUModuleOp moduleOp) {
if (auto targetEnvAttr = lookupTargetEnvInTargets(moduleOp))
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

also here

return targetEnvAttr;

return spirv::lookupTargetEnvOrDefault(moduleOp);
}

void GPUToSPIRVPass::runOnOperation() {
MLIRContext *context = &getContext();
ModuleOp module = getOperation();

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);
};
Expand All @@ -86,7 +104,7 @@ void GPUToSPIRVPass::runOnOperation() {
// TargetEnv attributes.
for (Operation *gpuModule : gpuModules) {
spirv::TargetEnvAttr targetAttr =
spirv::lookupTargetEnvOrDefault(gpuModule);
lookupTargetEnvOrDefault(cast<gpu::GPUModuleOp>(gpuModule));

// Map MemRef memory space to SPIR-V storage class first if requested.
if (mapMemorySpace) {
Expand Down
17 changes: 17 additions & 0 deletions mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
// 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

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What do you think about running convert-gpu-to-spirv only? In the current form, it's not very clear where the 'targets' attribute is materialized.


module attributes {gpu.container_module} {
// 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)>})
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
}
}
}