From aff5fee3c32bb033b6a93a5b691e56aba2df0f8e Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Wed, 28 Oct 2020 22:29:57 +0300 Subject: [PATCH 01/10] [SYCL] Implement SYCL_INTEL_mem_channel_property extension On some targets manual assignment of buffers to memory regions can improve memory bandwidth. This extension adds a buffer property to indicate in which memory channel a particular buffer should be allocated. This information is an optimization hint to the runtime and thus it is legal to ignore. Spec: https://github.com/intel/llvm/pull/2688 Signed-off-by: Dmitry Sidorov --- sycl/include/CL/sycl/aspects.hpp | 3 +- sycl/include/CL/sycl/detail/pi.h | 3 +- .../CL/sycl/detail/property_helper.hpp | 1 + sycl/include/CL/sycl/info/device_traits.def | 1 + sycl/include/CL/sycl/info/info_desc.hpp | 4 +- .../CL/sycl/properties/buffer_properties.hpp | 11 ++++ sycl/plugins/opencl/pi_opencl.cpp | 62 +++++++++++++++---- sycl/source/detail/device_info.hpp | 16 +++++ sycl/test/abi/sycl_symbols_linux.dump | 1 + sycl/test/basic_tests/property_list.cpp | 17 +++++ .../on-device/basic_tests/buffer/buffer.cpp | 19 ++++++ 11 files changed, 123 insertions(+), 15 deletions(-) diff --git a/sycl/include/CL/sycl/aspects.hpp b/sycl/include/CL/sycl/aspects.hpp index 4b66fcab94a79..1325f5ee13639 100644 --- a/sycl/include/CL/sycl/aspects.hpp +++ b/sycl/include/CL/sycl/aspects.hpp @@ -30,7 +30,8 @@ enum class aspect { usm_host_allocations, usm_shared_allocations, usm_restricted_shared_allocations, - usm_system_allocator + usm_system_allocator, + ext_intel_mem_channel }; } // namespace sycl diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index e0f16833e0eb7..40eefa0f70784 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -503,9 +503,8 @@ constexpr pi_map_flags PI_MAP_WRITE_INVALIDATE_REGION = // NOTE: this is made 64-bit to match the size of cl_mem_properties_intel to // make the translation to OpenCL transparent. -// TODO: populate -// using pi_mem_properties = pi_bitfield; +constexpr pi_mem_properties PI_MEM_CHANNEL_INTEL = CL_MEM_CHANNEL_INTEL; // NOTE: queue properties are implemented this way to better support bit // manipulations diff --git a/sycl/include/CL/sycl/detail/property_helper.hpp b/sycl/include/CL/sycl/detail/property_helper.hpp index 3c8a82bc00117..e55f23e856a70 100644 --- a/sycl/include/CL/sycl/detail/property_helper.hpp +++ b/sycl/include/CL/sycl/detail/property_helper.hpp @@ -40,6 +40,7 @@ enum PropWithDataKind { BufferContextBound, ImageUseMutex, ImageContextBound, + BufferMemChannel, PropWithDataKindSize }; diff --git a/sycl/include/CL/sycl/info/device_traits.def b/sycl/include/CL/sycl/info/device_traits.def index 6ac01d748c137..0526e939b0d16 100644 --- a/sycl/include/CL/sycl/info/device_traits.def +++ b/sycl/include/CL/sycl/info/device_traits.def @@ -85,3 +85,4 @@ __SYCL_PARAM_TRAITS_SPEC(device, usm_host_allocations, bool) __SYCL_PARAM_TRAITS_SPEC(device, usm_shared_allocations, bool) __SYCL_PARAM_TRAITS_SPEC(device, usm_restricted_shared_allocations, bool) __SYCL_PARAM_TRAITS_SPEC(device, usm_system_allocator, bool) +__SYCL_PARAM_TRAITS_SPEC(device, ext_intel_mem_channel, bool) diff --git a/sycl/include/CL/sycl/info/info_desc.hpp b/sycl/include/CL/sycl/info/info_desc.hpp index a1c368cdc5519..8bc456fb9d20b 100644 --- a/sycl/include/CL/sycl/info/info_desc.hpp +++ b/sycl/include/CL/sycl/info/info_desc.hpp @@ -130,7 +130,9 @@ enum class device : cl_device_info { usm_host_allocations = PI_USM_HOST_SUPPORT, usm_shared_allocations = PI_USM_SINGLE_SHARED_SUPPORT, usm_restricted_shared_allocations = PI_USM_CROSS_SHARED_SUPPORT, - usm_system_allocator = PI_USM_SYSTEM_SHARED_SUPPORT + usm_system_allocator = PI_USM_SYSTEM_SHARED_SUPPORT, + + ext_intel_mem_channel = PI_MEM_CHANNEL_INTEL }; enum class device_type : pi_uint64 { diff --git a/sycl/include/CL/sycl/properties/buffer_properties.hpp b/sycl/include/CL/sycl/properties/buffer_properties.hpp index 8905385cdc9f4..1fe593ef67a39 100644 --- a/sycl/include/CL/sycl/properties/buffer_properties.hpp +++ b/sycl/include/CL/sycl/properties/buffer_properties.hpp @@ -39,6 +39,17 @@ class context_bound private: sycl::context MCtx; }; + +class mem_channel : public detail::PropertyWithData< + detail::PropWithDataKind::BufferMemChannel> { +public: + mem_channel(cl_uint Channel) : m_Channel(Channel) {} + cl_uint get_channel() const { return m_Channel; } + +private: + cl_uint m_Channel; +}; + } // namespace buffer } // namespace property diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index ac15e9334e04a..ce72304882b53 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -523,22 +523,62 @@ pi_result piMemBufferCreate(pi_context context, pi_mem_flags flags, size_t size, void *host_ptr, pi_mem *ret_mem, const pi_mem_properties *properties) { pi_result ret_err = PI_INVALID_OPERATION; - clCreateBufferWithPropertiesINTEL_fn FuncPtr = nullptr; - - if (properties) + if (properties) { + clCreateBufferWithPropertiesINTEL_fn FuncPtr = nullptr; + const size_t propSize = sizeof(properties) / sizeof(pi_mem_properties); // First we need to look up the function pointer ret_err = getExtFuncFromContext( context, &FuncPtr); + if (FuncPtr) { + std::vector supported(properties, + properties + propSize); + // Go through buffer properties. If there is one, that shall be propagated + // to an OpenCL runtime - check if this property is being supported. + for (auto prop = supported.begin(); prop != supported.end(); ++prop) { + // Check if PI_MEM_CHANNEL_INTEL property is supported. If it's not - + // just ignore it, as it's an optimization hint. + if (*prop == PI_MEM_CHANNEL_INTEL) { + size_t deviceCount; + cl_int ret_err = + clGetContextInfo(cast(context), CL_CONTEXT_DEVICES, 0, + nullptr, &deviceCount); + if (ret_err != CL_SUCCESS || deviceCount < 1) + return PI_INVALID_CONTEXT; + std::vector devicesInCtx(deviceCount); + ret_err = clGetContextInfo( + cast(context), CL_CONTEXT_DEVICES, + deviceCount * sizeof(cl_device_id), devicesInCtx.data(), nullptr); + + size_t retSize; + ret_err = clGetDeviceInfo(devicesInCtx[0], CL_DEVICE_EXTENSIONS, 0, + nullptr, &retSize); + if (ret_err != CL_SUCCESS) + return PI_INVALID_DEVICE; + std::string extensions(retSize, '\0'); + ret_err = clGetDeviceInfo(devicesInCtx[0], CL_DEVICE_EXTENSIONS, + retSize, &extensions[0], nullptr); + if (ret_err != CL_SUCCESS) + return PI_INVALID_DEVICE; + + size_t pos = extensions.find("cl_intel_mem_channel_property"); + if (pos == std::string::npos) + supported.erase(prop); + } + } + if (!supported.empty()) { + *ret_mem = + cast(FuncPtr(cast(context), supported.data(), + cast(flags), size, host_ptr, + cast(&ret_err))); + return ret_err; + } + } + } - if (FuncPtr) - *ret_mem = cast(FuncPtr(cast(context), properties, - cast(flags), size, host_ptr, - cast(&ret_err))); - else - *ret_mem = cast(clCreateBuffer(cast(context), - cast(flags), size, - host_ptr, cast(&ret_err))); + *ret_mem = cast(clCreateBuffer(cast(context), + cast(flags), size, + host_ptr, cast(&ret_err))); return ret_err; } diff --git a/sycl/source/detail/device_info.hpp b/sycl/source/detail/device_info.hpp index 4fc8be379bb0e..0c35698cdbc61 100644 --- a/sycl/source/detail/device_info.hpp +++ b/sycl/source/detail/device_info.hpp @@ -936,6 +936,11 @@ inline bool get_device_info_host() { return true; } +template <> +inline bool get_device_info_host() { + return false; +} + cl_uint get_native_vector_width(size_t idx); // USM @@ -1003,6 +1008,17 @@ template <> struct get_device_info { } }; +// Specialization for memory channel query +template <> struct get_device_info { + static bool get(RT::PiDevice dev, const plugin &Plugin) { + pi_mem_properties caps; + pi_result Err = Plugin.call_nocheck( + dev, pi::cast(info::device::ext_intel_mem_channel), + sizeof(pi_mem_properties), &caps, nullptr); + return (Err != PI_SUCCESS) ? false : (caps & PI_MEM_CHANNEL_INTEL); + } +}; + } // namespace detail } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index e8677277fca0f..755c07ab29a73 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -4036,6 +4036,7 @@ _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE4168EEENS3_12param_traitsIS4_XT_E _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE4169EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE4188EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE4189EEENS3_12param_traitsIS4_XT_EE11return_typeEv +_ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE16915EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl6device9getNativeEv _ZNK2cl4sycl6kernel11get_contextEv _ZNK2cl4sycl6kernel11get_programEv diff --git a/sycl/test/basic_tests/property_list.cpp b/sycl/test/basic_tests/property_list.cpp index 0150fb92daabc..2dc390a521734 100644 --- a/sycl/test/basic_tests/property_list.cpp +++ b/sycl/test/basic_tests/property_list.cpp @@ -64,6 +64,23 @@ int main() { } } + { + cl::sycl::property_list MemChannelProp{ + sycl_property::buffer::mem_channel(2)}; + if (!MemChannelProp.has_property()) { + std::cerr << "Error: property list has no property while should have." + << std::endl; + Failed = true; + } + auto Prop = + MemChannelProp.get_property(); + if (Prop.get_channel() != 2) { + std::cerr << "Error: mem_channel property is not equal to 2." + << std::endl; + Failed = true; + } + } + std::cerr << "Test status : " << (Failed ? "FAILED" : "PASSED") << std::endl; return Failed; diff --git a/sycl/test/on-device/basic_tests/buffer/buffer.cpp b/sycl/test/on-device/basic_tests/buffer/buffer.cpp index ce357f188c9ff..de74d3f8760ea 100644 --- a/sycl/test/on-device/basic_tests/buffer/buffer.cpp +++ b/sycl/test/on-device/basic_tests/buffer/buffer.cpp @@ -40,6 +40,25 @@ int main() { assert(data1[i] == 0); } + { + int data1[10] = {-1, -1, -1, -1, -1, -1, -1, -1, -1, -1}; + { + buffer b(data1, range<1>(10), {property::buffer::mem_channel{3}}); + queue myQueue; + myQueue.submit([&](handler &cgh) { + auto B = b.get_access(cgh); + cgh.parallel_for(range<1>{10}, + [=](id<1> index) { B[index] = 0; }); + }); + assert(b.has_property()); + auto prop = b.get_property(); + assert(prop.get_channel() == 3 && "oops it's not 3"); + + } // Data is copied back because there is a user side shared_ptr + for (int i = 0; i < 10; i++) + assert(data1[i] == 0); + } + { std::vector data1(10, -1); { From 9d888a1b1e075ef25bc87cea285ee7a29dc484eb Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Wed, 25 Nov 2020 14:07:05 +0300 Subject: [PATCH 02/10] Make properties enum numbering be explicit Signed-off-by: Dmitry Sidorov --- .../CL/sycl/detail/property_helper.hpp | 24 +++++++++---------- 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/sycl/include/CL/sycl/detail/property_helper.hpp b/sycl/include/CL/sycl/detail/property_helper.hpp index e55f23e856a70..6c0f8d71ef40f 100644 --- a/sycl/include/CL/sycl/detail/property_helper.hpp +++ b/sycl/include/CL/sycl/detail/property_helper.hpp @@ -25,23 +25,23 @@ namespace detail { // List of all dataless properties' IDs enum DataLessPropKind { BufferUseHostPtr = 0, - ImageUseHostPtr, - QueueEnableProfiling, - InOrder, - NoInit, - BufferUsePinnedHostMemory, - UsePrimaryContext, - DataLessPropKindSize + ImageUseHostPtr = 1, + QueueEnableProfiling = 2, + InOrder = 3, + NoInit = 4, + BufferUsePinnedHostMemory = 5, + UsePrimaryContext = 6, + DataLessPropKindSize = 7 }; // List of all properties with data IDs enum PropWithDataKind { BufferUseMutex = 0, - BufferContextBound, - ImageUseMutex, - ImageContextBound, - BufferMemChannel, - PropWithDataKindSize + BufferContextBound = 1, + ImageUseMutex = 2, + ImageContextBound = 3, + BufferMemChannel = 4, + PropWithDataKindSize = 5 }; // Base class for dataless properties, needed to check that the type of an From e876fd139a92e09be8ea356c005f042183edbc0b Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Mon, 14 Dec 2020 15:53:30 +0300 Subject: [PATCH 03/10] Apply most of the comments Signed-off-by: Dmitry Sidorov --- sycl/include/CL/sycl/detail/pi.h | 2 +- sycl/include/CL/sycl/info/info_desc.hpp | 2 +- sycl/plugins/opencl/pi_opencl.cpp | 77 +++++++++++++++---------- sycl/source/detail/device_info.hpp | 2 +- 4 files changed, 51 insertions(+), 32 deletions(-) diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index 40eefa0f70784..d84cb815bee7f 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -504,7 +504,7 @@ constexpr pi_map_flags PI_MAP_WRITE_INVALIDATE_REGION = // NOTE: this is made 64-bit to match the size of cl_mem_properties_intel to // make the translation to OpenCL transparent. using pi_mem_properties = pi_bitfield; -constexpr pi_mem_properties PI_MEM_CHANNEL_INTEL = CL_MEM_CHANNEL_INTEL; +constexpr pi_mem_properties PI_MEM_PROPERTIES_CHANNEL = CL_MEM_CHANNEL_INTEL; // NOTE: queue properties are implemented this way to better support bit // manipulations diff --git a/sycl/include/CL/sycl/info/info_desc.hpp b/sycl/include/CL/sycl/info/info_desc.hpp index 8bc456fb9d20b..eeb860181e8da 100644 --- a/sycl/include/CL/sycl/info/info_desc.hpp +++ b/sycl/include/CL/sycl/info/info_desc.hpp @@ -132,7 +132,7 @@ enum class device : cl_device_info { usm_restricted_shared_allocations = PI_USM_CROSS_SHARED_SUPPORT, usm_system_allocator = PI_USM_SYSTEM_SHARED_SUPPORT, - ext_intel_mem_channel = PI_MEM_CHANNEL_INTEL + ext_intel_mem_channel = PI_MEM_PROPERTIES_CHANNEL }; enum class device_type : pi_uint64 { diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index ce72304882b53..09db37420290d 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -17,11 +17,14 @@ #include #include +#include #include #include #include #include #include +#include +#include #include #include @@ -33,6 +36,7 @@ } const char SupportedVersion[] = _PI_H_VERSION_STRING; +std::set SupportedExtensions; // Want all the needed casts be explicit, do not define conversion operators. template To cast(From value) { @@ -66,6 +70,39 @@ CONSTFIX char clSetProgramSpecializationConstantName[] = #undef CONSTFIX +// Helper to get extensions that are common for all devices within a context +pi_result getSupportedExtensionsWithinContext(pi_context context) { + size_t deviceCount; + cl_int ret_err = + clGetContextInfo(cast(context), CL_CONTEXT_DEVICES, 0, + nullptr, &deviceCount); + if (ret_err != CL_SUCCESS || deviceCount < 1) + return PI_INVALID_CONTEXT; + std::vector devicesInCtx(deviceCount); + ret_err = clGetContextInfo( + cast(context), CL_CONTEXT_DEVICES, + deviceCount * sizeof(cl_device_id), devicesInCtx.data(), nullptr); + + size_t retSize; + for (size_t i = 0; i != deviceCount; ++i) { + ret_err = clGetDeviceInfo(devicesInCtx[i], CL_DEVICE_EXTENSIONS, 0, + nullptr, &retSize); + if (ret_err != CL_SUCCESS) + return PI_INVALID_DEVICE; + std::string extensions(retSize, '\0'); + ret_err = clGetDeviceInfo(devicesInCtx[i], CL_DEVICE_EXTENSIONS, + retSize, &extensions[0], nullptr); + if (ret_err != CL_SUCCESS) + return PI_INVALID_DEVICE; + std::string extension; + std::stringstream ss(extensions); + while (getline(ss, extension, ' ')) + SupportedExtensions.insert(extension); + } + return cast(ret_err); +} + + // USM helper function to get an extension function pointer template static pi_result getExtFuncFromContext(pi_context context, T *fptr) { @@ -535,36 +572,18 @@ pi_result piMemBufferCreate(pi_context context, pi_mem_flags flags, size_t size, properties + propSize); // Go through buffer properties. If there is one, that shall be propagated // to an OpenCL runtime - check if this property is being supported. - for (auto prop = supported.begin(); prop != supported.end(); ++prop) { - // Check if PI_MEM_CHANNEL_INTEL property is supported. If it's not - - // just ignore it, as it's an optimization hint. - if (*prop == PI_MEM_CHANNEL_INTEL) { - size_t deviceCount; - cl_int ret_err = - clGetContextInfo(cast(context), CL_CONTEXT_DEVICES, 0, - nullptr, &deviceCount); - if (ret_err != CL_SUCCESS || deviceCount < 1) - return PI_INVALID_CONTEXT; - std::vector devicesInCtx(deviceCount); - ret_err = clGetContextInfo( - cast(context), CL_CONTEXT_DEVICES, - deviceCount * sizeof(cl_device_id), devicesInCtx.data(), nullptr); - - size_t retSize; - ret_err = clGetDeviceInfo(devicesInCtx[0], CL_DEVICE_EXTENSIONS, 0, - nullptr, &retSize); - if (ret_err != CL_SUCCESS) - return PI_INVALID_DEVICE; - std::string extensions(retSize, '\0'); - ret_err = clGetDeviceInfo(devicesInCtx[0], CL_DEVICE_EXTENSIONS, - retSize, &extensions[0], nullptr); - if (ret_err != CL_SUCCESS) - return PI_INVALID_DEVICE; - - size_t pos = extensions.find("cl_intel_mem_channel_property"); - if (pos == std::string::npos) + for (const auto &prop = supported.begin(); prop != supported.end(); + ++(*prop)) { + if (!SupportedExtensions.empty()) + ret_err = getSupportedExtensionsWithinContext(context); + // Check if PI_MEM_PROPERTIES_CHANNEL property is supported. If it's + // not - just ignore it, as it's an optimization hint. + if (*prop == PI_MEM_PROPERTIES_CHANNEL) { + if (SupportedExtensions.find("cl_intel_mem_channel_property") != + SupportedExtensions.end()) supported.erase(prop); - } + } else + assert("Unsupported property found"); } if (!supported.empty()) { *ret_mem = diff --git a/sycl/source/detail/device_info.hpp b/sycl/source/detail/device_info.hpp index 0c35698cdbc61..82e2a8833e213 100644 --- a/sycl/source/detail/device_info.hpp +++ b/sycl/source/detail/device_info.hpp @@ -1015,7 +1015,7 @@ template <> struct get_device_info { pi_result Err = Plugin.call_nocheck( dev, pi::cast(info::device::ext_intel_mem_channel), sizeof(pi_mem_properties), &caps, nullptr); - return (Err != PI_SUCCESS) ? false : (caps & PI_MEM_CHANNEL_INTEL); + return (Err != PI_SUCCESS) ? false : (caps & PI_MEM_PROPERTIES_CHANNEL); } }; From 950084eda9da9853fd8da190f450a3c004f7e3d2 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Mon, 14 Dec 2020 17:02:29 +0300 Subject: [PATCH 04/10] Fix clang-format Signed-off-by: Dmitry Sidorov --- sycl/plugins/opencl/pi_opencl.cpp | 20 +++++++++----------- 1 file changed, 9 insertions(+), 11 deletions(-) diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 09db37420290d..13b1531cc0a54 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -73,25 +73,24 @@ CONSTFIX char clSetProgramSpecializationConstantName[] = // Helper to get extensions that are common for all devices within a context pi_result getSupportedExtensionsWithinContext(pi_context context) { size_t deviceCount; - cl_int ret_err = - clGetContextInfo(cast(context), CL_CONTEXT_DEVICES, 0, - nullptr, &deviceCount); + cl_int ret_err = clGetContextInfo( + cast(context), CL_CONTEXT_DEVICES, 0, nullptr, &deviceCount); if (ret_err != CL_SUCCESS || deviceCount < 1) return PI_INVALID_CONTEXT; std::vector devicesInCtx(deviceCount); - ret_err = clGetContextInfo( - cast(context), CL_CONTEXT_DEVICES, - deviceCount * sizeof(cl_device_id), devicesInCtx.data(), nullptr); + ret_err = clGetContextInfo(cast(context), CL_CONTEXT_DEVICES, + deviceCount * sizeof(cl_device_id), + devicesInCtx.data(), nullptr); size_t retSize; for (size_t i = 0; i != deviceCount; ++i) { - ret_err = clGetDeviceInfo(devicesInCtx[i], CL_DEVICE_EXTENSIONS, 0, - nullptr, &retSize); + ret_err = clGetDeviceInfo(devicesInCtx[i], CL_DEVICE_EXTENSIONS, 0, nullptr, + &retSize); if (ret_err != CL_SUCCESS) return PI_INVALID_DEVICE; std::string extensions(retSize, '\0'); - ret_err = clGetDeviceInfo(devicesInCtx[i], CL_DEVICE_EXTENSIONS, - retSize, &extensions[0], nullptr); + ret_err = clGetDeviceInfo(devicesInCtx[i], CL_DEVICE_EXTENSIONS, retSize, + &extensions[0], nullptr); if (ret_err != CL_SUCCESS) return PI_INVALID_DEVICE; std::string extension; @@ -102,7 +101,6 @@ pi_result getSupportedExtensionsWithinContext(pi_context context) { return cast(ret_err); } - // USM helper function to get an extension function pointer template static pi_result getExtFuncFromContext(pi_context context, T *fptr) { From 4a3e3d4f7c555f9fc3217b67fe8f6f24503a8442 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Wed, 16 Dec 2020 17:14:44 +0300 Subject: [PATCH 05/10] cl_uint -> uint32_t Signed-off-by: Dmitry Sidorov --- sycl/include/CL/sycl/properties/buffer_properties.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/properties/buffer_properties.hpp b/sycl/include/CL/sycl/properties/buffer_properties.hpp index 1fe593ef67a39..8a5d898da280d 100644 --- a/sycl/include/CL/sycl/properties/buffer_properties.hpp +++ b/sycl/include/CL/sycl/properties/buffer_properties.hpp @@ -44,10 +44,10 @@ class mem_channel : public detail::PropertyWithData< detail::PropWithDataKind::BufferMemChannel> { public: mem_channel(cl_uint Channel) : m_Channel(Channel) {} - cl_uint get_channel() const { return m_Channel; } + uint32_t get_channel() const { return m_Channel; } private: - cl_uint m_Channel; + uint32_t m_Channel; }; } // namespace buffer From 541a905eadd8c662351c1452d9bf87ae210145cf Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Wed, 16 Dec 2020 17:15:58 +0300 Subject: [PATCH 06/10] One more Signed-off-by: Dmitry Sidorov --- sycl/include/CL/sycl/properties/buffer_properties.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/properties/buffer_properties.hpp b/sycl/include/CL/sycl/properties/buffer_properties.hpp index 8a5d898da280d..139e195557ef6 100644 --- a/sycl/include/CL/sycl/properties/buffer_properties.hpp +++ b/sycl/include/CL/sycl/properties/buffer_properties.hpp @@ -43,7 +43,7 @@ class context_bound class mem_channel : public detail::PropertyWithData< detail::PropWithDataKind::BufferMemChannel> { public: - mem_channel(cl_uint Channel) : m_Channel(Channel) {} + mem_channel(uint32_t Channel) : m_Channel(Channel) {} uint32_t get_channel() const { return m_Channel; } private: From 8f78d1fdf2d7477437b5b5e75ce025b7de2ca7b7 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Mon, 21 Dec 2020 14:50:55 +0300 Subject: [PATCH 07/10] Set -> Map Signed-off-by: Dmitry Sidorov --- sycl/plugins/opencl/pi_opencl.cpp | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 13b1531cc0a54..d88c39e766b8d 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -36,7 +36,7 @@ } const char SupportedVersion[] = _PI_H_VERSION_STRING; -std::set SupportedExtensions; +std::map> SupportedExtensions; // Want all the needed casts be explicit, do not define conversion operators. template To cast(From value) { @@ -83,6 +83,7 @@ pi_result getSupportedExtensionsWithinContext(pi_context context) { devicesInCtx.data(), nullptr); size_t retSize; + std::set commonExtensions; for (size_t i = 0; i != deviceCount; ++i) { ret_err = clGetDeviceInfo(devicesInCtx[i], CL_DEVICE_EXTENSIONS, 0, nullptr, &retSize); @@ -96,8 +97,9 @@ pi_result getSupportedExtensionsWithinContext(pi_context context) { std::string extension; std::stringstream ss(extensions); while (getline(ss, extension, ' ')) - SupportedExtensions.insert(extension); + commonExtensions.insert(extension); } + SupportedExtensions.emplace(context, commonExtensions); return cast(ret_err); } @@ -572,13 +574,14 @@ pi_result piMemBufferCreate(pi_context context, pi_mem_flags flags, size_t size, // to an OpenCL runtime - check if this property is being supported. for (const auto &prop = supported.begin(); prop != supported.end(); ++(*prop)) { - if (!SupportedExtensions.empty()) + if (SupportedExtensions.find(context) == SupportedExtensions.end()) ret_err = getSupportedExtensionsWithinContext(context); // Check if PI_MEM_PROPERTIES_CHANNEL property is supported. If it's // not - just ignore it, as it's an optimization hint. if (*prop == PI_MEM_PROPERTIES_CHANNEL) { - if (SupportedExtensions.find("cl_intel_mem_channel_property") != - SupportedExtensions.end()) + if (SupportedExtensions[context].find( + "cl_intel_mem_channel_property") != + SupportedExtensions[context].end()) supported.erase(prop); } else assert("Unsupported property found"); From 4b71b32fe5c89766509d3f770a167772c9672896 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Tue, 22 Dec 2020 11:46:27 +0300 Subject: [PATCH 08/10] Apply several comments Signed-off-by: Dmitry Sidorov --- sycl/include/CL/sycl/properties/buffer_properties.hpp | 6 +++--- sycl/plugins/opencl/pi_opencl.cpp | 9 ++++----- 2 files changed, 7 insertions(+), 8 deletions(-) diff --git a/sycl/include/CL/sycl/properties/buffer_properties.hpp b/sycl/include/CL/sycl/properties/buffer_properties.hpp index 139e195557ef6..a8f0353227b05 100644 --- a/sycl/include/CL/sycl/properties/buffer_properties.hpp +++ b/sycl/include/CL/sycl/properties/buffer_properties.hpp @@ -43,11 +43,11 @@ class context_bound class mem_channel : public detail::PropertyWithData< detail::PropWithDataKind::BufferMemChannel> { public: - mem_channel(uint32_t Channel) : m_Channel(Channel) {} - uint32_t get_channel() const { return m_Channel; } + mem_channel(uint32_t Channel) : MChannel(Channel) {} + uint32_t get_channel() const { return MChannel; } private: - uint32_t m_Channel; + uint32_t MChannel; }; } // namespace buffer diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index d88c39e766b8d..e2199f5c083c0 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -572,19 +572,18 @@ pi_result piMemBufferCreate(pi_context context, pi_mem_flags flags, size_t size, properties + propSize); // Go through buffer properties. If there is one, that shall be propagated // to an OpenCL runtime - check if this property is being supported. - for (const auto &prop = supported.begin(); prop != supported.end(); - ++(*prop)) { + for (auto prop = supported.begin(); prop != supported.end(); ++prop) { if (SupportedExtensions.find(context) == SupportedExtensions.end()) ret_err = getSupportedExtensionsWithinContext(context); // Check if PI_MEM_PROPERTIES_CHANNEL property is supported. If it's // not - just ignore it, as it's an optimization hint. if (*prop == PI_MEM_PROPERTIES_CHANNEL) { if (SupportedExtensions[context].find( - "cl_intel_mem_channel_property") != + "cl_intel_mem_channel_property") == SupportedExtensions[context].end()) - supported.erase(prop); + prop = supported.erase(prop); } else - assert("Unsupported property found"); + assert(!"Unsupported property found"); } if (!supported.empty()) { *ret_mem = From 1f1135b48f1b3f28b8709902411e53d09fe2b848 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Tue, 22 Dec 2020 12:09:41 +0300 Subject: [PATCH 09/10] Remove properties check Signed-off-by: Dmitry Sidorov --- sycl/plugins/opencl/pi_opencl.cpp | 66 +++---------------------------- 1 file changed, 6 insertions(+), 60 deletions(-) diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index e2199f5c083c0..d4af46cfc1f0a 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -23,7 +23,6 @@ #include #include #include -#include #include #include #include @@ -36,7 +35,6 @@ } const char SupportedVersion[] = _PI_H_VERSION_STRING; -std::map> SupportedExtensions; // Want all the needed casts be explicit, do not define conversion operators. template To cast(From value) { @@ -70,39 +68,6 @@ CONSTFIX char clSetProgramSpecializationConstantName[] = #undef CONSTFIX -// Helper to get extensions that are common for all devices within a context -pi_result getSupportedExtensionsWithinContext(pi_context context) { - size_t deviceCount; - cl_int ret_err = clGetContextInfo( - cast(context), CL_CONTEXT_DEVICES, 0, nullptr, &deviceCount); - if (ret_err != CL_SUCCESS || deviceCount < 1) - return PI_INVALID_CONTEXT; - std::vector devicesInCtx(deviceCount); - ret_err = clGetContextInfo(cast(context), CL_CONTEXT_DEVICES, - deviceCount * sizeof(cl_device_id), - devicesInCtx.data(), nullptr); - - size_t retSize; - std::set commonExtensions; - for (size_t i = 0; i != deviceCount; ++i) { - ret_err = clGetDeviceInfo(devicesInCtx[i], CL_DEVICE_EXTENSIONS, 0, nullptr, - &retSize); - if (ret_err != CL_SUCCESS) - return PI_INVALID_DEVICE; - std::string extensions(retSize, '\0'); - ret_err = clGetDeviceInfo(devicesInCtx[i], CL_DEVICE_EXTENSIONS, retSize, - &extensions[0], nullptr); - if (ret_err != CL_SUCCESS) - return PI_INVALID_DEVICE; - std::string extension; - std::stringstream ss(extensions); - while (getline(ss, extension, ' ')) - commonExtensions.insert(extension); - } - SupportedExtensions.emplace(context, commonExtensions); - return cast(ret_err); -} - // USM helper function to get an extension function pointer template static pi_result getExtFuncFromContext(pi_context context, T *fptr) { @@ -561,37 +526,18 @@ pi_result piMemBufferCreate(pi_context context, pi_mem_flags flags, size_t size, const pi_mem_properties *properties) { pi_result ret_err = PI_INVALID_OPERATION; if (properties) { + // TODO: need to check if all properties are supported by OpenCL RT and + // ignore unsupported clCreateBufferWithPropertiesINTEL_fn FuncPtr = nullptr; - const size_t propSize = sizeof(properties) / sizeof(pi_mem_properties); // First we need to look up the function pointer ret_err = getExtFuncFromContext( context, &FuncPtr); if (FuncPtr) { - std::vector supported(properties, - properties + propSize); - // Go through buffer properties. If there is one, that shall be propagated - // to an OpenCL runtime - check if this property is being supported. - for (auto prop = supported.begin(); prop != supported.end(); ++prop) { - if (SupportedExtensions.find(context) == SupportedExtensions.end()) - ret_err = getSupportedExtensionsWithinContext(context); - // Check if PI_MEM_PROPERTIES_CHANNEL property is supported. If it's - // not - just ignore it, as it's an optimization hint. - if (*prop == PI_MEM_PROPERTIES_CHANNEL) { - if (SupportedExtensions[context].find( - "cl_intel_mem_channel_property") == - SupportedExtensions[context].end()) - prop = supported.erase(prop); - } else - assert(!"Unsupported property found"); - } - if (!supported.empty()) { - *ret_mem = - cast(FuncPtr(cast(context), supported.data(), - cast(flags), size, host_ptr, - cast(&ret_err))); - return ret_err; - } + *ret_mem = cast(FuncPtr(cast(context), properties, + cast(flags), size, host_ptr, + cast(&ret_err))); + return ret_err; } } From e29fe01e9e4f3108d3453e9c1bc266c00d0a4651 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Tue, 22 Dec 2020 19:57:18 +0300 Subject: [PATCH 10/10] Fix merge errors Signed-off-by: Dmitry Sidorov --- sycl/include/CL/sycl/aspects.hpp | 2 +- sycl/include/CL/sycl/info/info_desc.hpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/aspects.hpp b/sycl/include/CL/sycl/aspects.hpp index ae2635c1b0b83..8d06ada94aa8c 100644 --- a/sycl/include/CL/sycl/aspects.hpp +++ b/sycl/include/CL/sycl/aspects.hpp @@ -37,7 +37,7 @@ enum class aspect { ext_intel_gpu_slices, ext_intel_gpu_subslices_per_slice, ext_intel_gpu_eu_count_per_subslice, - ext_intel_max_mem_bandwidth + ext_intel_max_mem_bandwidth, ext_intel_mem_channel }; diff --git a/sycl/include/CL/sycl/info/info_desc.hpp b/sycl/include/CL/sycl/info/info_desc.hpp index c0af377ad6548..f221638c9af99 100644 --- a/sycl/include/CL/sycl/info/info_desc.hpp +++ b/sycl/include/CL/sycl/info/info_desc.hpp @@ -140,7 +140,7 @@ enum class device : cl_device_info { ext_intel_gpu_subslices_per_slice = PI_DEVICE_INFO_GPU_SUBSLICES_PER_SLICE, ext_intel_gpu_eu_count_per_subslice = PI_DEVICE_INFO_GPU_EU_COUNT_PER_SUBSLICE, - ext_intel_max_mem_bandwidth = PI_DEVICE_INFO_MAX_MEM_BANDWIDTH + ext_intel_max_mem_bandwidth = PI_DEVICE_INFO_MAX_MEM_BANDWIDTH, ext_intel_mem_channel = PI_MEM_PROPERTIES_CHANNEL };