[SPIR-V] Fix types of internal intrinsic functions and add a test case for __builtin_alloca()

llvm/docs/SPIRVUsage.rst

@@ -298,7 +298,7 @@ SPIR-V backend, along with their descriptions and argument details.
      - `[Type, Type, Any Integer]`
      - Inserts an element into an aggregate type at a specified index. Allows for building and modifying arrays and vectors.
    * - `int_spv_const_composite`
-     - 32-bit Integer
+     - Type
      - `[Vararg]`
      - Constructs a composite type from given elements. Key for creating arrays, structs, and vectors from individual components.
    * - `int_spv_bitcast`

llvm/include/llvm/IR/IntrinsicsSPIRV.td

@@ -27,7 +27,7 @@ let TargetPrefix = "spv" in {
   def int_spv_insertv : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_any_ty, llvm_vararg_ty]>;
   def int_spv_extractelt : Intrinsic<[llvm_any_ty], [llvm_any_ty, llvm_anyint_ty]>;
   def int_spv_insertelt : Intrinsic<[llvm_any_ty], [llvm_any_ty, llvm_any_ty, llvm_anyint_ty]>;
-  def int_spv_const_composite : Intrinsic<[llvm_i32_ty], [llvm_vararg_ty]>;
+  def int_spv_const_composite : Intrinsic<[llvm_any_ty], [llvm_vararg_ty]>;
   def int_spv_bitcast : Intrinsic<[llvm_any_ty], [llvm_any_ty]>;
   def int_spv_ptrcast : Intrinsic<[llvm_any_ty], [llvm_any_ty, llvm_metadata_ty, llvm_i32_ty], [ImmArg<ArgIndex<2>>]>;
   def int_spv_switch : Intrinsic<[], [llvm_any_ty, llvm_vararg_ty]>;

llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp

@@ -167,8 +167,9 @@ static bool isMemInstrToReplace(Instruction *I) {
          isa<ExtractValueInst>(I) || isa<AtomicCmpXchgInst>(I);
 }
 
