[NFC][SYCL][SYCLLowerIR] Add cl::opt SpecConstantMode for LIT testing

This allows testing SpecConstantsPass with different modes.

Author: wenju-he

llvm/lib/SYCLLowerIR/SpecConstants.cpp

@@ -29,6 +29,19 @@
 
 using namespace llvm;
 
+static cl::opt<SpecConstantsPass::HandlingMode> SpecConstantMode(
+    "spec-constant-mode", cl::Optional, cl::Hidden,
+    cl::desc("Spec constant handling mode"),
+    cl::values(
+        clEnumValN(SpecConstantsPass::HandlingMode::default_values,
+                   "default_values",
+                   "spec constant uses are replaced by default values"),
+        clEnumValN(SpecConstantsPass::HandlingMode::emulation, "emulation",
+                   "spec constant intrinsics are replaced by RT buffers"),
+        clEnumValN(
+            SpecConstantsPass::HandlingMode::native, "native",
+            "spec constant intrinsics are lowered to spirv intrinsics")));
+
 namespace {
 
 // __sycl* intrinsic names are Itanium ABI-mangled; this is common prefix for
@@ -827,6 +840,9 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
   MapVector<StringRef, MDNode *> SCMetadata;
   SmallVector<MDNode *, 4> DefaultsMetadata;
 
+  if (SpecConstantMode.getNumOccurrences() > 0)
+    Mode = SpecConstantMode;
+
   // Iterate through all declarations of instances of function template
   // template <typename T> T __sycl_get*SpecConstantValue(const char *ID)
   // intrinsic to find its calls and lower them depending on the HandlingMode.

llvm/test/tools/sycl-post-link/spec-constants/default-value/SYCL-alloca.ll renamed to llvm/test/SYCLLowerIR/SpecConstants/SYCL-alloca.ll

@@ -1,7 +1,6 @@
-; RUN: sycl-post-link -properties -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts
-; RUN: FileCheck %s -input-file %t_1.ll --implicit-check-not="SpecConst"
+; RUN: opt -passes=spec-constants -spec-constant-mode=default_values %s -S -o - | FileCheck %s
 
-; This test checks that the post link tool is able to correctly transform
+; 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 }
@@ -20,9 +19,9 @@
 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: alloca float, i64 10, align 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: alloca %my_range, i16 1, align 64
+; 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
 }