[Mlir-commits] [mlir] 014f4e9 - [mlir][SPIR-V] Add lowering for gpu.lane_id op (#90873)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Sun May 25 08:41:37 PDT 2025
Author: Sang Ik Lee
Date: 2025-05-25T08:41:34-07:00
New Revision: 014f4e95e042e91f3219d084c8af4b13adc30f27
URL: https://github.com/llvm/llvm-project/commit/014f4e95e042e91f3219d084c8af4b13adc30f27
DIFF: https://github.com/llvm/llvm-project/commit/014f4e95e042e91f3219d084c8af4b13adc30f27.diff
LOG: [mlir][SPIR-V] Add lowering for gpu.lane_id op (#90873)
Add gpu.lane_id op lower for convert-gpu-to-spirv pass
Added:
Modified:
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp
mlir/test/Conversion/GPUToSPIRV/builtins-opencl.mlir
Removed:
################################################################################
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
index 3cc64b82950b5..78e6ebb523a46 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
@@ -747,6 +747,8 @@ void mlir::populateGPUToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
spirv::BuiltIn::NumSubgroups>,
SingleDimLaunchConfigConversion<gpu::SubgroupSizeOp,
spirv::BuiltIn::SubgroupSize>,
+ SingleDimLaunchConfigConversion<
+ gpu::LaneIdOp, spirv::BuiltIn::SubgroupLocalInvocationId>,
WorkGroupSizeConversion, GPUAllReduceConversion,
GPUSubgroupReduceConversion, GPUPrintfConversion>(typeConverter,
patterns.getContext());
diff --git a/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp b/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp
index 811f03abb3461..62a24646d0662 100644
--- a/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp
+++ b/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp
@@ -778,7 +778,8 @@ getOrInsertBuiltinVariable(Block &body, Location loc, spirv::BuiltIn builtin,
}
case spirv::BuiltIn::SubgroupId:
case spirv::BuiltIn::NumSubgroups:
- case spirv::BuiltIn::SubgroupSize: {
+ case spirv::BuiltIn::SubgroupSize:
+ case spirv::BuiltIn::SubgroupLocalInvocationId: {
auto ptrType =
spirv::PointerType::get(integerType, spirv::StorageClass::Input);
std::string name = getBuiltinVarName(builtin, prefix, suffix);
diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins-opencl.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins-opencl.mlir
index 8990d066e4e27..d4fe618b9df29 100644
--- a/mlir/test/Conversion/GPUToSPIRV/builtins-opencl.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/builtins-opencl.mlir
@@ -50,3 +50,25 @@ module attributes {
}
}
}
+
+// -----
+
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Addresses, Kernel, Int64], []>, #spirv.resource_limits<>>
+} {
+ // INDEX32-LABEL: spirv.module @{{.*}} Physical32 OpenCL
+ // INDEX32: spirv.GlobalVariable [[LANEID:@.*]] built_in("SubgroupLocalInvocationId") : !spirv.ptr<i32, Input>
+ // INDEX64-LABEL: spirv.module @{{.*}} Physical64 OpenCL
+ // INDEX64: spirv.GlobalVariable [[LANEID:@.*]] built_in("SubgroupLocalInvocationId") : !spirv.ptr<i32, Input>
+ gpu.module @kernels {
+ gpu.func @builtin_laneid() kernel
+ attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[LANEID]]
+ // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
+ // INDEX64: spirv.UConvert %{{.+}} : i32 to i64
+ %0 = gpu.lane_id
+ gpu.return
+ }
+ }
+}
More information about the Mlir-commits
mailing list