Skip to content

Commit

Permalink
Merge pull request #872 from fknorr/compat-half
Browse files Browse the repository at this point in the history
Compatibility: Only use `sycl::half` if SYCL_CTS_ENABLE_HALF_TESTS is set
  • Loading branch information
bader authored Jan 9, 2025
2 parents cc2ad1f + 9d923ce commit b9cc217
Show file tree
Hide file tree
Showing 28 changed files with 1,356 additions and 1,183 deletions.
3 changes: 3 additions & 0 deletions tests/accessor_basic/accessor_common.h
Original file line number Diff line number Diff line change
Expand Up @@ -143,13 +143,16 @@ inline std::string get_section_name(const std::string& type_name,
// FIXME: re-enable when marrray is implemented in adaptivecpp and type_coverage
// is enabled
#ifndef SYCL_CTS_COMPILING_WITH_ADAPTIVECPP

#if SYCL_CTS_ENABLE_HALF_TESTS
/**
* @brief Factory function for getting type_pack with fp16 type
*/
inline auto get_fp16_type() {
static const auto types = named_type_pack<sycl::half>::generate("sycl::half");
return types;
}
#endif

/**
* @brief Factory function for getting type_pack with fp64 type
Expand Down
2 changes: 2 additions & 0 deletions tests/common/common.h
Original file line number Diff line number Diff line change
Expand Up @@ -237,6 +237,7 @@ bool check_type_sign(bool expected_sign) {
return (std::is_signed<T>::value == expected_sign);
}

#if SYCL_CTS_ENABLE_HALF_TESTS
/**
* @brief Helper function to see if sycl::half is of the wrong sign
*/
Expand All @@ -245,6 +246,7 @@ inline bool check_type_sign<sycl::half>(bool expected_sign) {
bool is_signed = sycl::half(1) > sycl::half(-1);
return is_signed == expected_sign;
}
#endif

/**
* @brief Helper function to log a failure if a type is of the wrong size or
Expand Down
30 changes: 13 additions & 17 deletions tests/common/common_vec.h
Original file line number Diff line number Diff line change
Expand Up @@ -57,8 +57,7 @@ bool check_vector_size(sycl::vec<vecType, numOfElems> vector) {
* @brief Helper function to check vector values are correct.
*/
template <typename vecType, int numOfElems>
bool check_vector_values(sycl::vec<vecType, numOfElems> vector,
vecType* vals) {
bool check_vector_values(sycl::vec<vecType, numOfElems> vector, vecType* vals) {
for (int i = 0; i < numOfElems; i++) {
if ((vals[i] != vector[i])) {
return false;
Expand All @@ -72,14 +71,12 @@ bool check_vector_values(sycl::vec<vecType, numOfElems> vector,
* for division result are accurate enough
*/
template <typename vecType, int numOfElems>
typename std::enable_if<is_sycl_floating_point<vecType>::value, bool>::type
check_vector_values_div(sycl::vec<vecType, numOfElems> vector,
vecType *vals) {
typename std::enable_if_t<is_sycl_scalar_floating_point_v<vecType>, bool>
check_vector_values_div(sycl::vec<vecType, numOfElems> vector, vecType* vals) {
for (int i = 0; i < numOfElems; i++) {
vecType vectorValue = vector[i];
if (vals[i] == vectorValue)
continue;
const vecType ulpsExpected = 2.5; // Min Accuracy for x / y
if (vals[i] == vectorValue) continue;
const vecType ulpsExpected = 2.5; // Min Accuracy for x / y
const vecType difference = sycl::fabs(vectorValue - vals[i]);
// using sycl functions to get ulp because it used in kernel
const vecType differenceExpected = ulpsExpected * get_ulp_sycl(vals[i]);
Expand All @@ -95,9 +92,8 @@ check_vector_values_div(sycl::vec<vecType, numOfElems> vector,
* @brief Helper function to check that vector values for division are correct
*/
template <typename vecType, int numOfElems>
typename std::enable_if<!is_sycl_floating_point<vecType>::value, bool>::type
check_vector_values_div(sycl::vec<vecType, numOfElems> vector,
vecType *vals) {
typename std::enable_if_t<!is_sycl_scalar_floating_point_v<vecType>, bool>
check_vector_values_div(sycl::vec<vecType, numOfElems> vector, vecType* vals) {
return check_vector_values(vector, vals);
}

Expand All @@ -123,7 +119,8 @@ bool check_single_vector_op(vectorType vector1, lambdaFunc lambda) {

template <typename sourceType, typename targetType>
static constexpr bool if_FP_to_non_FP_conv_v =
is_sycl_floating_point<sourceType>::value && !is_sycl_floating_point<targetType>::value;
is_sycl_scalar_floating_point_v<sourceType> &&
!is_sycl_scalar_floating_point_v<targetType>;

template <typename vecType, int N, typename convertType>
sycl::vec<convertType, N> convert_vec(sycl::vec<vecType, N> inputVec) {
Expand Down Expand Up @@ -196,7 +193,7 @@ sycl::vec<convertType, N> rtn(sycl::vec<vecType, N> inputVec) {
// values instead.
template <typename vecType, int N, typename convertType>
void handleFPToUnsignedConv(sycl::vec<vecType, N>& inputVec) {
if constexpr (is_sycl_floating_point<vecType>::value &&
if constexpr (is_sycl_scalar_floating_point_v<vecType> &&
std::is_unsigned_v<convertType>) {
for (size_t i = 0; i < N; ++i) {
vecType elem = inputVec[i];
Expand Down Expand Up @@ -247,7 +244,7 @@ bool check_vector_convert_result_impl(sycl::vec<vecType, N> inputVec,
sycl::vec<convertType, N> expectedVec;
switch (mode) {
case sycl::rounding_mode::automatic:
if constexpr (is_sycl_floating_point<vecType>::value) {
if constexpr (is_sycl_scalar_floating_point_v<vecType>) {
expectedVec = rte<vecType, N, convertType>(inputVec);
} else {
expectedVec = rtz<vecType, N, convertType>(inputVec);
Expand Down Expand Up @@ -291,9 +288,8 @@ bool check_vector_convert_result(sycl::vec<vecType, N> inputVec) {
template <typename vecType, int N, typename convertType>
bool check_vector_convert_modes(sycl::vec<vecType, N> inputVec) {
bool flag = true;
flag &=
check_vector_convert_result<vecType, N, convertType,
sycl::rounding_mode::automatic>(inputVec);
flag &= check_vector_convert_result<vecType, N, convertType,
sycl::rounding_mode::automatic>(inputVec);
#if SYCL_CTS_ENABLE_FULL_CONFORMANCE
flag &= check_vector_convert_result<vecType, N, convertType,
sycl::rounding_mode::rte>(inputVec);
Expand Down
6 changes: 4 additions & 2 deletions tests/group_functions/group_functions_common.h
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,7 @@
//
*******************************************************************************/

#include "../../util/type_traits.h"
#include "../common/common.h"
#include "../common/get_group_range.h"
#include "../common/once_per_unit.h"
Expand Down Expand Up @@ -133,8 +134,10 @@ struct custom_type {
template <typename T>
inline constexpr uint64_t exact_max = std::numeric_limits<T>::max();

#if SYCL_CTS_ENABLE_HALF_TESTS
template <>
inline constexpr uint64_t exact_max<sycl::half> = 1ull << 11;
#endif
template <>
inline constexpr uint64_t exact_max<float> = 1ull << 24;
template <>
Expand Down Expand Up @@ -256,8 +259,7 @@ template <typename T>
inline auto get_op_types() {
#if SYCL_CTS_ENABLE_FULL_CONFORMANCE
static const auto types = []() {
if constexpr (std::is_floating_point_v<T> ||
std::is_same_v<std::remove_cv_t<T>, sycl::half>) {
if constexpr (is_sycl_scalar_floating_point_v<T>) {
// Bitwise operations are not defined for floating point types.
return named_type_pack<sycl::plus<T>, sycl::multiplies<T>,
sycl::logical_and<T>, sycl::logical_or<T>,
Expand Down
10 changes: 8 additions & 2 deletions tests/group_functions/group_joint_reduce.cpp.in
Original file line number Diff line number Diff line change
Expand Up @@ -32,12 +32,15 @@ TEST_CASE(CTS_TYPE_NAME + " group and sub-group joint reduce functions",
const auto Operators = get_op_types<CTS_TYPE>();
const auto RetType = unnamed_type_pack<CTS_TYPE>();

#if SYCL_CTS_ENABLE_HALF_TESTS
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, sycl::half>) {
if (!queue.get_device().has(sycl::aspect::fp16))
SKIP(
"Device does not support half precision floating point "
"operations.");
} else if (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
}
#endif
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
if (!queue.get_device().has(sycl::aspect::fp64))
SKIP(
"Device does not support double precision floating point "
Expand All @@ -57,12 +60,15 @@ TEMPLATE_LIST_TEST_CASE(
const auto RetType = unnamed_type_pack<CTS_TYPE>();
const auto ReducedType = unnamed_type_pack<TestType>();

#if SYCL_CTS_ENABLE_HALF_TESTS
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, sycl::half>) {
if (!queue.get_device().has(sycl::aspect::fp16))
SKIP(
"Device does not support half precision floating point "
"operations.");
} else if (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
}
#endif
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
if (!queue.get_device().has(sycl::aspect::fp64))
SKIP(
"Device does not support double precision floating point "
Expand Down
10 changes: 8 additions & 2 deletions tests/group_functions/group_joint_scan.cpp.in
Original file line number Diff line number Diff line change
Expand Up @@ -49,12 +49,15 @@ DISABLED_FOR_TEST_CASE(AdaptiveCpp)
#endif

auto queue = once_per_unit::get_queue();
#if SYCL_CTS_ENABLE_HALF_TESTS
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, sycl::half>) {
if (!queue.get_device().has(sycl::aspect::fp16))
SKIP(
"Device does not support half precision floating point "
"operations.");
} else if (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
}
#endif
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
if (!queue.get_device().has(sycl::aspect::fp64))
SKIP(
"Device does not support double precision floating point "
Expand Down Expand Up @@ -86,12 +89,15 @@ DISABLED_FOR_TEST_CASE(AdaptiveCpp)
#endif

auto queue = once_per_unit::get_queue();
#if SYCL_CTS_ENABLE_HALF_TESTS
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, sycl::half>) {
if (!queue.get_device().has(sycl::aspect::fp16))
SKIP(
"Device does not support half precision floating point "
"operations.");
} else if (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
}
#endif
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
if (!queue.get_device().has(sycl::aspect::fp64))
SKIP(
"Device does not support double precision floating point "
Expand Down
10 changes: 8 additions & 2 deletions tests/group_functions/group_reduce_over_group.cpp.in
Original file line number Diff line number Diff line change
Expand Up @@ -33,12 +33,15 @@ TEST_CASE(CTS_TYPE_NAME + " group and sub-group reduce functions",
const auto Operators = get_op_types<CTS_TYPE>();
const auto RetType = unnamed_type_pack<CTS_TYPE>();

#if SYCL_CTS_ENABLE_HALF_TESTS
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, sycl::half>) {
if (!queue.get_device().has(sycl::aspect::fp16))
SKIP(
"Device does not support half precision floating point "
"operations.");
} else if (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
}
#endif
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
if (!queue.get_device().has(sycl::aspect::fp64))
SKIP(
"Device does not support double precision floating point "
Expand All @@ -59,12 +62,15 @@ TEMPLATE_LIST_TEST_CASE(CTS_TYPE_NAME +
const auto RetType = unnamed_type_pack<CTS_TYPE>();
const auto ReducedType = unnamed_type_pack<TestType>();

#if SYCL_CTS_ENABLE_HALF_TESTS
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, sycl::half>) {
if (!queue.get_device().has(sycl::aspect::fp16))
SKIP(
"Device does not support half precision floating point "
"operations.");
} else if (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
}
#endif
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
if (!queue.get_device().has(sycl::aspect::fp64))
SKIP(
"Device does not support double precision floating point "
Expand Down
10 changes: 8 additions & 2 deletions tests/group_functions/group_scan_over_group.cpp.in
Original file line number Diff line number Diff line change
Expand Up @@ -39,12 +39,15 @@ DISABLED_FOR_TEST_CASE(AdaptiveCpp)
(CTS_TYPE_NAME + " group and sub-group scan functions",
"[group_func][type_list][dim]")({
auto queue = once_per_unit::get_queue();
#if SYCL_CTS_ENABLE_HALF_TESTS
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, sycl::half>) {
if (!queue.get_device().has(sycl::aspect::fp16))
SKIP(
"Device does not support half precision floating point "
"operations.");
} else if (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
}
#endif
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
if (!queue.get_device().has(sycl::aspect::fp64))
SKIP(
"Device does not support double precision floating point "
Expand All @@ -59,12 +62,15 @@ DISABLED_FOR_TEST_CASE(AdaptiveCpp)
(CTS_TYPE_NAME + " group and sub-group scan functions with init",
"[group_func][type_list][dim]")({
auto queue = once_per_unit::get_queue();
#if SYCL_CTS_ENABLE_HALF_TESTS
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, sycl::half>) {
if (!queue.get_device().has(sycl::aspect::fp16))
SKIP(
"Device does not support half precision floating point "
"operations.");
} else if (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
}
#endif
if constexpr (std::is_same_v<std::remove_cv_t<CTS_TYPE>, double>) {
if (!queue.get_device().has(sycl::aspect::fp64))
SKIP(
"Device does not support double precision floating point "
Expand Down
4 changes: 4 additions & 0 deletions tests/kernel_bundle/kernels.h
Original file line number Diff line number Diff line change
Expand Up @@ -221,6 +221,8 @@ struct kernel_atomic64_descriptor {
// fp16, fp64, atomic64 kernels without sycl::requires attribute but with
// explicit operations

#if SYCL_CTS_ENABLE_HALF_TESTS

struct kernel_fp16_no_attr : kernel_base {
void operator()(sycl::item<1> id) const {
if (id.get_linear_id() == 0) {
Expand All @@ -239,6 +241,8 @@ struct kernel_fp16_no_attr_descriptor {
}
};

#endif // SYCL_CTS_ENABLE_HALF_TESTS

struct kernel_fp64_no_attr : kernel_base {
void operator()(sycl::item<1> id) const {
if (id.get_linear_id() == 0) {
Expand Down
4 changes: 4 additions & 0 deletions tests/kernel_bundle/sycl_is_compatible.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -102,10 +102,12 @@ TEST_CASE("Check is_compatible for kernels with no kernel attributes",
CHECK(sycl::is_compatible(builtinKernelIds, device));
}

#if SYCL_CTS_ENABLE_HALF_TESTS
SECTION("for a kernel that uses `sycl::half`") {
check_with_optional_features<kernels::kernel_fp16_no_attr_descriptor>(
device, queue, device.has(sycl::aspect::fp16));
}
#endif

SECTION("for a kernel that uses `double`") {
check_with_optional_features<kernels::kernel_fp64_no_attr_descriptor>(
Expand Down Expand Up @@ -204,10 +206,12 @@ TEST_CASE(
const sycl::device device = sycl_cts::util::get_cts_object::device();
sycl::queue queue = sycl_cts::util::get_cts_object::queue();

#if SYCL_CTS_ENABLE_HALF_TESTS
SECTION("for a kernel that uses `sycl::half`") {
check_with_optional_features<kernels::kernel_fp16_descriptor>(
device, queue, device.has(sycl::aspect::fp16));
}
#endif

SECTION("for a kernel that uses `double`") {
check_with_optional_features<kernels::kernel_fp64_descriptor>(
Expand Down
17 changes: 9 additions & 8 deletions tests/marray_basic/marray_operators.h
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,7 @@
#ifndef SYCLCTS_TESTS_MARRAY_MARRAY_OPERATOR_H
#define SYCLCTS_TESTS_MARRAY_MARRAY_OPERATOR_H

#include "../../util/type_traits.h"
#include "../common/common.h"
#include "../common/section_name_builder.h"
#include "marray_common.h"
Expand Down Expand Up @@ -53,10 +54,10 @@ struct operators_helper {
// similar to native::divide we can skip checking them.
template <typename OpT, typename ElemT>
struct skip_result_check
: std::bool_constant<
(std::is_same_v<OpT, op_div> || std::is_same_v<OpT, op_assign_div>)&&(
std::is_same_v<ElemT, float> || std::is_same_v<ElemT, double> ||
std::is_same_v<ElemT, sycl::half>)> {};
: std::bool_constant<(
std::is_same_v<OpT, op_div> ||
std::is_same_v<
OpT, op_assign_div>)&&is_sycl_scalar_floating_point_v<ElemT>> {};

template <typename OpT, typename ElemT>
constexpr bool skip_result_check_v = skip_result_check<OpT, ElemT>::value;
Expand All @@ -67,10 +68,10 @@ bool are_equal_ignore_division(const T1& lhs, const T1& rhs) {
// similar to native::divide we can skip checking them here.
constexpr bool is_div =
std::is_same_v<OpT, op_div> || std::is_same_v<OpT, op_assign_div>;
constexpr bool is_sycl_floating_point = std::is_same_v<ElemT, float> ||
std::is_same_v<ElemT, double> ||
std::is_same_v<ElemT, sycl::half>;
if constexpr (is_div && is_sycl_floating_point) return true;
constexpr bool is_sycl_scalar_floating_point =
std::is_same_v<ElemT, float> || std::is_same_v<ElemT, double> ||
std::is_same_v<ElemT, sycl::half>;
if constexpr (is_div && is_sycl_scalar_floating_point) return true;
return value_operations::are_equal(lhs, rhs);
}

Expand Down
Loading

0 comments on commit b9cc217

Please sign in to comment.