SWDEV-493276 - [catch2][dtest] Additional tests for Scratch limit change (#46)
* SWDEV-493276 - [catch2][dtest] Additional tests for Scratch limit change * Removed the Kernel wait related scenarios * Modified kernel operation * Fixed PSDB NVIDIA build issue * Restricted for Windows * Used isSetScratchLimitSupported for all tests and modified isSetScratchLimitSupported
Šī revīzija ir iekļauta:
Swargam, Rambabu
revīziju iesūtīja
GitHub
GitHub
vecāks
b23ebe9a0e
revīzija
f7712bd5a8
@@ -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()
|
||||
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_process.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup hipDeviceSetLimit hipDeviceSetLimit
|
||||
@@ -187,4 +188,406 @@ TEST_CASE("Unit_hipDeviceGetLimit_Negative_Parameters") {
|
||||
size_t val;
|
||||
HIP_CHECK_ERROR(hipDeviceGetLimit(&val, static_cast<hipLimit_t>(-1)), hipErrorUnsupportedLimit);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#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
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
Atsaukties uz šo jaunā problēmā
Block a user