Fix merging of uniform and non-uniform allocas

Merging between uniform and non-uniform allocas was creating incorrect
address calculation since simdLaneId was being ignored for non-uniform
values.

Author: stefan-il

IGC/AdaptorCommon/LivenessUtils/MergeAllocas.cpp

@@ -50,7 +50,8 @@ static MergeAllocas::AllocaInfo GetAllocaInfo(AllocaInst *allocaI,
             allocationSize,
             static_cast<size_t>(
                 DL->getPrefTypeAlign(allocaI->getAllocatedType()).value()),
-            0};
+            0,
+            allocaI->getMetadata("uniform") != nullptr};
 }
 
 static size_t GetStartingOffset(size_t startOffset, size_t alignment) {
@@ -63,6 +64,9 @@ static size_t GetStartingOffset(size_t startOffset, size_t alignment) {
 
 static bool AddNonOverlappingAlloca(MergeAllocas::AllocaInfo* MergableAlloca,
                                     MergeAllocas::AllocaInfo* NewAlloca) {
+    if (MergableAlloca->isUniform != NewAlloca->isUniform) {
+        return false;
+    }
     if (MergableAlloca->addressSpace != NewAlloca->addressSpace) {
         return false;
     }

IGC/AdaptorCommon/LivenessUtils/MergeAllocas.h

@@ -34,6 +34,7 @@ namespace IGC
             std::size_t alignment;
             // start offset of this alloca in top level alloca (if any)
             std::size_t offset;
+            bool isUniform;
         };
 
         MergeAllocas(char& pid) : AllocationLivenessAnalyzer(pid) {}

IGC/Compiler/tests/PrivateMemoryResolution/MergeAllocas/mergable_allocas_mixed_uniform.ll

@@ -0,0 +1,22 @@
+;=========================== begin_copyright_notice ============================
+;
+; Copyright (C) 2025 Intel Corporation
+;
+; SPDX-License-Identifier: MIT
+;
+;============================ end_copyright_notice =============================
+; RUN: igc_opt --igc-ocl-merge-allocas -S %s --platformpvc | FileCheck %s
+; ------------------------------------------------
+; MergeAllocas
+; ------------------------------------------------
+
+; Check that allocas are not merged if they are not both uniform or non-uniform.
+define spir_kernel void @testFn() {
+; CHECK-LABEL: testFn
+; CHECK-NEXT: alloca [128 x float], i32 0, align 4
+; CHECK-NEXT: alloca [128 x float], i32 0, align 4
+  %1 = alloca [128 x float], i32 0, align 4, !uniform !0
+  %2 = alloca [128 x float], i32 0, align 4
+  ret void
+}
+!0 = !{i1 true}

IGC/Compiler/tests/PrivateMemoryResolution/MergeAllocas/merge_allocas.ll

@@ -24,8 +24,8 @@ declare void @llvm.assume(i1 noundef) #1
 
 declare spir_func void @__itt_offload_wi_finish_wrapper()
 
-define spir_kernel void @_ZTS43Kernel_NoReusePrivMem_SameFunc_AlwaysInline(float addrspace(1)* %0, i64 %1, i64 %2, i32 %3, i32 %4, i32 %5) {
-; CHECK-LABEL: _ZTS43Kernel_NoReusePrivMem_SameFunc_AlwaysInline
+define spir_kernel void @main(float addrspace(1)* %0, i64 %1, i64 %2, i32 %3, i32 %4, i32 %5) {
+; CHECK-LABEL: main
 ; CHECK-NEXT: alloca [128 x float], align 4
 ; CHECK-NOT: alloca [128 x float], align 4
   %7 = alloca [128 x float], align 4

IGC/Compiler/tests/PrivateMemoryResolution/MergeAllocas/merge_allocas_loop.ll

@@ -21,7 +21,7 @@ declare void @llvm.dbg.declare(metadata, metadata, metadata) #0
 declare void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* noalias nocapture writeonly, i8 addrspace(4)* noalias nocapture readonly, i64, i1 immarg) #1
 
 ; Function Attrs: noinline optnone
-define internal spir_func i1 @_ZN12_GLOBAL__N_117check_elems_equalIdLi3EEEbRKN4sycl3_V13vecIT_XT0_EEES7_.28(i1 %0) #2 {
+define internal spir_func i1 @testFn(i1 %0) #2 {
   %2 = alloca %"class.sycl::_V1::vec.73", i32 0, align 32
   %3 = addrspacecast %"class.sycl::_V1::vec.73"* null to %"class.sycl::_V1::vec.73" addrspace(4)*
   br label %4
@@ -45,8 +45,8 @@ define internal spir_func i1 @_ZN12_GLOBAL__N_117check_elems_equalIdLi3EEEbRKN4s
   ret i1 false
 }
 
-define spir_kernel void @_ZTSN16accessor_utility34buffer_accessor_get_pointer_kernelIN25accessor_api_local_fp64__11kernel_nameIN4sycl3_V13vecIdLi3EEEEELi0ELNS4_6access4modeE1026ELNS8_6targetE2016ELNS8_11placeholderE0EEE() {
-  %1 = call spir_func i1 @_ZN12_GLOBAL__N_117check_elems_equalIdLi3EEEbRKN4sycl3_V13vecIT_XT0_EEES7_.28(i1 undef)
+define spir_kernel void @main() {
+  %1 = call spir_func i1 @testFn(i1 undef)
   ret void
 }
 

IGC/Compiler/tests/PrivateMemoryResolution/MergeAllocas/merge_allocas_loop2.ll

@@ -19,7 +19,7 @@ target triple = "spir64-unknown-unknown"
 %"class.sycl::_V1::vec.73" = type { <3 x double> }
 
 ; Function Attrs: noinline optnone
-define internal spir_func i1 @_ZN12_GLOBAL__N_117check_elems_equalIdLi3EEEbRKN4sycl3_V13vecIT_XT0_EEES7_.28() #0 {
+define internal spir_func i1 @testFn() #0 {
   %1 = alloca %"class.sycl::_V1::vec.73", i32 0, align 32
   br label %2
 
@@ -39,8 +39,8 @@ define internal spir_func i1 @_ZN12_GLOBAL__N_117check_elems_equalIdLi3EEEbRKN4s
   ret i1 false
 }
 
-define spir_kernel void @_ZTSN16accessor_utility34buffer_accessor_get_pointer_kernelIN25accessor_api_local_fp64__11kernel_nameIN4sycl3_V13vecIdLi3EEEEELi0ELNS4_6access4modeE1026ELNS8_6targetE2016ELNS8_11placeholderE0EEE() {
-  %1 = call spir_func i1 @_ZN12_GLOBAL__N_117check_elems_equalIdLi3EEEbRKN4sycl3_V13vecIT_XT0_EEES7_.28()
+define spir_kernel void @main() {
+  %1 = call spir_func i1 @testFn()
   ret void
 }
 

IGC/Compiler/tests/PrivateMemoryResolution/MergeAllocas/merge_allocas_optnone.ll

@@ -12,8 +12,8 @@
 
 ; Check that allocas are not merged if optnone attribute is set
 ; Function Attrs: noinline optnone
-define spir_kernel void @_ZTS43Kernel_NoReusePrivMem_SameFunc_AlwaysInline() #0 {
-; CHECK-LABEL: _ZTS43Kernel_NoReusePrivMem_SameFunc_AlwaysInline
+define spir_kernel void @testFn() #0 {
+; CHECK-LABEL: testFn
 ; CHECK-NEXT: alloca [128 x float], i32 0, align 4
 ; CHECK-NEXT: alloca [128 x float], i32 0, align 4
   %1 = alloca [128 x float], i32 0, align 4