Add device descriptors: sub_group_independent_forward_progress and preferred_vector_width (#308)

1e-to · etotmeni · web-flow · commit 75aa0c1d1a21 · 2021-03-25T15:29:17.000-05:00
* Add sub_group_independent_forward_progress and preferred_vector_width_char

* Add tests

* Add python api

* Fixes + tests

* Add preferred_vector_width funcs

* Added property

Co-authored-by: etotmeni &lt;elena.totmenina@intel.com&gt;
diff --git 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<info::device::sub_group_independent_forward_progress>.
+ *
+ * @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<info::device::preferred_vector_width_char>.
+ *
+ * @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<info::device::preferred_vector_width_short>.
+ *
+ * @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<info::device::preferred_vector_width_int>.
+ *
+ * @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<info::device::preferred_vector_width_long>.
+ *
+ * @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<info::device::preferred_vector_width_float>.
+ *
+ * @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<info::device::preferred_vector_width_double>.
+ *
+ * @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<info::device::preferred_vector_width_half>.
+ *
+ * @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
@@ -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<info::device::preferred_vector_width_char>();
+        } 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<info::device::preferred_vector_width_short>();
+        } 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<info::device::preferred_vector_width_int>();
+        } 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<info::device::preferred_vector_width_long>();
+        } 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<info::device::preferred_vector_width_float>();
+        } 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<info::device::preferred_vector_width_double>();
+        } 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<info::device::preferred_vector_width_half>();
+        } 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
@@ -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<device *>(DRef);
+    auto get_sub_group_progress =
+        D->get_info<info::device::sub_group_independent_forward_progress>();
+    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
@@ -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
diff --git a/dpctl/tests/test_sycl_device.py b/dpctl/tests/test_sycl_device.py
