Fix of usage of lifetime instrinsic on llvm 22

LLVM 22 drops the size parameter and forces lifetime instrinsics to
directly refer to allocas rather through other objects. For this reason
we make the instrinsic usage uniform.

Author: Colin Davidson

modules/compiler/vecz/source/analysis/uniform_value_analysis.cpp

@@ -24,6 +24,7 @@
 #include <llvm/IR/Instructions.h>
 #include <llvm/IR/Module.h>
 #include <llvm/Support/Debug.h>
+#include <multi_llvm/lifetime_helper.h>
 
 #include <cstdlib>
 
@@ -358,6 +359,20 @@ void UniformValueResult::markVaryingValues(Value *V, Value *From) {
           return;
         }
       }
+#if LLVM_VERSION_GREATER_EQUAL(22, 0)
+      // LLVM 22 drops the size parameter and requires the object
+      // to be an alloca or a pointer to an alloca, so we force uniformity.
+      // can drop the LifeTimeHasSizeArg() check once intel/llvm has
+      // the size parameter removed from the intrinsic
+      else if (auto *intrinsic = dyn_cast<llvm::IntrinsicInst>(CI)) {
+        const auto intrinsicID = intrinsic->getIntrinsicID();
+        if (!multi_llvm::LifeTimeHasSizeArg() &&
+            (intrinsicID == llvm::Intrinsic::lifetime_end ||
+             intrinsicID == llvm::Intrinsic::lifetime_start)) {
+          return;
+        }
+      }
+#endif
     }
   }
 

modules/compiler/vecz/source/transform/packetizer.cpp

@@ -38,6 +38,7 @@
 #include <llvm/Support/Debug.h>
 #include <llvm/Support/raw_ostream.h>
 #include <llvm/Transforms/Utils/LoopUtils.h>
+#include <multi_llvm/lifetime_helper.h>
 #include <multi_llvm/llvm_version.h>
 #include <multi_llvm/multi_llvm.h>
 #include <multi_llvm/vector_type_helper.h>
