[Mlir-commits] [mlir] [MLIR][GPUToLLVMSPV] Fix subgroup ops mangling (PR #126111)

[llvmlistbot at llvm.org]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir-gpu

Author: Artem Kroviakov (akroviakov)

<details>
<summary>Changes</summary>

I try to use LLVM's SPIRV backend to compile a GPU module into a binary string that is later consumed by OpenCL to create a program like this:
```cpp
  const char *build_flags = "-cl-kernel-arg-info -x spir -cl-std=CL3.0";
  err = clBuildProgram(program, 1, &device, build_flags, NULL, NULL);
```
The current mangling does not seem to work:
```
unresolved external symbol _Z22get_sub_group_local_id
```
The proposed fix (add parameter list to mangling) resolves the issue.  Am I doing something wrong or why does the current mangled names not work for me?
_______________
Related question:
The similar call in OpenCL:
```cpp
__kernel void kernel(__global int *data) { 
*data = get_sub_group_local_id();
};
```
leads to the following call in IGC dump:
```mlir
call spir_func i32 @<!-- -->_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #<!-- -->0
```
Why doesn't llvm-spv pass lower to spirv built-ins (as the pass name suggests) and uses OpenCL instead?



---
Full diff: https://github.com/llvm/llvm-project/pull/126111.diff


2 Files Affected:

- (modified) mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp (+4-4) 
- (modified) mlir/test/Conversion/GPUToLLVMSPV/gpu-to-llvm-spv.mlir (+8-8) 


``````````diff
diff --git a/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp b/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
index 8b6b553f6eed054..26bfa97fa34d058 100644
--- a/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
+++ b/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
@@ -401,13 +401,13 @@ struct GPUSubgroupOpConversion final : ConvertOpToLLVMPattern<SubgroupOp> {
                   ConversionPatternRewriter &rewriter) const final {
     constexpr StringRef funcName = [] {
       if constexpr (std::is_same_v<SubgroupOp, gpu::SubgroupIdOp>) {
-        return "_Z16get_sub_group_id";
+        return "_Z16get_sub_group_idv";
       } else if constexpr (std::is_same_v<SubgroupOp, gpu::LaneIdOp>) {
-        return "_Z22get_sub_group_local_id";
+        return "_Z22get_sub_group_local_idv";
       } else if constexpr (std::is_same_v<SubgroupOp, gpu::NumSubgroupsOp>) {
-        return "_Z18get_num_sub_groups";
+        return "_Z18get_num_sub_groupsv";
       } else if constexpr (std::is_same_v<SubgroupOp, gpu::SubgroupSizeOp>) {
-        return "_Z18get_sub_group_size";
+        return "_Z18get_sub_group_sizev";
       }
     }();
 
diff --git a/mlir/test/Conversion/GPUToLLVMSPV/gpu-to-llvm-spv.mlir b/mlir/test/Conversion/GPUToLLVMSPV/gpu-to-llvm-spv.mlir
index c2930971dbcf9b9..b3a9c33eb66d911 100644
--- a/mlir/test/Conversion/GPUToLLVMSPV/gpu-to-llvm-spv.mlir
+++ b/mlir/test/Conversion/GPUToLLVMSPV/gpu-to-llvm-spv.mlir
@@ -546,28 +546,28 @@ gpu.module @kernels {
 
 // Lowering of subgroup query operations
 
-// CHECK-DAG: llvm.func spir_funccc @_Z18get_sub_group_size() -> i32 attributes {no_unwind, will_return}
-// CHECK-DAG: llvm.func spir_funccc @_Z18get_num_sub_groups() -> i32 attributes {no_unwind, will_return}
-// CHECK-DAG: llvm.func spir_funccc @_Z22get_sub_group_local_id() -> i32 attributes {no_unwind, will_return}
-// CHECK-DAG: llvm.func spir_funccc @_Z16get_sub_group_id() -> i32 attributes {no_unwind, will_return}
+// CHECK-DAG: llvm.func spir_funccc @_Z18get_sub_group_sizev() -> i32 attributes {no_unwind, will_return}
+// CHECK-DAG: llvm.func spir_funccc @_Z18get_num_sub_groupsv() -> i32 attributes {no_unwind, will_return}
+// CHECK-DAG: llvm.func spir_funccc @_Z22get_sub_group_local_idv() -> i32 attributes {no_unwind, will_return}
+// CHECK-DAG: llvm.func spir_funccc @_Z16get_sub_group_idv() -> i32 attributes {no_unwind, will_return}
 
 
 gpu.module @subgroup_operations {
 // CHECK-LABEL: @gpu_subgroup
   func.func @gpu_subgroup() {
-    // CHECK:       %[[SG_ID:.*]] = llvm.call spir_funccc @_Z16get_sub_group_id() {no_unwind, will_return} : () -> i32
+    // CHECK:       %[[SG_ID:.*]] = llvm.call spir_funccc @_Z16get_sub_group_idv() {no_unwind, will_return} : () -> i32
     // CHECK-32-NOT:                llvm.zext
     // CHECK-64           %{{.*}} = llvm.zext %[[SG_ID]] : i32 to i64
     %0 = gpu.subgroup_id : index
-    // CHECK: %[[SG_LOCAL_ID:.*]] = llvm.call spir_funccc @_Z22get_sub_group_local_id() {no_unwind, will_return}  : () -> i32
+    // CHECK: %[[SG_LOCAL_ID:.*]] = llvm.call spir_funccc @_Z22get_sub_group_local_idv() {no_unwind, will_return}  : () -> i32
     // CHECK-32-NOT:                llvm.zext
     // CHECK-64:          %{{.*}} = llvm.zext %[[SG_LOCAL_ID]] : i32 to i64
     %1 = gpu.lane_id
-    // CHECK:     %[[NUM_SGS:.*]] = llvm.call spir_funccc @_Z18get_num_sub_groups() {no_unwind, will_return} : () -> i32
+    // CHECK:     %[[NUM_SGS:.*]] = llvm.call spir_funccc @_Z18get_num_sub_groupsv() {no_unwind, will_return} : () -> i32
     // CHECK-32-NOT:                llvm.zext
     // CHECK-64:          %{{.*}} = llvm.zext %[[NUM_SGS]] : i32 to i64
     %2 = gpu.num_subgroups : index
-    // CHECK:     %[[SG_SIZE:.*]] = llvm.call spir_funccc @_Z18get_sub_group_size() {no_unwind, will_return} : () -> i32
+    // CHECK:     %[[SG_SIZE:.*]] = llvm.call spir_funccc @_Z18get_sub_group_sizev() {no_unwind, will_return} : () -> i32
     // CHECK-32-NOT:                llvm.zext
     // CHECK-64:          %{{.*}} = llvm.zext %[[SG_SIZE]] : i32 to i64
     %3 = gpu.subgroup_size : index

``````````

</details>


https://github.com/llvm/llvm-project/pull/126111