Skip to content

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

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

Merged
merged 2 commits into from
Oct 24, 2024
Merged

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

Show file tree
Hide file tree
Changes from 1 commit
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
16 changes: 16 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,19 @@

using namespace llvm;

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

@asudarsa asudarsa Oct 23, 2024

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

@wenju-he wenju-he Oct 24, 2024

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("Spec constant handling mode"),
Copy link
Contributor

@asudarsa asudarsa Oct 23, 2024

Choose a reason for hiding this comment

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

Suggested change
cl::desc("Spec constant handling mode"),
cl::desc("Specialization constant handling mode"),

Copy link
Contributor Author

@wenju-he wenju-he Oct 24, 2024

Choose a reason for hiding this comment

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

done

cl::values(
clEnumValN(SpecConstantsPass::HandlingMode::default_values,
"default_values",
"spec constant uses are replaced by default values"),
Copy link
Contributor

@asudarsa asudarsa Oct 23, 2024

Choose a reason for hiding this comment

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

Suggested change
"spec constant uses are replaced by default values"),
"Specialization constant uses are replaced by default values"),

Copy link
Contributor Author

@wenju-he wenju-he Oct 24, 2024

Choose a reason for hiding this comment

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

done

clEnumValN(SpecConstantsPass::HandlingMode::emulation, "emulation",
"spec constant intrinsics are replaced by RT buffers"),
Copy link
Contributor

@asudarsa asudarsa Oct 23, 2024

Choose a reason for hiding this comment

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

Suggested change
"spec constant intrinsics are replaced by RT buffers"),
"Specialization constant intrinsics are replaced by run-time buffers"),

Copy link
Contributor Author

@wenju-he wenju-he Oct 24, 2024

Choose a reason for hiding this comment

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

done

clEnumValN(
SpecConstantsPass::HandlingMode::native, "native",
"spec constant intrinsics are lowered to spirv intrinsics")));
Copy link
Contributor

@asudarsa asudarsa Oct 23, 2024

Choose a reason for hiding this comment

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

Suggested change
"spec constant intrinsics are lowered to spirv intrinsics")));
"Specialization constant intrinsics are lowered to SPIR-V intrinsics")));

Copy link
Contributor Author

@wenju-he wenju-he Oct 24, 2024

Choose a reason for hiding this comment

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

done


namespace {

// __sycl* intrinsic names are Itanium ABI-mangled; this is common prefix for
Expand Down Expand Up @@ -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.
Expand Down
9 changes: 4 additions & 5 deletions ...ec-constants/default-value/SYCL-alloca.ll → .../SYCLLowerIR/SpecConstants/SYCL-alloca.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Copy link
Contributor

@asudarsa asudarsa Oct 23, 2024

Choose a reason for hiding this comment

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

Is the sycl-post-link test not expected to work anymore? I somehow feel it will be better to add a new test here.

Copy link
Contributor Author

@wenju-he wenju-he Oct 24, 2024

Choose a reason for hiding this comment

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

sycl-post-link test is still expected to work.
This test was only checking output of SpecConstantPass, while other tests in llvm/test/tools/sycl-post-link/ folder are also checking sycl-post-link tool output. So I move this test into llvm/test/SYCLLowerIR/SpecConstants/SYCL-alloca.ll

After reading your comment, I also feel it is better to keep the test to check sycl-post-link is correctly passing parameters to SpecConstantPass. So I added it back in the second commit.

; 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 }
Expand All @@ -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
}
Expand Down
Loading