-static bool isAggrToReplace(const Value *V) {
-  return isa<ConstantAggregate>(V) || isa<ConstantDataArray>(V) ||
+static bool isAggrConstForceInt32(const Value *V) {
+  return isa<ConstantArray>(V) || isa<ConstantStruct>(V) ||
+         isa<ConstantDataArray>(V) ||
          (isa<ConstantAggregateZero>(V) && !V->getType()->isVectorTy());
 }
 
@@ -576,36 +577,42 @@ void SPIRVEmitIntrinsics::preprocessCompositeConstants(IRBuilder<> &B) {
     assert(I);
     bool KeepInst = false;
     for (const auto &Op : I->operands()) {
-      auto BuildCompositeIntrinsic =
-          [](Constant *AggrC, ArrayRef<Value *> Args, Value *Op, Instruction *I,
-             IRBuilder<> &B, std::queue<Instruction *> &Worklist,
-             bool &KeepInst, SPIRVEmitIntrinsics &SEI) {
-            B.SetInsertPoint(I);
-            auto *CCI =
-                B.CreateIntrinsic(Intrinsic::spv_const_composite, {}, {Args});
-            Worklist.push(CCI);
-            I->replaceUsesOfWith(Op, CCI);
-            KeepInst = true;
-            SEI.AggrConsts[CCI] = AggrC;
-            SEI.AggrConstTypes[CCI] = SEI.deduceNestedTypeHelper(AggrC);
-          };
-
-      if (auto *AggrC = dyn_cast<ConstantAggregate>(Op)) {
-        SmallVector<Value *> Args(AggrC->op_begin(), AggrC->op_end());
-        BuildCompositeIntrinsic(AggrC, Args, Op, I, B, Worklist, KeepInst,
-                                *this);
-      } else if (auto *AggrC = dyn_cast<ConstantDataArray>(Op)) {
+      Constant *AggrConst = nullptr;
+      Type *ResTy = nullptr;
+      if (auto *COp = dyn_cast<ConstantVector>(Op)) {
+        AggrConst = cast<Constant>(COp);
+        ResTy = COp->getType();
+      } else if (auto *COp = dyn_cast<ConstantArray>(Op)) {
+        AggrConst = cast<Constant>(COp);
+        ResTy = B.getInt32Ty();
+      } else if (auto *COp = dyn_cast<ConstantStruct>(Op)) {
+        AggrConst = cast<Constant>(COp);
+        ResTy = B.getInt32Ty();
+      } else if (auto *COp = dyn_cast<ConstantDataArray>(Op)) {
+        AggrConst = cast<Constant>(COp);
+        ResTy = B.getInt32Ty();
+      } else if (auto *COp = dyn_cast<ConstantAggregateZero>(Op)) {
+        if (!Op->getType()->isVectorTy()) {
+          AggrConst = cast<Constant>(COp);
+          ResTy = B.getInt32Ty();
+        }
+      }
+      if (AggrConst) {
         SmallVector<Value *> Args;
-        for (unsigned i = 0; i < AggrC->getNumElements(); ++i)
-          Args.push_back(AggrC->getElementAsConstant(i));
-        BuildCompositeIntrinsic(AggrC, Args, Op, I, B, Worklist, KeepInst,
-                                *this);
-      } else if (isa<ConstantAggregateZero>(Op) &&
-                 !Op->getType()->isVectorTy()) {
-        auto *AggrC = cast<ConstantAggregateZero>(Op);
-        SmallVector<Value *> Args(AggrC->op_begin(), AggrC->op_end());
-        BuildCompositeIntrinsic(AggrC, Args, Op, I, B, Worklist, KeepInst,
-                                *this);
+        if (auto *COp = dyn_cast<ConstantDataSequential>(Op))
+          for (unsigned i = 0; i < COp->getNumElements(); ++i)
+            Args.push_back(COp->getElementAsConstant(i));
+        else
+          for (auto &COp : AggrConst->operands())
+            Args.push_back(COp);
+        B.SetInsertPoint(I);
+        auto *CI =
+            B.CreateIntrinsic(Intrinsic::spv_const_composite, {ResTy}, {Args});
+        Worklist.push(CI);
+        I->replaceUsesOfWith(Op, CI);
+        KeepInst = true;
+        AggrConsts[CI] = AggrConst;
+        AggrConstTypes[CI] = deduceNestedTypeHelper(AggrConst);
       }
     }
     if (!KeepInst)
@@ -1054,8 +1061,8 @@ void SPIRVEmitIntrinsics::processGlobalValue(GlobalVariable &GV,
     // by llvm IR general logic.
     deduceElementTypeHelper(&GV);
     Constant *Init = GV.getInitializer();
-    Type *Ty = isAggrToReplace(Init) ? B.getInt32Ty() : Init->getType();
-    Constant *Const = isAggrToReplace(Init) ? B.getInt32(1) : Init;
+    Type *Ty = isAggrConstForceInt32(Init) ? B.getInt32Ty() : Init->getType();
+    Constant *Const = isAggrConstForceInt32(Init) ? B.getInt32(1) : Init;
     auto *InitInst = B.CreateIntrinsic(Intrinsic::spv_init_global,
                                        {GV.getType(), Ty}, {&GV, Const});
     InitInst->setArgOperand(1, Init);
@@ -1132,11 +1139,11 @@ void SPIRVEmitIntrinsics::processInstrAfterVisit(Instruction *I,
   if (II && II->getIntrinsicID() == Intrinsic::spv_const_composite &&
       TrackConstants) {
     B.SetInsertPoint(I->getNextNode());
-    Type *Ty = B.getInt32Ty();
     auto t = AggrConsts.find(I);
     assert(t != AggrConsts.end());
-    auto *NewOp = buildIntrWithMD(Intrinsic::spv_track_constant, {Ty, Ty},
-                                  t->second, I, {}, B);
+    auto *NewOp =
+        buildIntrWithMD(Intrinsic::spv_track_constant,
+                        {II->getType(), II->getType()}, t->second, I, {}, B);
     I->replaceAllUsesWith(NewOp);
     NewOp->setArgOperand(0, I);
   }

llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/builtin_alloca.ll

@@ -0,0 +1,48 @@
+; The goal of the test is to:
+; 1) check that composite constants of ConstantVector type preserve their
+;    type and can be successfully used further in LLVM intrinsic functions;
+; 2) demonstrate that a call to __builtin_alloca() maps to instructions
+;    from SPV_INTEL_variable_length_array when this extension is available.
+
+; Test LLVM IR is an artificial example, but it's similar to what can be
+; generated by DPC++ compiler from the code snippet:
+;   ...
+;   size_t Sz = ...;
+;   queue Q;
+;   Q.submit([&](sycl::handler &CGH) {
+;     ...
+;     CGH.single_task([=](sycl::kernel_handler KH) SYCL_ESIMD_KERNEL {
+;       int *PrivateArray = (int *)__builtin_alloca(sizeof(int) * Sz);
+;       ...
+;       simd<int, 8> InitVec(100, 10);
+;       InitVec.copy_to(PrivateArray);
+;       ...
+;     });
+;   }).wait();
+;   ...
+
+; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown %s -o %t.spvt 2>&1 | FileCheck %s --check-prefix=CHECK-ERROR
+
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_variable_length_array %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_variable_length_array %s -o - -filetype=obj | spirv-val %}
+
+; CHECK-ERROR: LLVM ERROR: array allocation: this instruction requires the following SPIR-V extension: SPV_INTEL_variable_length_array
+
+; CHECK-SPIRV: Capability VariableLengthArrayINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_variable_length_array"
+; CHECK-SPIRV: OpVariableLengthArrayINTEL %[[#]] %[[#]]
+
+define spir_kernel void @foo(i64 %_arg_sz) {
+entry:
+  %sz = shl i64 %_arg_sz, 2
+  %p1 = alloca i8, i64 %sz, align 8
+  %p4 = addrspacecast ptr %p1 to ptr addrspace(4)
+  %i = ptrtoint ptr addrspace(4) %p4 to i64
+  %splat_ins = insertelement <8 x i64> poison, i64 %i, i64 0
+  %splat_v = shufflevector <8 x i64> %splat_ins, <8 x i64> poison, <8 x i32> zeroinitializer
+  %sum_r = add <8 x i64> %splat_v, <i64 0, i64 4, i64 8, i64 12, i64 16, i64 20, i64 24, i64 28>
+  call void @llvm.genx.svm.scatter.v8i1.v8i64.v8i32(<8 x i1> <i1 true, i1 true, i1 true, i1 true, i1 true, i1 true, i1 true, i1 true>, i32 0, <8 x i64> %sum_r, <8 x i32> <i32 100, i32 110, i32 120, i32 130, i32 140, i32 150, i32 160, i32 170>)
+  ret void
+}
+
+declare void @llvm.genx.svm.scatter.v8i1.v8i64.v8i32(<8 x i1>, i32, <8 x i64>, <8 x i32>)