diff --git a/sycl/include/sycl/kernel_handler.hpp b/sycl/include/sycl/kernel_handler.hpp index 738c23b3eda97..50504fe5ae2cf 100644 --- a/sycl/include/sycl/kernel_handler.hpp +++ b/sycl/include/sycl/kernel_handler.hpp @@ -67,7 +67,7 @@ class __SYCL_TYPE(kernel_handler) kernel_handler { template < auto &S, typename T = typename std::remove_reference_t::value_type, - std::enable_if_t> * = nullptr> + std::enable_if_t> * = nullptr> __SYCL_ALWAYS_INLINE T getSpecializationConstantOnDevice() { const char *SymbolicID = __builtin_sycl_unique_stable_id(S); return __sycl_getScalar2020SpecConstantValue( @@ -76,7 +76,7 @@ class __SYCL_TYPE(kernel_handler) kernel_handler { template < auto &S, typename T = typename std::remove_reference_t::value_type, - std::enable_if_t> * = nullptr> + std::enable_if_t> * = nullptr> __SYCL_ALWAYS_INLINE T getSpecializationConstantOnDevice() { const char *SymbolicID = __builtin_sycl_unique_stable_id(S); return __sycl_getComposite2020SpecConstantValue( diff --git a/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp index e102773934438..b9d29507a8eac 100644 --- a/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp +++ b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp @@ -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 -------------------------------==// // @@ -61,6 +61,10 @@ struct composite { constexpr sycl::specialization_id composite_id(12, 13); +enum class enumeration { a, b, c }; + +constexpr sycl::specialization_id enumeration_id(enumeration::c); + class SpecializedKernel; int main() { @@ -88,11 +92,12 @@ int main() { auto f32 = kh.get_specialization_constant(); auto f64 = kh.get_specialization_constant(); auto c = kh.get_specialization_constant(); + auto e = kh.get_specialization_constant(); // 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(e); }); }); } @@ -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] @@ -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} @@ -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]]