Skip to content

Commit

Permalink
[SYCL] Correct handling of enumeration specialization constants (#12925)
Browse files Browse the repository at this point in the history 
Lower enumeration specialization constants to `SpecConstant` in SPIR-V
instead of `SpecConstantComposite`, thus handling them using their
underlying integral types.

---------

Signed-off-by: Victor Perez <[email protected]>
  • Loading branch information
@victor-eds
victor-eds authored Mar 8, 2024
1 parent 0d2fe31 commit 1f0dc36
Show file tree
Hide file tree
Showing 2 changed files with 49 additions and 5 deletions.
4 changes: 2 additions & 2 deletions sycl/include/sycl/kernel_handler.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -67,7 +67,7 @@ class __SYCL_TYPE(kernel_handler) kernel_handler {
template <
auto &S,
typename T = typename std::remove_reference_t<decltype(S)>::value_type,
std::enable_if_t<std::is_fundamental_v<T>> * = nullptr>
std::enable_if_t<std::is_scalar_v<T>> * = nullptr>
__SYCL_ALWAYS_INLINE T getSpecializationConstantOnDevice() {
const char *SymbolicID = __builtin_sycl_unique_stable_id(S);
return __sycl_getScalar2020SpecConstantValue<T>(
Expand All @@ -76,7 +76,7 @@ class __SYCL_TYPE(kernel_handler) kernel_handler {
template <
auto &S,
typename T = typename std::remove_reference_t<decltype(S)>::value_type,
std::enable_if_t<std::is_compound_v<T>> * = nullptr>
std::enable_if_t<!std::is_scalar_v<T>> * = nullptr>
__SYCL_ALWAYS_INLINE T getSpecializationConstantOnDevice() {
const char *SymbolicID = __builtin_sycl_unique_stable_id(S);
return __sycl_getComposite2020SpecConstantValue<T>(
Expand Down
50 changes: 47 additions & 3 deletions sycl/test/basic_tests/SYCL-2020-spec-constants.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
// RUN: sycl-post-link %t.bc -spec-const=emulation -o %t-split2.txt
// RUN: cat %t-split2_0.prop | FileCheck %s -check-prefixes=CHECK,CHECK-DEF
// RUN: llvm-spirv -o %t-split1_0.spv -spirv-max-version=1.1 -spirv-ext=+all %t-split1_0.bc
// RUN: llvm-spirv -o %t-split2_0.spv -spirv-max-version=1.1 -spirv-ext=+all %t-split2_0.bc
// RUN: llvm-spirv -o - --to-text %t-split1_0.spv | FileCheck %s -check-prefixes=CHECK-SPV
//
//==----------- SYCL-2020-spec-constants.cpp -------------------------------==//
//
Expand Down Expand Up @@ -61,6 +61,10 @@ struct composite {

constexpr sycl::specialization_id<composite> composite_id(12, 13);

enum class enumeration { a, b, c };

constexpr sycl::specialization_id<enumeration> enumeration_id(enumeration::c);

class SpecializedKernel;

int main() {
Expand Down Expand Up @@ -88,11 +92,12 @@ int main() {
auto f32 = kh.get_specialization_constant<float_id>();
auto f64 = kh.get_specialization_constant<double_id>();
auto c = kh.get_specialization_constant<composite_id>();
auto e = kh.get_specialization_constant<enumeration_id>();

// see FIXMEs above about bool and half types support
acc[0] = /*i1 +*/ i8 + u8 + i16 + u16 + i32 + u32 + i64 + u64 +
//f16.get() +
f32 + f64 + c.a + c.b;
// f16.get() +
f32 + f64 + c.a + c.b + static_cast<float>(e);
});
});
}
Expand All @@ -114,6 +119,7 @@ int main() {
// CHECK-DAG: [[UNIQUE_PREFIX]]____ZL9uint16_id=2|
// CHECK-DAG: [[UNIQUE_PREFIX]]____ZL9uint32_id=2|
// CHECK-DAG: [[UNIQUE_PREFIX]]____ZL9uint64_id=2|
// CHECK-DAG: [[UNIQUE_PREFIX]]____ZL14enumeration_id=2|
// FIXME: check line for half constant

// CHECK-RT: [SYCL/specialization constants default values]
Expand All @@ -133,6 +139,7 @@ int main() {
// CHECK-NEXT-LOG:[[UNIQUE_PREFIX]]____ZL9double_id={9, 0, 8}
// CHECK-NEXT-LOG:[[UNIQUE_PREFIX]]____ZL12composite_id={10, 0, 4}
// CHECK-NEXT-LOG:[[UNIQUE_PREFIX]]____ZL12composite_id={11, 4, 4}
// CHECK-NEXT-LOG:[[UNIQUE_PREFIX]]____ZL14enumeration_id={12, 0, 4}
// CHECK-NEXT-LOG:{0, 1, 42}
// CHECK-NEXT-LOG:{1, 1, 26}
// CHECK-NEXT-LOG:{2, 2, 34}
Expand All @@ -145,3 +152,40 @@ int main() {
// CHECK-NEXT-LOG:{34, 8, 1.100000e+01}
// CHECK-NEXT-LOG:{42, 4, 12}
// CHECK-NEXT-LOG:{46, 4, 13}
// CHECK-NEXT-LOG:{50, 4, 2}

// CHECK-SPV-DAG: Decorate [[#SPEC0:]] SpecId 0
// CHECK-SPV-DAG: Decorate [[#SPEC1:]] SpecId 1
// CHECK-SPV-DAG: Decorate [[#SPEC2:]] SpecId 2
// CHECK-SPV-DAG: Decorate [[#SPEC3:]] SpecId 3
// CHECK-SPV-DAG: Decorate [[#SPEC4:]] SpecId 4
// CHECK-SPV-DAG: Decorate [[#SPEC5:]] SpecId 5
// CHECK-SPV-DAG: Decorate [[#SPEC6:]] SpecId 6
// CHECK-SPV-DAG: Decorate [[#SPEC7:]] SpecId 7
// CHECK-SPV-DAG: Decorate [[#SPEC8:]] SpecId 8
// CHECK-SPV-DAG: Decorate [[#SPEC9:]] SpecId 9
// CHECK-SPV-DAG: Decorate [[#SPEC10:]] SpecId 10
// CHECK-SPV-DAG: Decorate [[#SPEC11:]] SpecId 11
// CHECK-SPV-DAG: Decorate [[#SPEC12:]] SpecId 12

// CHECK-SPV-DAG: TypeInt [[#I64TY:]] 64 0
// CHECK-SPV-DAG: TypeInt [[#I32TY:]] 32 0
// CHECK-SPV-DAG: TypeInt [[#I8TY:]] 8 0
// CHECK-SPV-DAG: TypeInt [[#I16TY:]] 16 0
// CHECK-SPV-DAG: TypeFloat [[#F32TY:]] 32
// CHECK-SPV-DAG: TypeFloat [[#F64TY:]] 64
// CHECK-SPV-DAG: TypeStruct [[#COMPTY:]] [[#I32TY]] [[#I32TY]]
// CHECK-SPV-DAG: SpecConstant [[#I8TY]] [[#SPEC0]] 42
// CHECK-SPV-DAG: SpecConstant [[#I8TY]] [[#SPEC1]] 26
// CHECK-SPV-DAG: SpecConstant [[#I16TY]] [[#SPEC2]] 34
// CHECK-SPV-DAG: SpecConstant [[#I16TY]] [[#SPEC3]] 14
// CHECK-SPV-DAG: SpecConstant [[#I32TY]] [[#SPEC4]] 52
// CHECK-SPV-DAG: SpecConstant [[#I32TY]] [[#SPEC5]] 46
// CHECK-SPV-DAG: SpecConstant [[#I64TY]] [[#SPEC6]] 27 0
// CHECK-SPV-DAG: SpecConstant [[#I64TY]] [[#SPEC7]] 81 0
// CHECK-SPV-DAG: SpecConstant [[#F32TY]] [[#SPEC8]] 1144094720
// CHECK-SPV-DAG: SpecConstant [[#F64TY]] [[#SPEC9]] 0 1076232192
// CHECK-SPV-DAG: SpecConstant [[#I32TY]] [[#SPEC10]] 12
// CHECK-SPV-DAG: SpecConstant [[#I32TY]] [[#SPEC11]] 13
// CHECK-SPV-DAG: SpecConstant [[#I32TY]] [[#SPEC12]] 2
// CHECK-SPV-DAG: SpecConstantComposite [[#COMPTY]] {{.*}} [[#SPEC10]] [[#SPEC11]]

0 comments on commit 1f0dc36

Please sign in to comment.