Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

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

Merged
merged 2 commits into from
Oct 24, 2024
+49 −0
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
87eb7a4
[NFC][SYCL][SYCLLowerIR] Add cl::opt SpecConstantMode for LIT testing
wenju-he Oct 23, 2024
98079ae
init value, update comment, add back sycl-post-link test
wenju-he Oct 24, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 18 additions & 0 deletions llvm/lib/SYCLLowerIR/SpecConstants.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,21 @@

using namespace llvm;

static cl::opt<SpecConstantsPass::HandlingMode> SpecConstantMode(
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

is there a default value that can be set here?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

added default value: cl::init(SpecConstantsPass::HandlingMode::emulation)

"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
Expand Down Expand Up @@ -827,6 +842,9 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
MapVector<StringRef, MDNode *> SCMetadata;
SmallVector<MDNode *, 4> DefaultsMetadata;

if (SpecConstantMode.getNumOccurrences() > 0)
asudarsa marked this conversation as resolved.
Show resolved Hide resolved
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.
Expand Down
31 changes: 31 additions & 0 deletions llvm/test/SYCLLowerIR/SpecConstants/SYCL-alloca.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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)
Loading