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