diff --git a/sycl/include/CL/sycl/aspects.hpp b/sycl/include/CL/sycl/aspects.hpp index 621f2764ba020..8d06ada94aa8c 100644 --- a/sycl/include/CL/sycl/aspects.hpp +++ b/sycl/include/CL/sycl/aspects.hpp @@ -37,7 +37,8 @@ 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 }; } // namespace sycl diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index 121af49969ab4..714afefedc281 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -517,9 +517,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_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/detail/property_helper.hpp b/sycl/include/CL/sycl/detail/property_helper.hpp index 3c8a82bc00117..6c0f8d71ef40f 100644 --- a/sycl/include/CL/sycl/detail/property_helper.hpp +++ b/sycl/include/CL/sycl/detail/property_helper.hpp @@ -25,22 +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, - PropWithDataKindSize + BufferContextBound = 1, + ImageUseMutex = 2, + ImageContextBound = 3, + BufferMemChannel = 4, + PropWithDataKindSize = 5 }; // Base class for dataless properties, needed to check that the type of an diff --git a/sycl/include/CL/sycl/info/device_traits.def b/sycl/include/CL/sycl/info/device_traits.def index e66990331f873..f3d4dd6b1d1b3 100644 --- a/sycl/include/CL/sycl/info/device_traits.def +++ b/sycl/include/CL/sycl/info/device_traits.def @@ -92,3 +92,4 @@ __SYCL_PARAM_TRAITS_SPEC(device, ext_intel_gpu_slices, pi_uint32) __SYCL_PARAM_TRAITS_SPEC(device, ext_intel_gpu_subslices_per_slice, pi_uint32) __SYCL_PARAM_TRAITS_SPEC(device, ext_intel_gpu_eu_count_per_subslice, pi_uint32) __SYCL_PARAM_TRAITS_SPEC(device, ext_intel_max_mem_bandwidth, pi_uint64) +__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 215c7e5be8292..f221638c9af99 100644 --- a/sycl/include/CL/sycl/info/info_desc.hpp +++ b/sycl/include/CL/sycl/info/info_desc.hpp @@ -131,6 +131,7 @@ enum class device : cl_device_info { 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, + // intel extensions ext_intel_pci_address = PI_DEVICE_INFO_PCI_ADDRESS, ext_intel_gpu_eu_count = PI_DEVICE_INFO_GPU_EU_COUNT, @@ -139,7 +140,8 @@ 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 }; 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..a8f0353227b05 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(uint32_t Channel) : MChannel(Channel) {} + uint32_t get_channel() const { return MChannel; } + +private: + uint32_t MChannel; +}; + } // namespace buffer } // namespace property diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 8ad8fba030d72..c84d1e2d1616d 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -17,11 +17,13 @@ #include #include +#include #include #include #include #include #include +#include #include #include @@ -546,22 +548,25 @@ 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) { + // TODO: need to check if all properties are supported by OpenCL RT and + // ignore unsupported + clCreateBufferWithPropertiesINTEL_fn FuncPtr = nullptr; // First we need to look up the function pointer ret_err = getExtFuncFromContext( context, &FuncPtr); + if (FuncPtr) { + *ret_mem = cast(FuncPtr(cast(context), properties, + 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 4942430a4859d..a5952ca60350b 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_PROPERTIES_CHANNEL); + } +}; + // Specializations for intel extensions for Level Zero low-level // detail device descriptors (not support on host). template <> diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index bc6b89bd5b3d8..921652e70037b 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 _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE65568EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE65569EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE65570EEENS3_12param_traitsIS4_XT_EE11return_typeEv 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); {