From d394977800285af11615b9a2653c36d572627ac7 Mon Sep 17 00:00:00 2001 From: Nicolas Miller Date: Fri, 17 Dec 2021 11:15:33 +0000 Subject: [PATCH 1/2] [SYCL][HIP] Fix max constant memory device query This fixes the `Basic/info.cpp` test for HIP AMD. The issue here is that AMD GPU report a very large constant memory (confirmed using `clinfo` as well). But the HIP entry point to query that information takes a `int*`. So the value will show up as negative, even though it's the correct value. So remove the assertion on if the value is positive and cast it back to `unsigned` before returning it. --- sycl/plugins/hip/pi_hip.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index dce26bc020715..47a0b30de27e5 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -1320,10 +1320,13 @@ pi_result hip_piDeviceGetInfo(pi_device device, pi_device_info param_name, hipDeviceGetAttribute(&constant_memory, hipDeviceAttributeTotalConstantMemory, device->get()) == hipSuccess); - cl::sycl::detail::pi::assertion(constant_memory >= 0); + // hipDeviceGetAttribute takes a int*, however the size of the constant + // memory on AMD GPU may be larger than what can fit in the positive part + // of a signed integer, so we need to cast it to unsigned before returning + // the value. return getInfo(param_value_size, param_value, param_value_size_ret, - pi_uint64(constant_memory)); + pi_uint64(static_cast(constant_memory))); } case PI_DEVICE_INFO_MAX_CONSTANT_ARGS: { // TODO: is there a way to retrieve this from HIP driver API? From b283623a780d551580c70c1e34a4475d5edd13f0 Mon Sep 17 00:00:00 2001 From: Nicolas Miller Date: Tue, 4 Jan 2022 10:40:13 +0000 Subject: [PATCH 2/2] [SYCL][HIP] Move cast around --- sycl/plugins/hip/pi_hip.cpp | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index 47a0b30de27e5..b9c6f65968303 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -1315,18 +1315,19 @@ pi_result hip_piDeviceGetInfo(pi_device device, pi_device_info param_name, pi_uint64{bytes}); } case PI_DEVICE_INFO_MAX_CONSTANT_BUFFER_SIZE: { - int constant_memory = 0; + unsigned int constant_memory = 0; + + // hipDeviceGetAttribute takes a int*, however the size of the constant + // memory on AMD GPU may be larger than what can fit in the positive part + // of a signed integer, so use an unsigned integer and cast the pointer to + // int*. cl::sycl::detail::pi::assertion( - hipDeviceGetAttribute(&constant_memory, + hipDeviceGetAttribute(reinterpret_cast(&constant_memory), hipDeviceAttributeTotalConstantMemory, device->get()) == hipSuccess); - // hipDeviceGetAttribute takes a int*, however the size of the constant - // memory on AMD GPU may be larger than what can fit in the positive part - // of a signed integer, so we need to cast it to unsigned before returning - // the value. return getInfo(param_value_size, param_value, param_value_size_ret, - pi_uint64(static_cast(constant_memory))); + pi_uint64(constant_memory)); } case PI_DEVICE_INFO_MAX_CONSTANT_ARGS: { // TODO: is there a way to retrieve this from HIP driver API?