@@ -2050,25 +2051,33 @@ ValuePacket Packetizer::Impl::packetizeCall(CallInst *CI) {
     auto IntrID = Intrinsic::ID(Callee->getIntrinsicID());
     if (IntrID == llvm::Intrinsic::lifetime_end ||
         IntrID == llvm::Intrinsic::lifetime_start) {
-      auto *ptr = CI->getOperand(1);
-      if (auto *const bcast = dyn_cast<BitCastInst>(ptr)) {
-        ptr = bcast->getOperand(0);
-      }
-
-      if (auto *const alloca = dyn_cast<AllocaInst>(ptr)) {
-        if (!needsInstantiation(Ctx, *alloca)) {
-          // If it's an alloca we can widen, we can just change the size
-          const llvm::TypeSize allocSize =
-              Ctx.dataLayout()->getTypeAllocSize(alloca->getAllocatedType());
-          const auto lifeSize =
-              allocSize.isScalable() || SimdWidth.isScalable()
-                  ? -1
-                  : allocSize.getKnownMinValue() * SimdWidth.getKnownMinValue();
-          CI->setOperand(
-              0, ConstantInt::get(CI->getOperand(0)->getType(), lifeSize));
-          results.push_back(CI);
+#if LLVM_VERSION_LESS(22, 0) || 1
+      // LLVM 22 drops the size parameter
+      // Remove LifetimeHasSizeArg once intel llvm no longer has the extra
+      // argument
+      if (multi_llvm::LifeTimeHasSizeArg()) {
+        auto *ptr = CI->getOperand(1);
+        if (auto *const bcast = dyn_cast<BitCastInst>(ptr)) {
+          ptr = bcast->getOperand(0);
+        }
+
+        if (auto *const alloca = dyn_cast<AllocaInst>(ptr)) {
+          if (!needsInstantiation(Ctx, *alloca)) {
+            // If it's an alloca we can widen, we can just change the size
+            const llvm::TypeSize allocSize =
+                Ctx.dataLayout()->getTypeAllocSize(alloca->getAllocatedType());
+            const auto lifeSize =
+                allocSize.isScalable() || SimdWidth.isScalable()
+                    ? -1
+                    : allocSize.getKnownMinValue() *
+                          SimdWidth.getKnownMinValue();
+            CI->setOperand(
+                0, ConstantInt::get(CI->getOperand(0)->getType(), lifeSize));
+            results.push_back(CI);
+          }
         }
       }
+#endif
       return results;
     }
 

modules/compiler/vecz/test/lit/llvm/divergent_loop_bug.ll

@@ -48,7 +48,7 @@ entry.if.end17_crit_edge:                          ; preds = %entry
 ; %or.cond branch.
 ; CHECK: if.then:
 ; CHECK: call void @__vecz_b_masked_store4_fu3ptrb(float 0.000000e+00, ptr %cosa, i1 [[CMP_NOT_NOT]])
-; CHECK: %1 = call spir_func float @__vecz_b_masked__Z6sincosfPf(float 0.000000e+00, ptr nonnull %cosa, i1 [[CMP_NOT_NOT]]) #9
+; CHECK: %1 = call spir_func float @__vecz_b_masked__Z6sincosfPf(float 0.000000e+00, ptr nonnull %cosa, i1 [[CMP_NOT_NOT]])
 ; CHECK: %2 = call float @__vecz_b_masked_load4_fu3ptrb(ptr %cosa, i1 [[CMP_NOT_NOT]])
 ; CHECK: %mul7 = fmul float %2, -2.950000e+01
 ; CHECK: %cmp11 = fcmp uge float %mul7, 0.000000e+00
@@ -60,7 +60,7 @@ entry.if.end17_crit_edge:                          ; preds = %entry
 if.then:                                           ; preds = %entry
   call void @llvm.lifetime.start.p0(i64 4, ptr nonnull %cosa) #6
   store float 0.000000e+00, ptr %cosa, align 4
-  %call4 = call spir_func float @_Z6sincosfPf(float 0.000000e+00, ptr nonnull %cosa) #7
+  %call4 = call spir_func float @_Z6sincosfPf(float 0.000000e+00, ptr nonnull %cosa)
   %1 = load float, ptr %cosa, align 4
   %mul7 = fmul float %1, -2.950000e+01
   %cmp11 = fcmp uge float %mul7, 0.000000e+00
@@ -113,7 +113,7 @@ entry.if.end17_crit_edge:                          ; preds = %entry
 ; %or.cond branch.
 ; CHECK: if.then:
 ; CHECK: call void @__vecz_b_masked_store4_fu3ptrb(float 0.000000e+00, ptr %cosa, i1 [[CMP_NOT_NOT]])
-; CHECK: %1 = call spir_func float @__vecz_b_masked__Z6sincosfPf(float 0.000000e+00, ptr nonnull %cosa, i1 [[CMP_NOT_NOT]]) #9
+; CHECK: %1 = call spir_func float @__vecz_b_masked__Z6sincosfPf(float 0.000000e+00, ptr nonnull %cosa, i1 [[CMP_NOT_NOT]])
 ; CHECK: %2 = call float @__vecz_b_masked_load4_fu3ptrb(ptr %cosa, i1 [[CMP_NOT_NOT]])
 ; CHECK: %mul7 = fmul float %2, -2.950000e+01
 ; CHECK: %cmp11 = fcmp uge float %mul7, 0.000000e+00

modules/compiler/vecz/test/lit/llvm/packetize_lifetime.ll

@@ -0,0 +1,44 @@
+; Copyright (C) Codeplay Software Limited
+;
+; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
+; Exceptions; you may not use this file except in compliance with the License.
+; You may obtain a copy of the License at
+;
+;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
+;
+; Unless required by applicable law or agreed to in writing, software
+; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
+; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
+; License for the specific language governing permissions and limitations
+; under the License.
+;
+; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+; RUN: %pp-llvm-ver -o %t < %s --llvm-ver %LLVMVER
+; RUN: veczc -vecz-passes=packetizer -S < %s | FileCheck %t
+; REQUIRES: llvm-22+
+
+target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target triple = "spir64-unknown-unknown"
+
+declare i64 @__mux_get_global_id(i32)
+
+define spir_kernel void @__vecz_v64_fract_double3() {
+entry:
+  %iout.i = alloca <3 x double>, align 32
+  %call.i = tail call i64 @__mux_get_global_id(i32 0)
+  %cmp.i = icmp ult i64 %call.i, 0
+  call void @llvm.lifetime.start.p0(ptr nonnull %iout.i)
+  %.splatinsert = insertelement <4 x i1> zeroinitializer, i1 %cmp.i, i64 0
+  call void null(<4 x double> zeroinitializer, ptr %iout.i, <4 x i1> %.splatinsert)
+  call void @llvm.lifetime.end.p0(ptr nonnull %iout.i)
+  ret void
+}
+; CHECK: spir_kernel void @__vecz_v4___vecz_v64_fract_double3
+; CHECK: %iout.i1 = alloca <3 x double>, i64 4, align 3
+; CHECK-GE22: call void @llvm.lifetime.start.p0(ptr nonnull %iout.i1)
+; CHECK-GE22: call void @llvm.lifetime.end.p0(ptr nonnull %iout.i1)
+; CHECK-LT22: call void @llvm.lifetime.start.p0(ptr nonnull %iout.i1)
+; CHECK-LT22: call void @llvm.lifetime.end.p0(ptr nonnull %iout.i1)
+; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
+