From 75a201f24fa0d30f418dc3dfc8cfe81f9b89f684 Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Fri, 31 Mar 2023 01:10:16 -0700 Subject: [PATCH] Returns minimum mandated capabilities for atomic_fence device queries in CUDA backend. Signed-off-by: Maronas, Marcos --- sycl/plugins/cuda/pi_cuda.cpp | 27 +++++++++++++++++++++------ 1 file changed, 21 insertions(+), 6 deletions(-) diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 83a986f193d2f..c93409c9422ec 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -1303,12 +1303,27 @@ pi_result cuda_piDeviceGetInfo(pi_device device, pi_device_info param_name, return getInfo(param_value_size, param_value, param_value_size_ret, capabilities); } - case PI_EXT_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES: - case PI_EXT_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES: - // There is no way to query this in the backend - setErrorMessage("CUDA backend does not support this query", - PI_ERROR_INVALID_ARG_VALUE); - return PI_ERROR_PLUGIN_SPECIFIC_ERROR; + case PI_EXT_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES: { + // SYCL2020 4.6.4.2 minimum mandated capabilities for + // atomic_fence_order_capabilities. + pi_memory_order_capabilities capabilities = + PI_MEMORY_ORDER_RELAXED | PI_MEMORY_ORDER_ACQUIRE | + PI_MEMORY_ORDER_RELEASE | PI_MEMORY_ORDER_ACQ_REL; + return getInfo(param_value_size, param_value, param_value_size_ret, + capabilities); + } + case PI_EXT_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES: { + // SYCL2020 4.6.4.2 minimum mandated capabilities for + // atomic_fence/memory_scope_capabilities. + // Because scopes are hierarchical, wider scopes support all narrower + // scopes. At a minimum, each device must support WORK_ITEM, SUB_GROUP and + // WORK_GROUP. (https://github.com/KhronosGroup/SYCL-Docs/pull/382) + pi_memory_scope_capabilities capabilities = PI_MEMORY_SCOPE_WORK_ITEM | + PI_MEMORY_SCOPE_SUB_GROUP | + PI_MEMORY_SCOPE_WORK_GROUP; + return getInfo(param_value_size, param_value, param_value_size_ret, + capabilities); + } case PI_EXT_ONEAPI_DEVICE_INFO_BFLOAT16_MATH_FUNCTIONS: { int major = 0; sycl::detail::pi::assertion(