Skip to content

Add device descriptors: sub_group_independent_forward_progress and preferred_vector_width #308

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 9 commits into from
Mar 25, 2021
Merged

Add device descriptors: sub_group_independent_forward_progress and preferred_vector_width #308

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
8507a50
Add sub_group_independent_forward_progress and preferred_vector_width…
Feb 26, 2021
9842db5
Merge master
Mar 2, 2021
952cb18
Add tests
Mar 2, 2021
e153c76
Add python api
Mar 2, 2021
7baa7a1
Fixes + tests
Mar 3, 2021
c857ecc
Merge master
Mar 10, 2021
68620f6
Add preferred_vector_width funcs
Mar 10, 2021
80bf1f0
Merge master
Mar 25, 2021
c0106e1
Added property
Mar 25, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
96 changes: 96 additions & 0 deletions dpctl-capi/include/dpctl_sycl_device_interface.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
136 changes: 136 additions & 0 deletions dpctl-capi/source/dpctl_sycl_device_interface.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}
123 changes: 123 additions & 0 deletions dpctl-capi/tests/test_sycl_device_interface.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down Expand Up @@ -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(
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nice! You have been reading the spec a lot :)

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",
Expand Down
8 changes: 8 additions & 0 deletions dpctl/_backend.pxd
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)

Expand Down
Loading