diff --git a/dpctl-capi/include/dpctl_sycl_device_interface.h b/dpctl-capi/include/dpctl_sycl_device_interface.h index 62d292ce31..293c511f07 100644 --- a/dpctl-capi/include/dpctl_sycl_device_interface.h +++ b/dpctl-capi/include/dpctl_sycl_device_interface.h @@ -266,3 +266,99 @@ bool DPCTLDevice_HasAspect(__dpctl_keep const DPCTLSyclDeviceRef DRef, DPCTLSyclAspectType AT); DPCTL_C_EXTERN_C_END + +/*! + * @brief Wrapper over + * device.get_info. + * + * @param DRef Opaque pointer to a sycl::device + * @return Returns true if the device supports independent forward progress of + * sub-groups with respect to other sub-groups in the same work-group. + */ +DPCTL_API +bool DPCTLDevice_GetSubGroupIndependentForwardProgress( + __dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Wrapper over + * device.get_info. + * + * @param DRef Opaque pointer to a sycl::device + * @return Returns the preferred native vector width size for built-in scalar + * types that can be put into vectors. + */ +DPCTL_API +uint32_t DPCTLDevice_GetPreferredVectorWidthChar( + __dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Wrapper over + * device.get_info. + * + * @param DRef Opaque pointer to a sycl::device + * @return Returns the preferred native vector width size for built-in scalar + * types that can be put into vectors. + */ +DPCTL_API +uint32_t DPCTLDevice_GetPreferredVectorWidthShort( + __dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Wrapper over + * device.get_info. + * + * @param DRef Opaque pointer to a sycl::device + * @return Returns the preferred native vector width size for built-in scalar + * types that can be put into vectors. + */ +DPCTL_API +uint32_t DPCTLDevice_GetPreferredVectorWidthInt( + __dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Wrapper over + * device.get_info. + * + * @param DRef Opaque pointer to a sycl::device + * @return Returns the preferred native vector width size for built-in scalar + * types that can be put into vectors. + */ +DPCTL_API +uint32_t DPCTLDevice_GetPreferredVectorWidthLong( + __dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Wrapper over + * device.get_info. + * + * @param DRef Opaque pointer to a sycl::device + * @return Returns the preferred native vector width size for built-in scalar + * types that can be put into vectors. + */ +DPCTL_API +uint32_t DPCTLDevice_GetPreferredVectorWidthFloat( + __dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Wrapper over + * device.get_info. + * + * @param DRef Opaque pointer to a sycl::device + * @return Returns the preferred native vector width size for built-in scalar + * types that can be put into vectors. + */ +DPCTL_API +uint32_t DPCTLDevice_GetPreferredVectorWidthDouble( + __dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Wrapper over + * device.get_info. + * + * @param DRef Opaque pointer to a sycl::device + * @return Returns the preferred native vector width size for built-in scalar + * types that can be put into vectors. + */ +DPCTL_API +uint32_t DPCTLDevice_GetPreferredVectorWidthHalf( + __dpctl_keep const DPCTLSyclDeviceRef DRef); diff --git a/dpctl-capi/source/dpctl_sycl_device_interface.cpp b/dpctl-capi/source/dpctl_sycl_device_interface.cpp index 7f6405f667..ff37d4834b 100644 --- a/dpctl-capi/source/dpctl_sycl_device_interface.cpp +++ b/dpctl-capi/source/dpctl_sycl_device_interface.cpp @@ -387,3 +387,139 @@ bool DPCTLDevice_HasAspect(__dpctl_keep const DPCTLSyclDeviceRef DRef, } return hasAspect; } + +bool DPCTLDevice_GetSubGroupIndependentForwardProgress( + __dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + bool SubGroupProgress = false; + auto D = unwrap(DRef); + if (D) { + try { + SubGroupProgress = D->get_info< + info::device::sub_group_independent_forward_progress>(); + } catch (runtime_error const &re) { + // \todo log error + std::cerr << re.what() << '\n'; + } + } + return SubGroupProgress; +} + +uint32_t DPCTLDevice_GetPreferredVectorWidthChar( + __dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + size_t vector_width_char = 0; + auto D = unwrap(DRef); + if (D) { + try { + vector_width_char = + D->get_info(); + } catch (runtime_error const &re) { + // \todo log error + std::cerr << re.what() << '\n'; + } + } + return vector_width_char; +} + +uint32_t DPCTLDevice_GetPreferredVectorWidthShort( + __dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + size_t vector_width_short = 0; + auto D = unwrap(DRef); + if (D) { + try { + vector_width_short = + D->get_info(); + } catch (runtime_error const &re) { + // \todo log error + std::cerr << re.what() << '\n'; + } + } + return vector_width_short; +} + +uint32_t DPCTLDevice_GetPreferredVectorWidthInt( + __dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + size_t vector_width_int = 0; + auto D = unwrap(DRef); + if (D) { + try { + vector_width_int = + D->get_info(); + } catch (runtime_error const &re) { + // \todo log error + std::cerr << re.what() << '\n'; + } + } + return vector_width_int; +} + +uint32_t DPCTLDevice_GetPreferredVectorWidthLong( + __dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + size_t vector_width_long = 0; + auto D = unwrap(DRef); + if (D) { + try { + vector_width_long = + D->get_info(); + } catch (runtime_error const &re) { + // \todo log error + std::cerr << re.what() << '\n'; + } + } + return vector_width_long; +} + +uint32_t DPCTLDevice_GetPreferredVectorWidthFloat( + __dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + size_t vector_width_float = 0; + auto D = unwrap(DRef); + if (D) { + try { + vector_width_float = + D->get_info(); + } catch (runtime_error const &re) { + // \todo log error + std::cerr << re.what() << '\n'; + } + } + return vector_width_float; +} + +uint32_t DPCTLDevice_GetPreferredVectorWidthDouble( + __dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + size_t vector_width_double = 0; + auto D = unwrap(DRef); + if (D) { + try { + vector_width_double = + D->get_info(); + } catch (runtime_error const &re) { + // \todo log error + std::cerr << re.what() << '\n'; + } + } + return vector_width_double; +} + +uint32_t DPCTLDevice_GetPreferredVectorWidthHalf( + __dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + size_t vector_width_half = 0; + auto D = unwrap(DRef); + if (D) { + try { + vector_width_half = + D->get_info(); + } catch (runtime_error const &re) { + // \todo log error + std::cerr << re.what() << '\n'; + } + } + return vector_width_half; +} diff --git a/dpctl-capi/tests/test_sycl_device_interface.cpp b/dpctl-capi/tests/test_sycl_device_interface.cpp index a25f35d926..1d18f8a9b4 100644 --- a/dpctl-capi/tests/test_sycl_device_interface.cpp +++ b/dpctl-capi/tests/test_sycl_device_interface.cpp @@ -24,6 +24,7 @@ /// //===----------------------------------------------------------------------===// +#include "../helper/include/dpctl_utils_helper.h" #include "dpctl_sycl_device_interface.h" #include "dpctl_sycl_device_selector_interface.h" #include "dpctl_sycl_platform_interface.h" @@ -268,6 +269,128 @@ TEST_P(TestDPCTLSyclDeviceInterface, Chk_IsHost) EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef)); } +TEST_P(TestDPCTLSyclDeviceInterface, Chk_GetSubGroupIndependentForwardProgress) +{ + DPCTLSyclDeviceRef DRef = nullptr; + bool sub_group_progress = 0; + EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef)); + if (!DRef) + GTEST_SKIP_("Device not found"); + EXPECT_NO_FATAL_FAILURE( + sub_group_progress = + DPCTLDevice_GetSubGroupIndependentForwardProgress(DRef)); + auto D = reinterpret_cast(DRef); + auto get_sub_group_progress = + D->get_info(); + EXPECT_TRUE(get_sub_group_progress == sub_group_progress); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef)); +} + +TEST_P(TestDPCTLSyclDeviceInterface, Chk_GetPreferredVectorWidthChar) +{ + DPCTLSyclDeviceRef DRef = nullptr; + size_t vector_width_char = 0; + EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef)); + if (!DRef) + GTEST_SKIP_("Device not found"); + EXPECT_NO_FATAL_FAILURE(vector_width_char = + DPCTLDevice_GetPreferredVectorWidthChar(DRef)); + EXPECT_TRUE(vector_width_char != 0); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef)); +} + +TEST_P(TestDPCTLSyclDeviceInterface, Chk_GetPreferredVectorWidthShort) +{ + DPCTLSyclDeviceRef DRef = nullptr; + size_t vector_width_short = 0; + EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef)); + if (!DRef) + GTEST_SKIP_("Device not found"); + EXPECT_NO_FATAL_FAILURE(vector_width_short = + DPCTLDevice_GetPreferredVectorWidthShort(DRef)); + EXPECT_TRUE(vector_width_short != 0); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef)); +} + +TEST_P(TestDPCTLSyclDeviceInterface, Chk_GetPreferredVectorWidthInt) +{ + DPCTLSyclDeviceRef DRef = nullptr; + size_t vector_width_int = 0; + EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef)); + if (!DRef) + GTEST_SKIP_("Device not found"); + EXPECT_NO_FATAL_FAILURE(vector_width_int = + DPCTLDevice_GetPreferredVectorWidthInt(DRef)); + EXPECT_TRUE(vector_width_int != 0); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef)); +} + +TEST_P(TestDPCTLSyclDeviceInterface, Chk_GetPreferredVectorWidthLong) +{ + DPCTLSyclDeviceRef DRef = nullptr; + size_t vector_width_long = 0; + EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef)); + if (!DRef) + GTEST_SKIP_("Device not found"); + EXPECT_NO_FATAL_FAILURE(vector_width_long = + DPCTLDevice_GetPreferredVectorWidthLong(DRef)); + EXPECT_TRUE(vector_width_long != 0); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef)); +} + +TEST_P(TestDPCTLSyclDeviceInterface, Chk_GetPreferredVectorWidthFloat) +{ + DPCTLSyclDeviceRef DRef = nullptr; + size_t vector_width_float = 0; + EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef)); + if (!DRef) + GTEST_SKIP_("Device not found"); + EXPECT_NO_FATAL_FAILURE(vector_width_float = + DPCTLDevice_GetPreferredVectorWidthFloat(DRef)); + EXPECT_TRUE(vector_width_float != 0); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef)); +} + +TEST_P(TestDPCTLSyclDeviceInterface, Chk_GetPreferredVectorWidthDouble) +{ + DPCTLSyclDeviceRef DRef = nullptr; + size_t vector_width_double = 0; + EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef)); + if (!DRef) + GTEST_SKIP_("Device not found"); + EXPECT_NO_FATAL_FAILURE( + vector_width_double = DPCTLDevice_GetPreferredVectorWidthDouble(DRef)); + if (DPCTLDevice_HasAspect(DRef, DPCTL_SyclAspectToDPCTLAspectType( + DPCTL_StrToAspectType("fp64")))) + { + EXPECT_TRUE(vector_width_double != 0); + } + else { + EXPECT_TRUE(vector_width_double == 0); + } + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef)); +} + +TEST_P(TestDPCTLSyclDeviceInterface, Chk_GetPreferredVectorWidthHalf) +{ + DPCTLSyclDeviceRef DRef = nullptr; + size_t vector_width_half = 0; + EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef)); + if (!DRef) + GTEST_SKIP_("Device not found"); + EXPECT_NO_FATAL_FAILURE(vector_width_half = + DPCTLDevice_GetPreferredVectorWidthHalf(DRef)); + if (DPCTLDevice_HasAspect(DRef, DPCTL_SyclAspectToDPCTLAspectType( + DPCTL_StrToAspectType("fp16")))) + { + EXPECT_TRUE(vector_width_half != 0); + } + else { + EXPECT_TRUE(vector_width_half == 0); + } + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef)); +} + INSTANTIATE_TEST_SUITE_P(DPCTLDevice_Fns, TestDPCTLSyclDeviceInterface, ::testing::Values("opencl", diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 7aab8ca8d5..4b206a0086 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -153,6 +153,14 @@ cdef extern from "dpctl_sycl_device_interface.h": cdef bool DPCTLDevice_IsGPU(const DPCTLSyclDeviceRef DRef) cdef bool DPCTLDevice_IsHost(const DPCTLSyclDeviceRef DRef) cdef bool DPCTLDevice_IsHostUnifiedMemory(const DPCTLSyclDeviceRef DRef) + cdef bool DPCTLDevice_GetSubGroupIndependentForwardProgress(const DPCTLSyclDeviceRef DRef) + cdef uint32_t DPCTLDevice_GetPreferredVectorWidthChar(const DPCTLSyclDeviceRef DRef) + cdef uint32_t DPCTLDevice_GetPreferredVectorWidthShort(const DPCTLSyclDeviceRef DRef) + cdef uint32_t DPCTLDevice_GetPreferredVectorWidthInt(const DPCTLSyclDeviceRef DRef) + cdef uint32_t DPCTLDevice_GetPreferredVectorWidthLong(const DPCTLSyclDeviceRef DRef) + cdef uint32_t DPCTLDevice_GetPreferredVectorWidthFloat(const DPCTLSyclDeviceRef DRef) + cdef uint32_t DPCTLDevice_GetPreferredVectorWidthDouble(const DPCTLSyclDeviceRef DRef) + cdef uint32_t DPCTLDevice_GetPreferredVectorWidthHalf(const DPCTLSyclDeviceRef DRef) cpdef bool DPCTLDevice_HasAspect( const DPCTLSyclDeviceRef DRef, DPCTLSyclAspectType AT) diff --git a/dpctl/_sycl_device.pyx b/dpctl/_sycl_device.pyx index cdfcfca4dc..7a53566707 100644 --- a/dpctl/_sycl_device.pyx +++ b/dpctl/_sycl_device.pyx @@ -53,6 +53,14 @@ from ._backend cimport ( DPCTLSyclDeviceSelectorRef, DPCTLDevice_HasAspect, DPCTLSyclDeviceType, + DPCTLDevice_GetSubGroupIndependentForwardProgress, + DPCTLDevice_GetPreferredVectorWidthChar, + DPCTLDevice_GetPreferredVectorWidthShort, + DPCTLDevice_GetPreferredVectorWidthInt, + DPCTLDevice_GetPreferredVectorWidthLong, + DPCTLDevice_GetPreferredVectorWidthFloat, + DPCTLDevice_GetPreferredVectorWidthDouble, + DPCTLDevice_GetPreferredVectorWidthHalf, ) from . import backend_type, device_type from libc.stdint cimport uint32_t @@ -457,6 +465,62 @@ cdef class SyclDevice(_SyclDevice): ) return max_num_sub_groups + @property + def sub_group_independent_forward_progress(self): + """ Returns true if the device supports independent forward progress of + sub-groups with respect to other sub-groups in the same work-group. + """ + return DPCTLDevice_GetSubGroupIndependentForwardProgress(self._device_ref) + + @property + def preferred_vector_width_char(self): + """ Returns the preferred native vector width size for built-in scalar + types that can be put into vectors. + """ + return DPCTLDevice_GetPreferredVectorWidthChar(self._device_ref) + + @property + def preferred_vector_width_short(self): + """ Returns the preferred native vector width size for built-in scalar + types that can be put into vectors. + """ + return DPCTLDevice_GetPreferredVectorWidthShort(self._device_ref) + + @property + def preferred_vector_width_int(self): + """ Returns the preferred native vector width size for built-in scalar + types that can be put into vectors. + """ + return DPCTLDevice_GetPreferredVectorWidthInt(self._device_ref) + + @property + def preferred_vector_width_long(self): + """ Returns the preferred native vector width size for built-in scalar + types that can be put into vectors. + """ + return DPCTLDevice_GetPreferredVectorWidthLong(self._device_ref) + + @property + def preferred_vector_width_float(self): + """ Returns the preferred native vector width size for built-in scalar + types that can be put into vectors. + """ + return DPCTLDevice_GetPreferredVectorWidthFloat(self._device_ref) + + @property + def preferred_vector_width_double(self): + """ Returns the preferred native vector width size for built-in scalar + types that can be put into vectors. + """ + return DPCTLDevice_GetPreferredVectorWidthDouble(self._device_ref) + + @property + def preferred_vector_width_half(self): + """ Returns the preferred native vector width size for built-in scalar + types that can be put into vectors. + """ + return DPCTLDevice_GetPreferredVectorWidthHalf(self._device_ref) + @property def vendor_name(self): """ Returns the device vendor name as a string. diff --git a/dpctl/tests/test_sycl_device.py b/dpctl/tests/test_sycl_device.py index 22fa118277..4d2211c323 100644 --- a/dpctl/tests/test_sycl_device.py +++ b/dpctl/tests/test_sycl_device.py @@ -240,6 +240,62 @@ def check_is_host(device): pytest.fail("is_hostcall failed") +def check_get_sub_group_independent_forward_progress(device): + try: + device.sub_group_independent_forward_progress + except Exception: + pytest.fail("sub_group_independent_forward_progress call failed") + + +def check_get_preferred_vector_width_char(device): + try: + device.preferred_vector_width_char + except Exception: + pytest.fail("preferred_vector_width_char call failed") + + +def check_get_preferred_vector_width_short(device): + try: + device.preferred_vector_width_short + except Exception: + pytest.fail("preferred_vector_width_short call failed") + + +def check_get_preferred_vector_width_int(device): + try: + device.preferred_vector_width_int + except Exception: + pytest.fail("preferred_vector_width_int call failed") + + +def check_get_preferred_vector_width_long(device): + try: + device.preferred_vector_width_long + except Exception: + pytest.fail("preferred_vector_width_long call failed") + + +def check_get_preferred_vector_width_float(device): + try: + device.preferred_vector_width_float + except Exception: + pytest.fail("preferred_vector_width_float call failed") + + +def check_get_preferred_vector_width_double(device): + try: + device.preferred_vector_width_double + except Exception: + pytest.fail("preferred_vector_width_double call failed") + + +def check_get_preferred_vector_width_half(device): + try: + device.preferred_vector_width_half + except Exception: + pytest.fail("preferred_vector_width_half call failed") + + list_of_checks = [ check_get_max_compute_units, check_get_max_work_item_dims, @@ -250,6 +306,14 @@ def check_is_host(device): check_is_cpu, check_is_gpu, check_is_host, + check_get_sub_group_independent_forward_progress, + check_get_preferred_vector_width_char, + check_get_preferred_vector_width_short, + check_get_preferred_vector_width_int, + check_get_preferred_vector_width_long, + check_get_preferred_vector_width_float, + check_get_preferred_vector_width_double, + check_get_preferred_vector_width_half, check_has_aspect_host, check_has_aspect_cpu, check_has_aspect_gpu,