Skip to content

Commit 9c0dc7f

Browse files
nyallocnyalloc
authored
[SYCL][CUDA] PI API CUDA piDeviceGetInfo with Image / Sampler support (#1951)
Signed-off-by: Stuart Adams <stuart.adams@codeplay.com>
1 parent 03ef819 commit 9c0dc7f

File tree

1 file changed

+114
-16
lines changed

1 file changed

+114
-16
lines changed

‎sycl/plugins/cuda/pi_cuda.cpp‎

Lines changed: 114 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -931,44 +931,142 @@ pi_result cuda_piDeviceGetInfo(pi_device device, pi_device_info param_name,
931931
}
932932
case PI_DEVICE_INFO_IMAGE_SUPPORT: {
933933
return getInfo(param_value_size, param_value, param_value_size_ret,
934-
PI_FALSE);
934+
PI_TRUE);
935935
}
936936
case PI_DEVICE_INFO_MAX_READ_IMAGE_ARGS: {
937-
return getInfo(param_value_size, param_value, param_value_size_ret, 0);
937+
// This call doesn't match to CUDA as it doesn't have images, but instead
938+
// surfaces and textures. No clear call in the CUDA API to determine this,
939+
// but some searching found as of SM 2.x 128 are supported.
940+
return getInfo(param_value_size, param_value, param_value_size_ret, 128u);
938941
}
939942
case PI_DEVICE_INFO_MAX_WRITE_IMAGE_ARGS: {
940-
return getInfo(param_value_size, param_value, param_value_size_ret, 0u);
943+
// This call doesn't match to CUDA as it doesn't have images, but instead
944+
// surfaces and textures. No clear call in the CUDA API to determine this,
945+
// but some searching found as of SM 2.x 128 are supported.
946+
return getInfo(param_value_size, param_value, param_value_size_ret, 128u);
941947
}
942948
case PI_DEVICE_INFO_IMAGE2D_MAX_HEIGHT: {
943-
return getInfo(param_value_size, param_value, param_value_size_ret,
944-
size_t(0));
949+
// Take the smaller of maximum surface and maximum texture height.
950+
int tex_height = 0;
951+
cl::sycl::detail::pi::assertion(
952+
cuDeviceGetAttribute(&tex_height,
953+
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT,
954+
device->get()) == CUDA_SUCCESS);
955+
cl::sycl::detail::pi::assertion(tex_height >= 0);
956+
int surf_height = 0;
957+
cl::sycl::detail::pi::assertion(
958+
cuDeviceGetAttribute(&surf_height,
959+
CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_HEIGHT,
960+
device->get()) == CUDA_SUCCESS);
961+
cl::sycl::detail::pi::assertion(surf_height >= 0);
962+
963+
int min = std::min(tex_height, surf_height);
964+
965+
return getInfo(param_value_size, param_value, param_value_size_ret, min);
945966
}
946967
case PI_DEVICE_INFO_IMAGE2D_MAX_WIDTH: {
947-
return getInfo(param_value_size, param_value, param_value_size_ret,
948-
size_t(0));
968+
// Take the smaller of maximum surface and maximum texture width.
969+
int tex_width = 0;
970+
cl::sycl::detail::pi::assertion(
971+
cuDeviceGetAttribute(&tex_width,
972+
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH,
973+
device->get()) == CUDA_SUCCESS);
974+
cl::sycl::detail::pi::assertion(tex_width >= 0);
975+
int surf_width = 0;
976+
cl::sycl::detail::pi::assertion(
977+
cuDeviceGetAttribute(&surf_width,
978+
CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_WIDTH,
979+
device->get()) == CUDA_SUCCESS);
980+
cl::sycl::detail::pi::assertion(surf_width >= 0);
981+
982+
int min = std::min(tex_width, surf_width);
983+
984+
return getInfo(param_value_size, param_value, param_value_size_ret, min);
949985
}
950986
case PI_DEVICE_INFO_IMAGE3D_MAX_HEIGHT: {
951-
return getInfo(param_value_size, param_value, param_value_size_ret,
952-
size_t(0));
987+
// Take the smaller of maximum surface and maximum texture height.
988+
int tex_height = 0;
989+
cl::sycl::detail::pi::assertion(
990+
cuDeviceGetAttribute(&tex_height,
991+
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT,
992+
device->get()) == CUDA_SUCCESS);
993+
cl::sycl::detail::pi::assertion(tex_height >= 0);
994+
int surf_height = 0;
995+
cl::sycl::detail::pi::assertion(
996+
cuDeviceGetAttribute(&surf_height,
997+
CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_HEIGHT,
998+
device->get()) == CUDA_SUCCESS);
999+
cl::sycl::detail::pi::assertion(surf_height >= 0);
1000+
1001+
int min = std::min(tex_height, surf_height);
1002+
1003+
return getInfo(param_value_size, param_value, param_value_size_ret, min);
9531004
}
9541005
case PI_DEVICE_INFO_IMAGE3D_MAX_WIDTH: {
955-
return getInfo(param_value_size, param_value, param_value_size_ret,
956-
size_t(0));
1006+
// Take the smaller of maximum surface and maximum texture width.
1007+
int tex_width = 0;
1008+
cl::sycl::detail::pi::assertion(
1009+
cuDeviceGetAttribute(&tex_width,
1010+
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH,
1011+
device->get()) == CUDA_SUCCESS);
1012+
cl::sycl::detail::pi::assertion(tex_width >= 0);
1013+
int surf_width = 0;
1014+
cl::sycl::detail::pi::assertion(
1015+
cuDeviceGetAttribute(&surf_width,
1016+
CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_WIDTH,
1017+
device->get()) == CUDA_SUCCESS);
1018+
cl::sycl::detail::pi::assertion(surf_width >= 0);
1019+
1020+
int min = std::min(tex_width, surf_width);
1021+
1022+
return getInfo(param_value_size, param_value, param_value_size_ret, min);
9571023
}
9581024
case PI_DEVICE_INFO_IMAGE3D_MAX_DEPTH: {
959-
return getInfo(param_value_size, param_value, param_value_size_ret,
960-
size_t(0));
1025+
// Take the smaller of maximum surface and maximum texture depth.
1026+
int tex_depth = 0;
1027+
cl::sycl::detail::pi::assertion(
1028+
cuDeviceGetAttribute(&tex_depth,
1029+
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH,
1030+
device->get()) == CUDA_SUCCESS);
1031+
cl::sycl::detail::pi::assertion(tex_depth >= 0);
1032+
int surf_depth = 0;
1033+
cl::sycl::detail::pi::assertion(
1034+
cuDeviceGetAttribute(&surf_depth,
1035+
CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_DEPTH,
1036+
device->get()) == CUDA_SUCCESS);
1037+
cl::sycl::detail::pi::assertion(surf_depth >= 0);
1038+
1039+
int min = std::min(tex_depth, surf_depth);
1040+
1041+
return getInfo(param_value_size, param_value, param_value_size_ret, min);
9611042
}
9621043
case PI_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE: {
963-
return getInfo(param_value_size, param_value, param_value_size_ret,
964-
size_t(0));
1044+
// Take the smaller of maximum surface and maximum texture width.
1045+
int tex_width = 0;
1046+
cl::sycl::detail::pi::assertion(
1047+
cuDeviceGetAttribute(&tex_width,
1048+
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH,
1049+
device->get()) == CUDA_SUCCESS);
1050+
cl::sycl::detail::pi::assertion(tex_width >= 0);
1051+
int surf_width = 0;
1052+
cl::sycl::detail::pi::assertion(
1053+
cuDeviceGetAttribute(&surf_width,
1054+
CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_WIDTH,
1055+
device->get()) == CUDA_SUCCESS);
1056+
cl::sycl::detail::pi::assertion(surf_width >= 0);
1057+
1058+
int min = std::min(tex_width, surf_width);
1059+
1060+
return getInfo(param_value_size, param_value, param_value_size_ret, min);
9651061
}
9661062
case PI_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE: {
9671063
return getInfo(param_value_size, param_value, param_value_size_ret,
9681064
size_t(0));
9691065
}
9701066
case PI_DEVICE_INFO_MAX_SAMPLERS: {
971-
return getInfo(param_value_size, param_value, param_value_size_ret, 0u);
1067+
// This call is kind of meaningless for cuda, as samplers don't exist.
1068+
// Closest thing is textures, which is 128.
1069+
return getInfo(param_value_size, param_value, param_value_size_ret, 128u);
9721070
}
9731071
case PI_DEVICE_INFO_MAX_PARAMETER_SIZE: {
9741072
// https://docs.nvidia.com/cuda/cuda-c-programming-guide/#function-parameters

0 commit comments

Comments
 (0)