diff --git a/catch/unit/device/CMakeLists.txt b/catch/unit/device/CMakeLists.txt index 59f64028f6..65d4167262 100644 --- a/catch/unit/device/CMakeLists.txt +++ b/catch/unit/device/CMakeLists.txt @@ -90,3 +90,9 @@ add_dependencies(build_tests getDeviceCount chkUUIDFrmChildProc_Exe chkUUIDInGra if(0) add_dependencies(build_tests passUUIDToGrandChild_Exe ResetUUIDInChild_Exe) endif() + +if(HIP_PLATFORM MATCHES "amd") + add_executable(hipDeviceSetGetScratchExe EXCLUDE_FROM_ALL hipDeviceSetGetScratchExe.cc) + add_dependencies(DeviceTest hipDeviceSetGetScratchExe) + set_property(GLOBAL APPEND PROPERTY G_INSTALL_EXE_TARGETS hipDeviceSetGetScratchExe) +endif() diff --git a/catch/unit/device/hipDeviceSetGetLimit.cc b/catch/unit/device/hipDeviceSetGetLimit.cc index 87b417d942..7071f979ef 100644 --- a/catch/unit/device/hipDeviceSetGetLimit.cc +++ b/catch/unit/device/hipDeviceSetGetLimit.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 ~ 2025 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -21,6 +21,7 @@ THE SOFTWARE. */ #include +#include /** * @addtogroup hipDeviceSetLimit hipDeviceSetLimit @@ -187,4 +188,406 @@ TEST_CASE("Unit_hipDeviceGetLimit_Negative_Parameters") { size_t val; HIP_CHECK_ERROR(hipDeviceGetLimit(&val, static_cast(-1)), hipErrorUnsupportedLimit); } -} \ No newline at end of file +} + +#if HT_AMD +bool isSetScratchLimitSupported() { + #if __linux__ + int deviceId; + HIP_CHECK(hipGetDevice(&deviceId)); + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + std::cout << "Device Id = " << deviceId + << " props.major = " << props.major + << " props.minor = " << props.minor << std::endl; + return ((props.major == 9 && props.minor >= 4) || + (props.major == 12 && props.minor >= 5)) + ? true + : false; + #else + std::cout << "Only Supported for Linux" << std::endl; + return false; + #endif +} + +/** + * Test Description + * ------------------------ + * - This testcase checks following Negative scenarios + * - 1) Call hipDeviceGetLimit, hipDeviceSetLimit with hipLimitRange + * - 2) Call hipDeviceSetLimit with hipExtLimitScratchMin + * - 3) Call hipDeviceSetLimit with hipExtLimitScratchMax + * Test source + * ------------------------ + * - unit/device/hipDeviceSetGetLimit.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDeviceGetSetLimit_Scratch_Negative") { + size_t value = 0; + SECTION("With hipLimitRange") { + HIP_CHECK_ERROR(hipDeviceGetLimit(&value, hipLimitRange), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipDeviceSetLimit(hipLimitRange, value), + hipErrorInvalidValue); + } + + SECTION("Set the minimum value") { + HIP_CHECK_ERROR(hipDeviceSetLimit(hipExtLimitScratchMin, value), + hipErrorUnsupportedLimit); + } + + SECTION("Set the Maximum value") { + HIP_CHECK_ERROR(hipDeviceSetLimit(hipExtLimitScratchMax, value), + hipErrorUnsupportedLimit); + } +} + +/** + * Test Description + * ------------------------ + * - This testcase checks following scenarios + * - 1) Set the Current Scratch limit to Minimum(0) + * - 2) Get the Maximum Scratch limit and set that as Current Scratch limit + * - Set back original Scratch limit value + * Test source + * ------------------------ + * - unit/device/hipDeviceSetGetLimit.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDeviceGetSetLimit_Scratch_SetMinAndMaxAsCurrent") { + if (!isSetScratchLimitSupported()) { + HipTest::HIP_SKIP_TEST("Set Scratch Limit Not Supported on Current Device."\ + " Only Mi300+ and Linux supports"); + return; + } + + size_t scratchLimitCurrent = 0; + HIP_CHECK(hipDeviceGetLimit(&scratchLimitCurrent, hipExtLimitScratchCurrent)); + + SECTION("Set the hipExtLimitScratchCurrent to hipExtLimitScratchMin") { + size_t scratchLimitMin = 0; + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, scratchLimitMin)); + } + + SECTION("Set the hipExtLimitScratchCurrent to hipExtLimitScratchMax") { + size_t scratchLimitMax = 0; + HIP_CHECK(hipDeviceGetLimit(&scratchLimitMax, hipExtLimitScratchMax)); + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, scratchLimitMax)); + } + + // Set back the original value + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, scratchLimitCurrent)); +} + +/** + * Test Description + * ------------------------ + * - This testcase checks following scenarios + * - 1) Set the Current Scratch limit to 1KB less than Current Scratch limit + * - 2) Set the Current Scratch limit to 1KB greater than Current Scratch limit + * - Set back original Scratch limit value + * Test source + * ------------------------ + * - unit/device/hipDeviceSetGetLimit.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDeviceGetSetLimit_Scratch_DecreaseIncrease") { + if (!isSetScratchLimitSupported()) { + HipTest::HIP_SKIP_TEST("Set Scratch Limit Not Supported on Current Device."\ + " Only Mi300+ and Linux supports"); + return; + } + + // Get the current scratch + size_t orgValue = 0; + HIP_CHECK(hipDeviceGetLimit(&orgValue, hipExtLimitScratchCurrent)); + + SECTION("Decrease from the current value") { + size_t setValue = orgValue - 1024; + setValue = (setValue < 0) ? 0 : setValue; + + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, setValue)); + + size_t getValue = 0; + HIP_CHECK(hipDeviceGetLimit(&getValue, hipExtLimitScratchCurrent)); + REQUIRE(getValue == setValue); + } + + SECTION("Increase from the current value") { + size_t setValue = orgValue + 1024; + + size_t max = 0; + HIP_CHECK(hipDeviceGetLimit(&max, hipExtLimitScratchMax)); + setValue = (setValue > max) ? (max*0.9) : setValue; + + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, setValue)); + + size_t getValue = 0; + HIP_CHECK(hipDeviceGetLimit(&getValue, hipExtLimitScratchCurrent)); + REQUIRE(getValue == setValue); + } + + // Set back original value + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, orgValue)); +} + +static constexpr size_t SIZE = 4 * 1024; +static constexpr size_t SIZE_BYTES = SIZE * sizeof(int); +static constexpr int N_BYTES = sizeof(int); + +/* + * Kernel function uses the scratch memory and fill the value + */ +__global__ void addOneKernelUseScratch(int *arr) { + int localArr[SIZE]; + for (int i = 0; i < SIZE; i++) { + localArr[i] = i; + } + + int sum = 0; + for (int i = 0; i < SIZE; i += 1) { + sum += localArr[i]; + } + + *arr = sum; +} + +/** + * Test Description + * ------------------------ + * - This testcase checks following scenario + * - 1) Get the Original Current Scratch limit + * - 2) Set the Current Scratch limit to 16KB + * - 3) Get the Current Scratch limit and validate it + * - 4) Launch the kernel with using scratch + * - 5) Validate the kernel output + * - 6) Set back original Scratch limit value + * Test source + * ------------------------ + * - unit/device/hipDeviceSetGetLimit.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDeviceGetSetLimit_Scratch_SetBeforeKernelLaunch") { + if (!isSetScratchLimitSupported()) { + HipTest::HIP_SKIP_TEST("Set Scratch Limit Not Supported on Current Device."\ + " Only Mi300+ and Linux supports"); + return; + } + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + int *devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, sizeof(int))); + REQUIRE(devMem != nullptr); + + size_t orgValue = 0; + HIP_CHECK(hipDeviceGetLimit(&orgValue, hipExtLimitScratchCurrent)); + + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, SIZE_BYTES)); + + size_t getValue = 0; + HIP_CHECK(hipDeviceGetLimit(&getValue, hipExtLimitScratchCurrent)); + REQUIRE(getValue == SIZE_BYTES); + + addOneKernelUseScratch<<<1, 1, 0, stream>>>(devMem); + HIP_CHECK(hipStreamSynchronize(stream)); + + int hostMem = 0, expectedValue = ((SIZE-1)*(SIZE))/2; + HIP_CHECK(hipMemcpy(&hostMem, devMem, N_BYTES, hipMemcpyDeviceToHost)); + REQUIRE(hostMem == expectedValue); + + // Set back original value + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, orgValue)); + + HIP_CHECK(hipFree(devMem)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +/* + * This helper funtion perform the below operations, + * 1) Get the Minimum Scratch limit and validate + * 2) Get the Maximum Scratch limit and validate + * 3) Get the Current Scratch limit and validate + * 4) Set the Current Scratch limit to 50% of Maximum Scratch limit + * 5) Get the Current Scratch limit and validate it + * 6) Set the Current Scratch limit to original value + */ +void getMinMaxCurrentAndSetCurrent() { + size_t min = 0, max = 0, orgCurrent = 0; + + HIP_CHECK(hipDeviceGetLimit(&min, hipExtLimitScratchMin)); + REQUIRE(min == 0); + + HIP_CHECK(hipDeviceGetLimit(&max, hipExtLimitScratchMax)); + REQUIRE(max > 0); + + HIP_CHECK(hipDeviceGetLimit(&orgCurrent, hipExtLimitScratchCurrent)); + REQUIRE(orgCurrent >= 0); + + size_t setCurrent = 0.5 * max; + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, setCurrent)); + + size_t getCurrent = 0; + HIP_CHECK(hipDeviceGetLimit(&getCurrent, hipExtLimitScratchCurrent)); + REQUIRE(getCurrent == setCurrent); + + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, orgCurrent)); +} + +/** + * Test Description + * ------------------------ + * - This testcase checks the behaviour of + * - hipDeviceGetLimit with hipExtLimitScratchMin, hipExtLimitScratchMax, + * - hipExtLimitScratchCurrent and hipDeviceSetLimit with + * - hipExtLimitScratchCurrent in all the devices. + * Test source + * ------------------------ + * - unit/device/hipDeviceSetGetLimit.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDeviceGetSetLimit_Scratch_MultiDevice") { + int deviceCount = 0; + HIP_CHECK(hipGetDeviceCount(&deviceCount)); + if (deviceCount < 2) { + HipTest::HIP_SKIP_TEST("Skipping because this machine has total GPUs < 2"); + return; + } + + if (!isSetScratchLimitSupported()) { + HipTest::HIP_SKIP_TEST("Set Scratch Limit Not Supported on Current Device."\ + " Only Mi300+ and Linux supports"); + return; + } + + for (int deviceId = 0; deviceId < deviceCount; deviceId++) { + HIP_CHECK(hipSetDevice(deviceId)); + getMinMaxCurrentAndSetCurrent(); + } + HIP_CHECK(hipSetDevice(0)); +} + +/** + * Test Description + * ------------------------ + * - This testcase checks the behaviour of + * - hipDeviceGetLimit with hipExtLimitScratchMin, hipExtLimitScratchMax, + * - hipExtLimitScratchCurrent and hipDeviceSetLimit with + * - hipExtLimitScratchCurrent in the thread. + * Test source + * ------------------------ + * - unit/device/hipDeviceSetGetLimit.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDeviceGetSetLimit_Scratch_InThread") { + if (!isSetScratchLimitSupported()) { + HipTest::HIP_SKIP_TEST("Set Scratch Limit Not Supported on Current Device."\ + " Only Mi300+ and Linux supports"); + return; + } + + std::thread threadObj(getMinMaxCurrentAndSetCurrent); + threadObj.join(); +} + +/** + * Test Description + * ------------------------ + * - This test creates the child process and checks the functionality of + * - hipDeviceGetLimit with hipExtLimitScratchMin, hipExtLimitScratchMax, + * - hipExtLimitScratchCurrent and hipDeviceSetLimit with + * - hipExtLimitScratchCurrent in the child process. + * Test source + * ------------------------ + * - unit/device/hipDeviceSetGetLimit.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDeviceGetSetLimit_Scratch_InChildProcess") { + if (!isSetScratchLimitSupported()) { + HipTest::HIP_SKIP_TEST("Set Scratch Limit Not Supported on Current Device."\ + " Only Mi300+ and Linux supports"); + return; + } + + hip::SpawnProc proc("hipDeviceSetGetScratchExe", true); + REQUIRE(proc.run() == 0); +} + +/* + * Local function to set the given current scratch limit + */ +void setScratchCurrent(size_t setValue) { + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, setValue)); +} + +/* + * Local function to get current scratch limit and validate with + * given reference value + */ +void getScratchCurrent(size_t checkValue) { + size_t getCurrent = 0; + HIP_CHECK(hipDeviceGetLimit(&getCurrent, hipExtLimitScratchCurrent)); + REQUIRE(getCurrent == checkValue); +} + +/** + * Test Description + * ------------------------ + * - This test tests the below scenario + * - 1) Get the Maximum Scratch limit + * - 2) Get the Current Scratch limit + * - 3) Set the Current Scratch limit to 50% of Maximum Scratch limit in thread + * - 4) Get the Current Scratch limit and validate it in thread + * - 5) Set the Current Scratch limit to original value in thread + * - 6) Get the Current Scratch limit and validate it in thread + * Test source + * ------------------------ + * - unit/device/hipDeviceSetGetLimit.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDeviceGetSetLimit_Scratch_SetGetThreads") { + if (!isSetScratchLimitSupported()) { + HipTest::HIP_SKIP_TEST("Set Scratch Limit Not Supported on Current Device."\ + " Only Mi300+ and Linux supports"); + return; + } + + size_t max = 0, orgCurrent = 0; + + HIP_CHECK(hipDeviceGetLimit(&max, hipExtLimitScratchMax)); + REQUIRE(max > 0); + + HIP_CHECK(hipDeviceGetLimit(&orgCurrent, hipExtLimitScratchCurrent)); + REQUIRE(orgCurrent >= 0); + + std::thread setThread1(setScratchCurrent, max * 0.5); + setThread1.join(); + + std::thread getThread1(getScratchCurrent, max * 0.5); + getThread1.join(); + + std::thread setThread2(setScratchCurrent, orgCurrent); + setThread2.join(); + + std::thread getThread2(getScratchCurrent, orgCurrent); + getThread2.join(); +} + +#endif diff --git a/catch/unit/device/hipDeviceSetGetScratchExe.cc b/catch/unit/device/hipDeviceSetGetScratchExe.cc new file mode 100644 index 0000000000..ad447a62bf --- /dev/null +++ b/catch/unit/device/hipDeviceSetGetScratchExe.cc @@ -0,0 +1,74 @@ +/* +Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip/hip_runtime.h" + +#define HIP_CHECK(error) \ + { \ + hipError_t localError = error; \ + if ((localError != hipSuccess) && \ + (localError != hipErrorPeerAccessAlreadyEnabled)) { \ + printf("error: '%s'(%d) from %s at %s:%d\n", \ + hipGetErrorString(localError), localError, #error, __FUNCTION__, \ + __LINE__); \ + return -1; \ + } \ + } + +#define REQUIRE(res) \ + { \ + if (res) { \ + return 0; \ + } else { \ + return -1; \ + } \ + } + +/* + * This funtion perform the below operations, + * 1) Get the Minimum Scratch limit and validate + * 2) Get the Maximum Scratch limit and validate + * 3) Get the Current Scratch limit and validate + * 4) Set the Current Scratch limit to 50% of Maximum Scratch limit + * 5) Get the Current Scratch limit and validate it + * 6) Set the Current Scratch limit to original value + */ +int main() { + size_t min = 0, max = 0, orgCurrent = 0; + + HIP_CHECK(hipDeviceGetLimit(&min, hipExtLimitScratchMin)); + REQUIRE(min == 0); + + HIP_CHECK(hipDeviceGetLimit(&max, hipExtLimitScratchMax)); + REQUIRE(max > 0); + + HIP_CHECK(hipDeviceGetLimit(&orgCurrent, hipExtLimitScratchCurrent)); + REQUIRE(orgCurrent >= 0); + + size_t setCurrent = 0.5 * max; + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, setCurrent)); + + size_t getCurrent = 0; + HIP_CHECK(hipDeviceGetLimit(&getCurrent, hipExtLimitScratchCurrent)); + REQUIRE(getCurrent == setCurrent); + + HIP_CHECK(hipDeviceSetLimit(hipExtLimitScratchCurrent, orgCurrent)); + + return 0; +}