diff --git a/llvm/lib/SYCLLowerIR/SpecConstants.cpp b/llvm/lib/SYCLLowerIR/SpecConstants.cpp index ab31be8ec831a..c00fe479ad365 100644 --- a/llvm/lib/SYCLLowerIR/SpecConstants.cpp +++ b/llvm/lib/SYCLLowerIR/SpecConstants.cpp @@ -29,6 +29,21 @@ using namespace llvm; +static cl::opt SpecConstantMode( + "spec-constant-mode", cl::Optional, cl::Hidden, + cl::desc("Specialization constant handling mode"), + cl::init(SpecConstantsPass::HandlingMode::emulation), + cl::values( + clEnumValN( + SpecConstantsPass::HandlingMode::default_values, "default_values", + "Specialization constant uses are replaced by default values"), + clEnumValN( + SpecConstantsPass::HandlingMode::emulation, "emulation", + "Specialization constant intrinsic is replaced by run-time buffer"), + clEnumValN(SpecConstantsPass::HandlingMode::native, "native", + "Specialization constant intrinsic is lowered to SPIR-V " + "intrinsic"))); + namespace { // __sycl* intrinsic names are Itanium ABI-mangled; this is common prefix for @@ -827,6 +842,9 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, MapVector SCMetadata; SmallVector DefaultsMetadata; + if (SpecConstantMode.getNumOccurrences() > 0) + Mode = SpecConstantMode; + // Iterate through all declarations of instances of function template // template T __sycl_get*SpecConstantValue(const char *ID) // intrinsic to find its calls and lower them depending on the HandlingMode. diff --git a/llvm/test/SYCLLowerIR/SpecConstants/SYCL-alloca.ll b/llvm/test/SYCLLowerIR/SpecConstants/SYCL-alloca.ll new file mode 100644 index 0000000000000..f1442c1be2154 --- /dev/null +++ b/llvm/test/SYCLLowerIR/SpecConstants/SYCL-alloca.ll @@ -0,0 +1,31 @@ +; RUN: opt -passes=spec-constants -spec-constant-mode=default_values %s -S -o - | FileCheck %s + +; This test checks that SpecConstantsPass is able to correctly transform +; SYCL alloca intrinsics in SPIR-V devices when using default values. + +%"class.sycl::_V1::specialization_id" = type { i64 } +%"class.sycl::_V1::specialization_id.0" = type { i32 } +%"class.sycl::_V1::specialization_id.1" = type { i16 } +%my_range = type { ptr addrspace(4), ptr addrspace(4) } + +@size_i64 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { i64 10 }, align 8 +@size_i32 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.0" { i32 120 }, align 4 +@size_i16 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.1" { i16 1 }, align 2 + +@size_i64_stable_name = private unnamed_addr constant [36 x i8] c"_ZTS14name_generatorIL_Z8size_i64EE\00", align 1 +@size_i32_stable_name = private unnamed_addr constant [36 x i8] c"_ZTS14name_generatorIL_Z8size_i32EE\00", align 1 +@size_i16_stable_name = private unnamed_addr constant [36 x i8] c"_ZTS14name_generatorIL_Z8size_i16EE\00", align 1 + +define spir_kernel void @private_alloca() { +; CHECK: alloca double, i32 120, align 8 + call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64(ptr addrspace(4) addrspacecast (ptr @size_i32_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i32 to ptr addrspace(4)), ptr addrspace(4) null, double 0.000000e+00, i64 8) +; CHECK-NEXT: alloca float, i64 10, align 8 + call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f32(ptr addrspace(4) addrspacecast (ptr @size_i64_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i64 to ptr addrspace(4)), ptr addrspace(4) null, float 0.000000e+00, i64 8) +; CHECK-NEXT: alloca %my_range, i16 1, align 64 + call ptr @llvm.sycl.alloca.p0.p4.p4.p4.s_my_range(ptr addrspace(4) addrspacecast (ptr @size_i16_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i16 to ptr addrspace(4)), ptr addrspace(4) null, %my_range zeroinitializer, i64 64) + ret void +} + +declare ptr @llvm.sycl.alloca.p0.p4.p4.p4.f32(ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), float, i64) +declare ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64(ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), double, i64) +declare ptr @llvm.sycl.alloca.p0.p4.p4.p4.s_my_range(ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), %my_range, i64)