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

; 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 -verify-machineinstrs -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>)