From 0099ebd52658e6146b54bc0b7865d6bc7a1edc4d Mon Sep 17 00:00:00 2001 From: Stuart Adams Date: Mon, 22 Jun 2020 22:17:08 +0100 Subject: [PATCH] Updated cuda_piDeviceGetInfo to prepare for image implementation Signed-off-by: Stuart Adams --- sycl/plugins/cuda/pi_cuda.cpp | 130 +++++++++++++++++++++++++++++----- 1 file changed, 114 insertions(+), 16 deletions(-) diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index d1e0722e153c7..b8dc8f71a636a 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -904,44 +904,142 @@ pi_result cuda_piDeviceGetInfo(pi_device device, pi_device_info param_name, } case PI_DEVICE_INFO_IMAGE_SUPPORT: { return getInfo(param_value_size, param_value, param_value_size_ret, - PI_FALSE); + PI_TRUE); } case PI_DEVICE_INFO_MAX_READ_IMAGE_ARGS: { - return getInfo(param_value_size, param_value, param_value_size_ret, 0); + // This call doesn't match to CUDA as it doesn't have images, but instead + // surfaces and textures. No clear call in the CUDA API to determine this, + // but some searching found as of SM 2.x 128 are supported. + return getInfo(param_value_size, param_value, param_value_size_ret, 128u); } case PI_DEVICE_INFO_MAX_WRITE_IMAGE_ARGS: { - return getInfo(param_value_size, param_value, param_value_size_ret, 0u); + // This call doesn't match to CUDA as it doesn't have images, but instead + // surfaces and textures. No clear call in the CUDA API to determine this, + // but some searching found as of SM 2.x 128 are supported. + return getInfo(param_value_size, param_value, param_value_size_ret, 128u); } case PI_DEVICE_INFO_IMAGE2D_MAX_HEIGHT: { - return getInfo(param_value_size, param_value, param_value_size_ret, - size_t(0)); + // Take the smaller of maximum surface and maximum texture height. + int tex_height = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&tex_height, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(tex_height >= 0); + int surf_height = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&surf_height, + CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_HEIGHT, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(surf_height >= 0); + + int min = std::min(tex_height, surf_height); + + return getInfo(param_value_size, param_value, param_value_size_ret, min); } case PI_DEVICE_INFO_IMAGE2D_MAX_WIDTH: { - return getInfo(param_value_size, param_value, param_value_size_ret, - size_t(0)); + // Take the smaller of maximum surface and maximum texture width. + int tex_width = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&tex_width, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(tex_width >= 0); + int surf_width = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&surf_width, + CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_WIDTH, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(surf_width >= 0); + + int min = std::min(tex_width, surf_width); + + return getInfo(param_value_size, param_value, param_value_size_ret, min); } case PI_DEVICE_INFO_IMAGE3D_MAX_HEIGHT: { - return getInfo(param_value_size, param_value, param_value_size_ret, - size_t(0)); + // Take the smaller of maximum surface and maximum texture height. + int tex_height = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&tex_height, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(tex_height >= 0); + int surf_height = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&surf_height, + CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_HEIGHT, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(surf_height >= 0); + + int min = std::min(tex_height, surf_height); + + return getInfo(param_value_size, param_value, param_value_size_ret, min); } case PI_DEVICE_INFO_IMAGE3D_MAX_WIDTH: { - return getInfo(param_value_size, param_value, param_value_size_ret, - size_t(0)); + // Take the smaller of maximum surface and maximum texture width. + int tex_width = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&tex_width, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(tex_width >= 0); + int surf_width = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&surf_width, + CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_WIDTH, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(surf_width >= 0); + + int min = std::min(tex_width, surf_width); + + return getInfo(param_value_size, param_value, param_value_size_ret, min); } case PI_DEVICE_INFO_IMAGE3D_MAX_DEPTH: { - return getInfo(param_value_size, param_value, param_value_size_ret, - size_t(0)); + // Take the smaller of maximum surface and maximum texture depth. + int tex_depth = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&tex_depth, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(tex_depth >= 0); + int surf_depth = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&surf_depth, + CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_DEPTH, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(surf_depth >= 0); + + int min = std::min(tex_depth, surf_depth); + + return getInfo(param_value_size, param_value, param_value_size_ret, min); } case PI_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE: { - return getInfo(param_value_size, param_value, param_value_size_ret, - size_t(0)); + // Take the smaller of maximum surface and maximum texture width. + int tex_width = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&tex_width, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(tex_width >= 0); + int surf_width = 0; + cl::sycl::detail::pi::assertion( + cuDeviceGetAttribute(&surf_width, + CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_WIDTH, + device->get()) == CUDA_SUCCESS); + cl::sycl::detail::pi::assertion(surf_width >= 0); + + int min = std::min(tex_width, surf_width); + + return getInfo(param_value_size, param_value, param_value_size_ret, min); } case PI_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE: { return getInfo(param_value_size, param_value, param_value_size_ret, size_t(0)); } case PI_DEVICE_INFO_MAX_SAMPLERS: { - return getInfo(param_value_size, param_value, param_value_size_ret, 0u); + // This call is kind of meaningless for cuda, as samplers don't exist. + // Closest thing is textures, which is 128. + return getInfo(param_value_size, param_value, param_value_size_ret, 128u); } case PI_DEVICE_INFO_MAX_PARAMETER_SIZE: { // https://docs.nvidia.com/cuda/cuda-c-programming-guide/#function-parameters