[Mlir-commits] [mlir] 014f4e9 - [mlir][SPIR-V] Add lowering for gpu.lane_id op (#90873)

[llvmlistbot at llvm.org]

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
+    }
+  }
+}