SWDEV-292637 - [dtest] Catch2 unit and multiprocess tests for Memset3d,HostMalloc and MallocConcurrency tests (#2348)
Change-Id: I9025bc13735c1d9fb0f0811a9c9d6ad304adc134
[ROCm/hip commit: 8c0558c448]
This commit is contained in:
sumanthtg
GitHub
@@ -69,6 +69,7 @@ add_subdirectory(unit)
|
||||
add_subdirectory(ABM)
|
||||
add_subdirectory(hipTestMain)
|
||||
add_subdirectory(stress)
|
||||
add_subdirectory(TypeQualifiers)
|
||||
|
||||
if(UNIX)
|
||||
add_subdirectory(multiproc)
|
||||
|
||||
@@ -0,0 +1,12 @@
|
||||
# Common Tests
|
||||
set(TEST_SRC
|
||||
hipManagedKeyword.cc
|
||||
)
|
||||
|
||||
|
||||
# Create shared lib of all tests
|
||||
add_library(TypeQualifiers SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests TypeQualifiers)
|
||||
|
||||
@@ -71,3 +71,14 @@ target_link_libraries(StressTest PRIVATE memory stdc++fs)
|
||||
add_dependencies(build_stress_test StressTest)
|
||||
add_custom_target(stress_test COMMAND StressTest)
|
||||
|
||||
# Space Specifiers/Qualifiers exe
|
||||
add_executable(TypeQualifierTests EXCLUDE_FROM_ALL main.cc hip_test_context.cc)
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set_property(TARGET TypeQualifierTests PROPERTY CXX_STANDARD 17)
|
||||
else()
|
||||
target_compile_options(TypeQualifierTests PUBLIC -std=c++17)
|
||||
endif()
|
||||
target_link_libraries(TypeQualifierTests PRIVATE TypeQualifiers stdc++fs)
|
||||
|
||||
catch_discover_tests(TypeQualifierTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
|
||||
add_dependencies(build_tests TypeQualifierTests)
|
||||
|
||||
@@ -0,0 +1,88 @@
|
||||
/*
|
||||
Copyright (c) 2021 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 WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
/**
|
||||
* @brief Error codes retured by rocm_smi_lib functions
|
||||
*/
|
||||
typedef enum {
|
||||
RSMI_STATUS_SUCCESS = 0x0, //!< Operation was successful
|
||||
RSMI_STATUS_INVALID_ARGS, //!< Passed in arguments are not valid
|
||||
RSMI_STATUS_NOT_SUPPORTED, //!< The requested information or
|
||||
//!< action is not available for the
|
||||
//!< given input, on the given system
|
||||
RSMI_STATUS_FILE_ERROR, //!< Problem accessing a file. This
|
||||
//!< may because the operation is not
|
||||
//!< supported by the Linux kernel
|
||||
//!< version running on the executing
|
||||
//!< machine
|
||||
RSMI_STATUS_PERMISSION, //!< Permission denied/EACCESS file
|
||||
//!< error. Many functions require
|
||||
//!< root access to run.
|
||||
RSMI_STATUS_OUT_OF_RESOURCES, //!< Unable to acquire memory or other
|
||||
//!< resource
|
||||
RSMI_STATUS_INTERNAL_EXCEPTION, //!< An internal exception was caught
|
||||
RSMI_STATUS_INPUT_OUT_OF_BOUNDS, //!< The provided input is out of
|
||||
//!< allowable or safe range
|
||||
RSMI_STATUS_INIT_ERROR, //!< An error occurred when rsmi
|
||||
//!< initializing internal data
|
||||
//!< structures
|
||||
RSMI_INITIALIZATION_ERROR = RSMI_STATUS_INIT_ERROR,
|
||||
RSMI_STATUS_NOT_YET_IMPLEMENTED, //!< The requested function has not
|
||||
//!< yet been implemented in the
|
||||
//!< current system for the current
|
||||
//!< devices
|
||||
RSMI_STATUS_NOT_FOUND, //!< An item was searched for but not
|
||||
//!< found
|
||||
RSMI_STATUS_INSUFFICIENT_SIZE, //!< Not enough resources were
|
||||
//!< available for the operation
|
||||
RSMI_STATUS_INTERRUPT, //!< An interrupt occurred during
|
||||
//!< execution of function
|
||||
RSMI_STATUS_UNEXPECTED_SIZE, //!< An unexpected amount of data
|
||||
//!< was read
|
||||
RSMI_STATUS_NO_DATA, //!< No data was found for a given
|
||||
//!< input
|
||||
RSMI_STATUS_UNEXPECTED_DATA, //!< The data read or provided to
|
||||
//!< function is not what was expected
|
||||
RSMI_STATUS_BUSY, //!< A resource or mutex could not be
|
||||
//!< acquired because it is already
|
||||
//!< being used
|
||||
RSMI_STATUS_REFCOUNT_OVERFLOW, //!< An internal reference counter
|
||||
//!< exceeded INT32_MAX
|
||||
|
||||
RSMI_STATUS_UNKNOWN_ERROR = 0xFFFFFFFF, //!< An unknown error occurred
|
||||
} rsmi_status_t;
|
||||
|
||||
|
||||
/**
|
||||
* @brief Types of memory
|
||||
*/
|
||||
typedef enum {
|
||||
RSMI_MEM_TYPE_FIRST = 0,
|
||||
|
||||
RSMI_MEM_TYPE_VRAM = RSMI_MEM_TYPE_FIRST, //!< VRAM memory
|
||||
RSMI_MEM_TYPE_VIS_VRAM, //!< VRAM memory that is visible
|
||||
RSMI_MEM_TYPE_GTT, //!< GTT memory
|
||||
|
||||
RSMI_MEM_TYPE_LAST = RSMI_MEM_TYPE_GTT
|
||||
} rsmi_memory_type_t;
|
||||
@@ -1,6 +1,5 @@
|
||||
# AMD Tests
|
||||
# Common Tests
|
||||
set(LINUX_TEST_SRC
|
||||
hipMallocConcurrency.cc
|
||||
childMalloc.cc
|
||||
hipDeviceComputeCapabilityMproc.cc
|
||||
hipDeviceGetPCIBusIdMproc.cc
|
||||
@@ -10,12 +9,16 @@ set(LINUX_TEST_SRC
|
||||
hipGetDevicePropertiesMproc.cc
|
||||
hipSetGetDeviceMproc.cc
|
||||
hipIpcMemAccessTest.cc
|
||||
hipHostMallocTestsMproc.cc
|
||||
hipMallocConcurrencyMproc.cc
|
||||
)
|
||||
|
||||
if(UNIX)
|
||||
# Create shared lib of all tests
|
||||
add_library(MultiProc SHARED EXCLUDE_FROM_ALL ${LINUX_TEST_SRC})
|
||||
|
||||
target_link_libraries(MultiProc ${CMAKE_DL_LIBS})
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests MultiProc)
|
||||
endif()
|
||||
|
||||
@@ -0,0 +1,314 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Validate page table allocation of hipHostMalloc() api when
|
||||
HIP_VISIBLE_DEVICES set to single device.
|
||||
2) Validate page table allocation of hipHostMalloc() api when
|
||||
HIP_VISIBLE_DEVICES set to list of multiple devices.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_smi.hh>
|
||||
|
||||
#if defined(__linux__)
|
||||
#include <unistd.h>
|
||||
#include <stdlib.h>
|
||||
#include <sys/wait.h>
|
||||
#include <dlfcn.h>
|
||||
|
||||
#if HT_AMD
|
||||
/* AMD specific as below tests use rocm-smi */
|
||||
|
||||
/**
|
||||
* Defines
|
||||
*/
|
||||
#define LIB_ROCMSMI "librocm_smi64.so"
|
||||
#define ALLOC_SIZE (1*1024*1024)
|
||||
|
||||
/**
|
||||
* Global variables
|
||||
*/
|
||||
static rsmi_status_t (*rsmi_dev_memory_usage_get_fp)(uint32_t,
|
||||
rsmi_memory_type_t,
|
||||
uint64_t *);
|
||||
static rsmi_status_t (*rsmi_init_fp)(uint64_t);
|
||||
static rsmi_status_t (*rsmi_shut_down_fp)();
|
||||
static void *rocm_smi_h;
|
||||
|
||||
|
||||
/**
|
||||
* Fetches Gpu device count
|
||||
*/
|
||||
static void getDeviceCount(int *pdevCnt) {
|
||||
int fd[2], val = 0;
|
||||
pid_t childpid;
|
||||
|
||||
// create pipe descriptors
|
||||
pipe(fd);
|
||||
|
||||
// disable visible_devices env from shell
|
||||
unsetenv("ROCR_VISIBLE_DEVICES");
|
||||
unsetenv("HIP_VISIBLE_DEVICES");
|
||||
|
||||
childpid = fork();
|
||||
|
||||
if (childpid > 0) { // Parent
|
||||
close(fd[1]);
|
||||
// parent will wait to read the device cnt
|
||||
read(fd[0], &val, sizeof(val));
|
||||
|
||||
// close the read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
// wait for child exit
|
||||
wait(nullptr);
|
||||
|
||||
*pdevCnt = val;
|
||||
} else if (!childpid) { // Child
|
||||
int devCnt = 1;
|
||||
// writing only, no need for read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
// send the value on the write-descriptor:
|
||||
write(fd[1], &devCnt, sizeof(devCnt));
|
||||
|
||||
// close the write descriptor:
|
||||
close(fd[1]);
|
||||
exit(0);
|
||||
} else { // failure
|
||||
*pdevCnt = 0;
|
||||
return;
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Initializes rocm smi library handles
|
||||
*/
|
||||
static bool rocm_smi_init() {
|
||||
// Open ROCm SMI Library
|
||||
if (!(rocm_smi_h = dlopen(LIB_ROCMSMI, RTLD_LAZY))) {
|
||||
printf("Error opening rocm smi library!\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
void* fnsym = dlsym(rocm_smi_h, "rsmi_dev_memory_usage_get");
|
||||
if (!fnsym) {
|
||||
printf("Error getting rsmi_dev_memory_usage_get() function\n");
|
||||
dlclose(rocm_smi_h);
|
||||
return false;
|
||||
}
|
||||
rsmi_dev_memory_usage_get_fp = reinterpret_cast<rsmi_status_t (*)(uint32_t,
|
||||
rsmi_memory_type_t, uint64_t *)>(fnsym);
|
||||
|
||||
fnsym = dlsym(rocm_smi_h, "rsmi_init");
|
||||
if (!fnsym) {
|
||||
printf("Error getting rsmi_init() function\n");
|
||||
dlclose(rocm_smi_h);
|
||||
return false;
|
||||
}
|
||||
rsmi_init_fp = reinterpret_cast<rsmi_status_t (*)(uint64_t)>(fnsym);
|
||||
|
||||
fnsym = dlsym(rocm_smi_h, "rsmi_shut_down");
|
||||
if (!fnsym) {
|
||||
printf("Error getting rsmi_shut_down() function\n");
|
||||
dlclose(rocm_smi_h);
|
||||
return false;
|
||||
}
|
||||
rsmi_shut_down_fp = reinterpret_cast<rsmi_status_t (*)()>(fnsym);
|
||||
|
||||
uint64_t init_flags = 0;
|
||||
rsmi_status_t retsmi_init;
|
||||
retsmi_init = rsmi_init_fp(init_flags);
|
||||
if (RSMI_STATUS_SUCCESS != retsmi_init) {
|
||||
printf("Error when initializing rocm_smi\n");
|
||||
dlclose(rocm_smi_h);
|
||||
return false;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* Exits rocm smi library
|
||||
*/
|
||||
static void rocm_smi_exit() {
|
||||
rsmi_shut_down_fp();
|
||||
dlclose(rocm_smi_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates page table memory allocations
|
||||
* by setting visible devices selected.
|
||||
*/
|
||||
static bool validatePageTableAllocations(const char *devList, int visDevCnt) {
|
||||
int fd[2];
|
||||
bool testResult = false;
|
||||
pid_t pid;
|
||||
int numdev = 0;
|
||||
|
||||
getDeviceCount(&numdev);
|
||||
REQUIRE(numdev > 0);
|
||||
|
||||
if (pipe(fd) < 0) {
|
||||
printf("Pipe system call failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
pid = fork();
|
||||
|
||||
if (!pid) { // Child process
|
||||
rsmi_status_t ret;
|
||||
std::vector<int> prev, current;
|
||||
uint64_t used = 0;
|
||||
int tmpdev = 0, changeCnt = 0, indx = 0;
|
||||
char *ptr = nullptr;
|
||||
bool testPassed = true;
|
||||
|
||||
// Disable visible_devices env from shell
|
||||
unsetenv("ROCR_VISIBLE_DEVICES");
|
||||
unsetenv("HIP_VISIBLE_DEVICES");
|
||||
|
||||
setenv("HIP_VISIBLE_DEVICES", devList, 1);
|
||||
|
||||
// First Call to initialize hip api
|
||||
hipGetDeviceCount(&tmpdev);
|
||||
|
||||
|
||||
// Get memory snapshot before hostmalloc
|
||||
for (indx = 0; indx < numdev; indx++) {
|
||||
ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used);
|
||||
if (RSMI_STATUS_SUCCESS != ret) {
|
||||
printf("Error while running rsmi_dev_memory_usage_get func\n");
|
||||
dlclose(rocm_smi_h);
|
||||
rsmi_shut_down_fp();
|
||||
return false;
|
||||
}
|
||||
prev.push_back(used);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipHostMalloc(&ptr, ALLOC_SIZE));
|
||||
|
||||
// Get memory snapshot after hostmalloc
|
||||
for (indx = 0; indx < numdev; indx++) {
|
||||
ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used);
|
||||
if (RSMI_STATUS_SUCCESS != ret) {
|
||||
printf("Error while running rsmi_dev_memory_usage_get func\n");
|
||||
dlclose(rocm_smi_h);
|
||||
rsmi_shut_down_fp();
|
||||
hipHostFree(ptr);
|
||||
return false;
|
||||
}
|
||||
current.push_back(used);
|
||||
}
|
||||
|
||||
for (indx = 0; indx < numdev; indx++) {
|
||||
if (current[indx] - prev[indx])
|
||||
changeCnt++;
|
||||
}
|
||||
|
||||
// Check if memory allocation happened only for visible devices
|
||||
if (changeCnt == visDevCnt) {
|
||||
testPassed = true;
|
||||
} else {
|
||||
testPassed = false;
|
||||
}
|
||||
|
||||
hipHostFree(ptr);
|
||||
|
||||
// writing only, no need for read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
// send the value on the write-descriptor:
|
||||
write(fd[1], &testPassed, sizeof(testPassed));
|
||||
|
||||
// close the write descriptor:
|
||||
close(fd[1]);
|
||||
exit(0);
|
||||
} else if (pid > 0) { // parent
|
||||
close(fd[1]);
|
||||
read(fd[0], &testResult, sizeof(testResult));
|
||||
close(fd[0]);
|
||||
wait(nullptr);
|
||||
} else {
|
||||
printf("fork() failed\n");
|
||||
HIP_ASSERT(false);
|
||||
}
|
||||
|
||||
return testResult;
|
||||
}
|
||||
|
||||
/**
|
||||
* Test page table allocation when HIP_VISIBLE_DEVICES set to
|
||||
* single device
|
||||
*/
|
||||
TEST_CASE("Unit_hipHostMalloc_SingleVisibleDevicePageAlloc") {
|
||||
int devCnt;
|
||||
std::string str;
|
||||
|
||||
if (!rocm_smi_init()) {
|
||||
WARN("Testcase skipped as rocm smi not initialized/present");
|
||||
return;
|
||||
}
|
||||
|
||||
getDeviceCount(&devCnt);
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
// Select single visible device and validate memory usage
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
str = std::to_string(i);
|
||||
REQUIRE(validatePageTableAllocations(str.c_str(), 1) == true);
|
||||
}
|
||||
|
||||
rocm_smi_exit();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test page table allocation when HIP_VISIBLE_DEVICES set to
|
||||
* multiple devices
|
||||
*/
|
||||
TEST_CASE("Unit_hipHostMalloc_MultipleVisibleDevicesPageAlloc") {
|
||||
int devCnt = 0, vdCnt = 0;
|
||||
std::string str;
|
||||
|
||||
if (!rocm_smi_init()) {
|
||||
WARN("Testcase skipped as rocm smi not initialized/present");
|
||||
return;
|
||||
}
|
||||
|
||||
getDeviceCount(&devCnt);
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
// Select multiple visible devices and validate memory usage
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
if (i == 0)
|
||||
str += std::to_string(i);
|
||||
else
|
||||
str += "," + std::to_string(i);
|
||||
|
||||
vdCnt++;
|
||||
REQUIRE(validatePageTableAllocations(str.c_str(), vdCnt) == true);
|
||||
}
|
||||
|
||||
rocm_smi_exit();
|
||||
}
|
||||
#endif // HT_AMD
|
||||
#endif // __linux__
|
||||
@@ -1,171 +0,0 @@
|
||||
#include <sys/types.h>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
#ifdef __linux__
|
||||
#include <sys/wait.h>
|
||||
#include <unistd.h>
|
||||
#endif
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <limits>
|
||||
#include <atomic>
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
|
||||
static constexpr size_t N = 4 * 1024 * 1024;
|
||||
static constexpr unsigned blocksPerCU = 6; // to hide latency
|
||||
static constexpr unsigned threadsPerBlock = 256;
|
||||
/**
|
||||
* Validates data consitency on supplied gpu
|
||||
*/
|
||||
bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
size_t prevAvl, prevTot, curAvl, curTot;
|
||||
bool TestPassed = true;
|
||||
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
|
||||
printf("tgs allocating..\n");
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
|
||||
printf("Validation PASSED for gpu %d from pid %d\n", gpu, getpid());
|
||||
} else {
|
||||
printf("%s : Validation FAILED for gpu %d from pid %d\n", __func__, gpu, getpid());
|
||||
TestPassed &= false;
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
|
||||
|
||||
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
|
||||
// In concurrent calls on one GPU, we cannot verify leaking in this way
|
||||
printf(
|
||||
"%s : Memory allocation mismatch observed."
|
||||
"Possible memory leak.\n",
|
||||
__func__);
|
||||
TestPassed &= false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
|
||||
#if 1
|
||||
/**
|
||||
* Fetches Gpu device count
|
||||
*/
|
||||
void getDeviceCount1(int* pdevCnt) {
|
||||
#ifdef __linux__
|
||||
int fd[2], val = 0;
|
||||
pid_t childpid;
|
||||
|
||||
// create pipe descriptors
|
||||
pipe(fd);
|
||||
|
||||
// disable visible_devices env from shell
|
||||
unsetenv("ROCR_VISIBLE_DEVICES");
|
||||
unsetenv("HIP_VISIBLE_DEVICES");
|
||||
|
||||
childpid = fork();
|
||||
|
||||
if (childpid > 0) { // Parent
|
||||
close(fd[1]);
|
||||
// parent will wait to read the device cnt
|
||||
read(fd[0], &val, sizeof(val));
|
||||
|
||||
// close the read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
// wait for child exit
|
||||
wait(NULL);
|
||||
|
||||
*pdevCnt = val;
|
||||
} else if (!childpid) { // Child
|
||||
int devCnt = 1;
|
||||
// writing only, no need for read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
// send the value on the write-descriptor:
|
||||
write(fd[1], &devCnt, sizeof(devCnt));
|
||||
|
||||
// close the write descriptor:
|
||||
close(fd[1]);
|
||||
exit(0);
|
||||
} else { // failure
|
||||
*pdevCnt = 1;
|
||||
return;
|
||||
}
|
||||
|
||||
#else
|
||||
HIP_CHECK(hipGetDeviceCount(pdevCnt));
|
||||
#endif
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
TEST_CASE("hipMallocChild_Concurrency_MultiGpu") {
|
||||
bool TestPassed = false;
|
||||
#ifdef __linux__
|
||||
// Parallel execution on multiple gpus from different child processes
|
||||
int devCnt = 1, pid = 0;
|
||||
|
||||
// Get GPU count
|
||||
getDeviceCount1(&devCnt);
|
||||
|
||||
// Spawn child for each GPU
|
||||
for (int gpu = 0; gpu < devCnt; gpu++) {
|
||||
if ((pid = fork()) < 0) {
|
||||
INFO("Child_Concurrency_MultiGpu : fork() returned error" << pid);
|
||||
REQUIRE(false);
|
||||
|
||||
} else if (!pid) { // Child process
|
||||
bool TestPassedChild = false;
|
||||
TestPassedChild = validateMemoryOnGPU(gpu);
|
||||
|
||||
if (TestPassedChild) {
|
||||
printf("returning exit(1) for success\n");
|
||||
exit(1); // child exit with success status
|
||||
} else {
|
||||
printf("Child_Concurrency_MultiGpu : childpid %d failed\n", getpid());
|
||||
exit(2); // child exit with failure status
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Parent shall wait for child to complete
|
||||
int cnt = 0;
|
||||
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
int pidwait = 0, exitStatus;
|
||||
pidwait = wait(&exitStatus);
|
||||
|
||||
printf("exitStatus for iter:%d is %d\n", i, exitStatus);
|
||||
if (pidwait < 0) {
|
||||
break;
|
||||
}
|
||||
|
||||
if (WEXITSTATUS(exitStatus) == 1) cnt++;
|
||||
}
|
||||
|
||||
if (cnt && (cnt == devCnt)) TestPassed = true;
|
||||
|
||||
#else
|
||||
INFO("Test hipMallocChild_Concurrency_MultiGpu skipped on non-linux");
|
||||
#endif
|
||||
REQUIRE(TestPassed == true);
|
||||
}
|
||||
@@ -0,0 +1,228 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
|
||||
1) Run hipMalloc() api/kernel code on same gpu parallely from parent and child
|
||||
processes, validate the results.
|
||||
|
||||
2) Execute hipMalloc() api simultaneously on all the gpus by spawning multiple
|
||||
child processes. Validate buffers allocated after running kernel code.
|
||||
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
|
||||
#ifdef __linux__
|
||||
#include <sys/wait.h>
|
||||
#include <sys/types.h>
|
||||
#include <unistd.h>
|
||||
|
||||
/**
|
||||
* Fetches Gpu device count
|
||||
*/
|
||||
static void getDeviceCount(int* pdevCnt) {
|
||||
int fd[2], val = 0;
|
||||
pid_t childpid;
|
||||
|
||||
// create pipe descriptors
|
||||
pipe(fd);
|
||||
|
||||
// disable visible_devices env from shell
|
||||
#ifdef HT_NVIDIA
|
||||
unsetenv("CUDA_VISIBLE_DEVICES");
|
||||
#else
|
||||
unsetenv("ROCR_VISIBLE_DEVICES");
|
||||
unsetenv("HIP_VISIBLE_DEVICES");
|
||||
#endif
|
||||
|
||||
childpid = fork();
|
||||
|
||||
if (childpid > 0) { // Parent
|
||||
close(fd[1]);
|
||||
// parent will wait to read the device cnt
|
||||
read(fd[0], &val, sizeof(val));
|
||||
|
||||
// close the read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
// wait for child exit
|
||||
wait(nullptr);
|
||||
|
||||
*pdevCnt = val;
|
||||
} else if (!childpid) { // Child
|
||||
int devCnt = 1;
|
||||
// writing only, no need for read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
// send the value on the write-descriptor:
|
||||
write(fd[1], &devCnt, sizeof(devCnt));
|
||||
|
||||
// close the write descriptor:
|
||||
close(fd[1]);
|
||||
exit(0);
|
||||
} else { // failure
|
||||
*pdevCnt = 0;
|
||||
return;
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates data consistency on supplied gpu
|
||||
*/
|
||||
static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
size_t prevAvl, prevTot, curAvl, curTot;
|
||||
bool TestPassed = true;
|
||||
constexpr auto N = 4 * 1024 * 1024;
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, 0, static_cast<const int*>(A_d),
|
||||
static_cast<const int*>(B_d), C_d, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
|
||||
printf("Validation PASSED for gpu %d from pid %d\n", gpu, getpid());
|
||||
} else {
|
||||
printf("Validation FAILED for gpu %d from pid %d\n", gpu, getpid());
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
|
||||
|
||||
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
|
||||
// In concurrent calls on one GPU, we cannot verify leaking in this way
|
||||
printf(
|
||||
"%s : Memory allocation mismatch observed."
|
||||
"Possible memory leak.\n",
|
||||
__func__);
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Parallel execution of parent and child on gpu0
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc_ChildConcurrencyDefaultGpu") {
|
||||
int devCnt = 0, pid = 0;
|
||||
constexpr auto resSuccess = 1, resFailure = 2;
|
||||
bool TestPassed = true;
|
||||
|
||||
// Get GPU count
|
||||
getDeviceCount(&devCnt);
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
if ((pid = fork()) < 0) {
|
||||
INFO("Child_Concurrency_DefaultGpu : fork() returned error : " << pid);
|
||||
HIP_ASSERT(false);
|
||||
|
||||
} else if (!pid) { // Child process
|
||||
bool TestPassedChild = false;
|
||||
|
||||
// Allocates and validates memory on Gpu0 simultaneously with parent
|
||||
TestPassedChild = validateMemoryOnGPU(0, true);
|
||||
|
||||
if (TestPassedChild) {
|
||||
exit(resSuccess); // child exit with success status
|
||||
} else {
|
||||
exit(resFailure); // child exit with failure status
|
||||
}
|
||||
|
||||
} else { // Parent process
|
||||
int exitStatus;
|
||||
|
||||
// Allocates and validates memory on Gpu0 simultaneously with child
|
||||
TestPassed = validateMemoryOnGPU(0, true);
|
||||
|
||||
// Wait and get result from child
|
||||
pid = wait(&exitStatus);
|
||||
if ((WEXITSTATUS(exitStatus) == resFailure) || (pid < 0))
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
REQUIRE(TestPassed == true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Parallel execution of api on multiple gpus from
|
||||
* different child processes.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc_ChildConcurrencyMultiGpu") {
|
||||
int devCnt = 0, pid = 0;
|
||||
constexpr auto resSuccess = 1, resFailure = 2;
|
||||
|
||||
// Get GPU count
|
||||
getDeviceCount(&devCnt);
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
// Spawn child for each GPU
|
||||
for (int gpu = 0; gpu < devCnt; gpu++) {
|
||||
if ((pid = fork()) < 0) {
|
||||
INFO("Child_Concurrency_MultiGpu : fork() returned error : " << pid);
|
||||
REQUIRE(false);
|
||||
|
||||
} else if (!pid) { // Child process
|
||||
bool TestPassedChild = false;
|
||||
TestPassedChild = validateMemoryOnGPU(gpu);
|
||||
|
||||
if (TestPassedChild) {
|
||||
exit(resSuccess); // child exit with success status
|
||||
} else {
|
||||
exit(resFailure); // child exit with failure status
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Parent shall wait for child to complete
|
||||
int passCnt = 0;
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
int pidwait = 0, exitStatus;
|
||||
pidwait = wait(&exitStatus);
|
||||
|
||||
printf("exitStatus for dev:%d is %d\n", i, WEXITSTATUS(exitStatus));
|
||||
if (pidwait < 0) {
|
||||
break;
|
||||
}
|
||||
|
||||
if (WEXITSTATUS(exitStatus) == resSuccess) passCnt++;
|
||||
}
|
||||
REQUIRE(passCnt == devCnt);
|
||||
}
|
||||
#endif // __linux__
|
||||
@@ -40,6 +40,11 @@ set(TEST_SRC
|
||||
hipMemset3D.cc
|
||||
hipMemset2D.cc
|
||||
hipMemset2DAsyncMultiThreadAndKernel.cc
|
||||
hipHostMallocTests.cc
|
||||
hipMallocConcurrency.cc
|
||||
hipMemset3DFunctional.cc
|
||||
hipMemset3DNegative.cc
|
||||
hipMemset3DRegressMultiThread.cc
|
||||
)
|
||||
else()
|
||||
set(TEST_SRC
|
||||
@@ -80,6 +85,11 @@ set(TEST_SRC
|
||||
hipMemset3D.cc
|
||||
hipMemset2D.cc
|
||||
hipMemset2DAsyncMultiThreadAndKernel.cc
|
||||
hipHostMallocTests.cc
|
||||
hipMallocConcurrency.cc
|
||||
hipMemset3DFunctional.cc
|
||||
hipMemset3DNegative.cc
|
||||
hipMemset3DRegressMultiThread.cc
|
||||
)
|
||||
endif()
|
||||
# Create shared lib of all tests
|
||||
|
||||
@@ -0,0 +1,61 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
|
||||
1) Test hipHostMalloc() api with ptr as nullptr and check for return value.
|
||||
2) Test hipHostMalloc() api with size as max(size_t) and check for OOM error.
|
||||
3) Test hipHostMalloc() api with flags as max(unsigned int) and validate
|
||||
return value.
|
||||
4) Pass size as zero for hipHostMalloc() api and check ptr is reset with
|
||||
with return value success.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
/**
|
||||
* Performs argument validation of hipHostMalloc api.
|
||||
*/
|
||||
TEST_CASE("Unit_hipHostMalloc_ArgValidation") {
|
||||
hipError_t ret;
|
||||
constexpr size_t allocSize = 1000;
|
||||
char *ptr;
|
||||
|
||||
SECTION("Pass ptr as nullptr") {
|
||||
ret = hipHostMalloc(static_cast<void **>(nullptr), allocSize);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Size as max(size_t)") {
|
||||
ret = hipHostMalloc(&ptr, std::numeric_limits<std::size_t>::max());
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Flags as max(uint)") {
|
||||
ret = hipHostMalloc(&ptr, allocSize,
|
||||
std::numeric_limits<unsigned int>::max());
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Pass size as zero and check ptr reset") {
|
||||
HIP_CHECK(hipHostMalloc(&ptr, 0));
|
||||
REQUIRE(ptr == nullptr);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,412 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
|
||||
1) Test hipMalloc() api passing zero size and confirming *ptr returning
|
||||
nullptr. Also pass nullptr to hipFree() api.
|
||||
|
||||
2) Pass maximum value of size_t for hipMalloc() api and make sure appropriate
|
||||
error is returned.
|
||||
|
||||
3) Check for hipMalloc() error code, passing invalid/null pointer.
|
||||
|
||||
4) Regress hipMalloc()/hipFree() in loop for bigger chunk of allocation
|
||||
with adequate number of iterations and later test for kernel execution on
|
||||
default gpu.
|
||||
|
||||
5) Regress hipMalloc()/hipFree() in loop while allocating smaller chunks
|
||||
keeping maximum number of iterations and then run kernel code on default
|
||||
gpu, perfom data validation.
|
||||
|
||||
6) Check hipMalloc() api adaptability when app creates small chunks of memory
|
||||
continuously, stores it for later use and then frees it at later point
|
||||
of time.
|
||||
|
||||
7) Multithread Scenario : Exercise hipMalloc() api parellely on all gpus from
|
||||
multiple threads and regress the api.
|
||||
|
||||
8) Validate memory usage with hipMemGetInfo() while regressing hipMalloc()
|
||||
api. Check for any possible memory leaks.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
|
||||
|
||||
#include <vector>
|
||||
#include <limits>
|
||||
#include <atomic>
|
||||
|
||||
|
||||
/* Buffer size for bigger chunks in alloc/free cycles */
|
||||
static constexpr auto BuffSizeBC = 5*1024*1024;
|
||||
|
||||
/* Buffer size for smaller chunks in alloc/free cycles */
|
||||
static constexpr auto BuffSizeSC = 16;
|
||||
|
||||
/* You may change it for individual test.
|
||||
* But default 100 is for quick return in Jenkin Build */
|
||||
static constexpr auto NumDiv = 100;
|
||||
|
||||
/* Max alloc/free iterations for smaller chunks */
|
||||
static constexpr auto MaxAllocFree_SmallChunks = (5000000/NumDiv);
|
||||
|
||||
/* Max alloc/free iterations for bigger chunks */
|
||||
static constexpr auto MaxAllocFree_BigChunks = 10000;
|
||||
|
||||
/* Max alloc and pool iterations */
|
||||
static constexpr auto MaxAllocPoolIter = (2000000/NumDiv);
|
||||
|
||||
/* Test status shared across threads */
|
||||
static std::atomic<bool> g_thTestPassed{true};
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* Validates data consistency on supplied gpu
|
||||
*/
|
||||
static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
size_t prevAvl, prevTot, curAvl, curTot;
|
||||
bool TestPassed = true;
|
||||
constexpr auto N = 4 * 1024 * 1024;
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, 0, static_cast<const int*>(A_d),
|
||||
static_cast<const int*>(B_d), C_d, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
|
||||
UNSCOPED_INFO("Validation PASSED for gpu " << gpu);
|
||||
} else {
|
||||
UNSCOPED_INFO("Validation FAILED for gpu " << gpu);
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
|
||||
|
||||
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
|
||||
// In concurrent calls on one GPU, we cannot verify leaking in this way
|
||||
UNSCOPED_INFO(
|
||||
"validateMemoryOnGPU : Memory allocation mismatch observed."
|
||||
<< "Possible memory leak.");
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Regress memory allocation and free in loop
|
||||
*/
|
||||
static bool regressAllocInLoop(int gpu) {
|
||||
bool TestPassed = true;
|
||||
size_t tot, avail, ptot, pavail, numBytes;
|
||||
int i = 0;
|
||||
int *ptr;
|
||||
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
numBytes = BuffSizeBC;
|
||||
|
||||
// Exercise allocation in loop with bigger chunks
|
||||
for (i = 0; i < MaxAllocFree_BigChunks; i++) {
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
HIP_CHECK(hipMalloc(&ptr, numBytes));
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
HIP_CHECK(hipFree(ptr));
|
||||
|
||||
if (pavail-avail < numBytes) { // We expect pavail-avail >= numBytes
|
||||
UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " <<
|
||||
numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" <<
|
||||
pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" <<
|
||||
tot << ", pavail-avail=" << pavail-avail);
|
||||
TestPassed = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
// Exercise allocation in loop with smaller chunks and maximum iters
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
numBytes = BuffSizeSC;
|
||||
|
||||
for (i = 0; i < MaxAllocFree_SmallChunks; i++) {
|
||||
HIP_CHECK(hipMalloc(&ptr, numBytes));
|
||||
|
||||
HIP_CHECK(hipFree(ptr));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
if ((pavail != avail) || (ptot != tot)) {
|
||||
UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." <<
|
||||
"Possible memory leak.");
|
||||
TestPassed &= false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates data consistency on supplied gpu
|
||||
* In Multithreaded Environment
|
||||
*/
|
||||
static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) {
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
size_t prevAvl, prevTot, curAvl, curTot;
|
||||
bool TestPassed = true;
|
||||
constexpr auto N = 4 * 1024 * 1024;
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
HIPCHECK(hipSetDevice(gpu));
|
||||
HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot));
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, 0, static_cast<const int*>(A_d),
|
||||
static_cast<const int*>(B_d), C_d, N);
|
||||
|
||||
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
|
||||
UNSCOPED_INFO("Validation PASSED for gpu " << gpu);
|
||||
} else {
|
||||
UNSCOPED_INFO("Validation FAILED for gpu " << gpu);
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipMemGetInfo(&curAvl, &curTot));
|
||||
|
||||
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
|
||||
// In concurrent calls on one GPU, we cannot verify leaking in this way
|
||||
UNSCOPED_INFO(
|
||||
"validateMemoryOnGpuMThread : Memory allocation mismatch observed."
|
||||
"Possible memory leak.");
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Regress memory allocation and free in loop
|
||||
* In Multithreaded Environment
|
||||
*/
|
||||
static bool regressAllocInLoopMthread(int gpu) {
|
||||
bool TestPassed = true;
|
||||
size_t tot, avail, ptot, pavail, numBytes;
|
||||
int i = 0;
|
||||
int *ptr;
|
||||
|
||||
HIPCHECK(hipSetDevice(gpu));
|
||||
numBytes = BuffSizeBC;
|
||||
|
||||
// Exercise allocation in loop with bigger chunks
|
||||
for (i = 0; i < MaxAllocFree_BigChunks; i++) {
|
||||
HIPCHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
HIPCHECK(hipMalloc(&ptr, numBytes));
|
||||
HIPCHECK(hipMemGetInfo(&avail, &tot));
|
||||
HIPCHECK(hipFree(ptr));
|
||||
|
||||
if (pavail-avail < numBytes) { // We expect pavail-avail >= numBytes
|
||||
UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " <<
|
||||
numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" <<
|
||||
pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" <<
|
||||
tot << ", pavail-avail=" << pavail-avail);
|
||||
TestPassed = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
// Exercise allocation in loop with smaller chunks and maximum iters
|
||||
HIPCHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
numBytes = BuffSizeSC;
|
||||
|
||||
for (i = 0; i < MaxAllocFree_SmallChunks; i++) {
|
||||
HIPCHECK(hipMalloc(&ptr, numBytes));
|
||||
|
||||
HIPCHECK(hipFree(ptr));
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
if ((pavail != avail) || (ptot != tot)) {
|
||||
UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." <<
|
||||
"Possible memory leak.");
|
||||
TestPassed &= false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/*
|
||||
* Thread func to regress alloc and check data consistency
|
||||
*/
|
||||
static void threadFunc(int gpu) {
|
||||
g_thTestPassed = regressAllocInLoopMthread(gpu);
|
||||
g_thTestPassed = g_thTestPassed & validateMemoryOnGpuMThread(gpu);
|
||||
|
||||
UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" <<
|
||||
g_thTestPassed.load());
|
||||
}
|
||||
|
||||
|
||||
/* Performs Argument Validation of api */
|
||||
TEST_CASE("Unit_hipMalloc_ArgumentValidation") {
|
||||
int *ptr;
|
||||
hipError_t ret;
|
||||
|
||||
SECTION("hipMalloc() when size(0)") {
|
||||
HIP_CHECK(hipMalloc(&ptr, 0));
|
||||
// ptr expected to be reset to null ptr
|
||||
REQUIRE(ptr == nullptr);
|
||||
}
|
||||
|
||||
SECTION("hipFree() when freeing nullptr ") {
|
||||
ptr = nullptr;
|
||||
// api should return success and shudnt crash
|
||||
HIP_CHECK(hipFree(ptr));
|
||||
}
|
||||
|
||||
SECTION("hipMalloc() with invalid argument") {
|
||||
constexpr auto sizeBytes = 100;
|
||||
ret = hipMalloc(nullptr, sizeBytes);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("hipMalloc() with max size_t") {
|
||||
ret = hipMalloc(&ptr, std::numeric_limits<std::size_t>::max());
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Regress hipMalloc()/hipFree() in loop for bigger chunks and
|
||||
* smaller chunks of memory allocation
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc_LoopRegressionAllocFreeCycles") {
|
||||
int devCnt = 0;
|
||||
|
||||
// Get GPU count
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
CHECK(regressAllocInLoop(0) == true);
|
||||
CHECK(validateMemoryOnGPU(0) == true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Application Behavior Modelling.
|
||||
* Check hipMalloc() api adaptability when app creates small chunks of memory
|
||||
* continuously, stores it for later use and then frees it at later point
|
||||
* of time.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
|
||||
size_t avail, tot, pavail, ptot;
|
||||
bool ret;
|
||||
hipError_t err;
|
||||
std::vector<int *> ptrlist;
|
||||
constexpr auto BuffSize = 10;
|
||||
int devCnt, *ptr;
|
||||
|
||||
// Get GPU count
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
|
||||
// Allocate small chunks of memory million times
|
||||
for (int i = 0; i < MaxAllocPoolIter ; i++) {
|
||||
if ((err = hipMalloc(&ptr, BuffSize)) != hipSuccess) {
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
INFO("Loop regression pool allocation failure. " <<
|
||||
"Total gpu memory " << tot/(1024.0*1024.0) <<", Free memory " <<
|
||||
avail/(1024.0*1024.0) << " iter " << i << " error "
|
||||
<< hipGetErrorString(err));
|
||||
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
// Store pointers allocated to emulate memory pool of app
|
||||
ptrlist.push_back(ptr);
|
||||
}
|
||||
|
||||
// Free ptrs at later point of time
|
||||
for ( auto &t : ptrlist ) {
|
||||
HIP_CHECK(hipFree(t));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
ret = validateMemoryOnGPU(0);
|
||||
REQUIRE(ret == true);
|
||||
REQUIRE(pavail == avail);
|
||||
REQUIRE(ptot == tot);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Exercise hipMalloc() api parellely on all gpus from
|
||||
* multiple threads and regress the api.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc_Multithreaded_MultiGPU") {
|
||||
std::vector<std::thread> threadlist;
|
||||
int devCnt;
|
||||
|
||||
// Get GPU count
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
threadlist.push_back(std::thread(threadFunc, i));
|
||||
}
|
||||
|
||||
for (auto &t : threadlist) {
|
||||
t.join();
|
||||
}
|
||||
|
||||
REQUIRE(g_thTestPassed == true);
|
||||
}
|
||||
@@ -0,0 +1,515 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
|
||||
1) Passing width as 0 in extent, verify hipMemset3D api returns success and
|
||||
doesn't modify the buffer passed.
|
||||
2) Passing width as 0 in extent, verify hipMemset3DAsync api returns success
|
||||
and doesn't modify the buffer passed.
|
||||
|
||||
3) Passing height as 0 in extent, verify hipMemset3D api returns success and
|
||||
doesn't modify the buffer passed.
|
||||
4) Passing height as 0 in extent, verify hipMemset3DAsync api returns success
|
||||
and doesn't modify the buffer passed.
|
||||
|
||||
5) Passing depth as 0 in extent, verify hipMemset3D api returns success and
|
||||
doesn't modify the buffer passed.
|
||||
6) Passing depth as 0 in extent, verify hipMemset3DAsync api returns success
|
||||
and doesn't modify the buffer passed.
|
||||
|
||||
7) When extent passed with width, height and depth all as zeroes, verify
|
||||
hipMemset3D api returns success and doesn't modify the buffer passed.
|
||||
8) When extent passed with width, height and depth all as zeroes, verify
|
||||
hipMemset3DAsync api returns success and doesn't modify the buffer passed.
|
||||
|
||||
9) Validate data after performing memory set operation with max memset value
|
||||
for hipMemset3D api.
|
||||
10) Validate data after performing memory set operation with max memset value
|
||||
for hipMemset3DAsync api.
|
||||
|
||||
11) Select random slice of 3d array and Memset complete slice with
|
||||
hipMemset3D api.
|
||||
12) Select random slice of 3d array and Memset complete slice with
|
||||
hipMemset3DAsync api.
|
||||
|
||||
13) Seek device pitched ptr to desired portion of 3d array and memset the
|
||||
portion with hipMemset3D api.
|
||||
14) Seek device pitched ptr to desired portion of 3d array and memset the
|
||||
portion with hipMemset3DAsync api.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
/*
|
||||
* Defines
|
||||
*/
|
||||
#define MEMSETVAL 1
|
||||
#define TESTVAL 2
|
||||
#define NUMH_EXT 256
|
||||
#define NUMW_EXT 100
|
||||
#define DEPTH_EXT 10
|
||||
#define NUMH_MAX 256
|
||||
#define NUMW_MAX 256
|
||||
#define DEPTH_MAX 10
|
||||
#define ZSIZE_S 32
|
||||
#define YSIZE_S 32
|
||||
#define XSIZE_S 32
|
||||
#define ZSIZE_P 30
|
||||
#define YSIZE_P 30
|
||||
#define XSIZE_P 30
|
||||
#define ZPOS_START 10
|
||||
#define ZSET_LEN 10
|
||||
#define ZPOS_END 19
|
||||
#define YPOS_START 10
|
||||
#define YSET_LEN 10
|
||||
#define YPOS_END 19
|
||||
#define XPOS_START 10
|
||||
#define XSET_LEN 10
|
||||
#define XPOS_END 19
|
||||
|
||||
|
||||
/**
|
||||
* Memset with extent passed and verify data to be intact
|
||||
*/
|
||||
static void testMemsetWithExtent(bool bAsync, hipExtent tstExtent) {
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
hipError_t ret;
|
||||
char *A_h;
|
||||
size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT;
|
||||
size_t width = numW * sizeof(char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
|
||||
A_h = reinterpret_cast<char *>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
memset(A_h, 0, sizeElements);
|
||||
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
if (bAsync) {
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
ret = hipMemset3DAsync(devPitchedPtr, MEMSETVAL, extent, stream);
|
||||
INFO("testMemsetWithExtent(" << extent.width << "," << extent.height
|
||||
<< "," << extent.depth << ") memset "
|
||||
<< MEMSETVAL << ", ret : " << ret);
|
||||
REQUIRE(ret == hipSuccess);
|
||||
|
||||
ret = hipMemset3DAsync(devPitchedPtr, TESTVAL, tstExtent, stream);
|
||||
INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height
|
||||
<< "," << tstExtent.depth << ") memset "
|
||||
<< TESTVAL << "ret : " << ret);
|
||||
REQUIRE(ret == hipSuccess);
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
} else {
|
||||
ret = hipMemset3D(devPitchedPtr, MEMSETVAL, extent);
|
||||
INFO("testMemsetWithExtent(" << extent.width << "," << extent.height
|
||||
<< "," << extent.depth << ") memset "
|
||||
<< MEMSETVAL << ",ret : " << ret);
|
||||
REQUIRE(ret == hipSuccess);
|
||||
|
||||
ret = hipMemset3D(devPitchedPtr, TESTVAL, tstExtent);
|
||||
INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height
|
||||
<< "," << tstExtent.depth << ") memset "
|
||||
<< TESTVAL << ",ret : " << ret);
|
||||
REQUIRE(ret == hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
|
||||
myparms.srcPtr = devPitchedPtr;
|
||||
myparms.extent = extent;
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
if (A_h[i] != MEMSETVAL) {
|
||||
INFO("testMemsetWithExtent: index:" << i << ",computed:"
|
||||
<< std::hex << static_cast<int>(A_h[i]) << ",memsetval:"
|
||||
<< std::hex << MEMSETVAL);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(devPitchedPtr.ptr));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Validates data after performing memory set operation with max memset value
|
||||
*/
|
||||
static void testMemsetMaxValue(bool bAsync) {
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
unsigned char *A_h;
|
||||
int memsetval = std::numeric_limits<unsigned char>::max();
|
||||
size_t numH = NUMH_MAX, numW = NUMW_MAX, depth = DEPTH_MAX;
|
||||
size_t width = numW * sizeof(unsigned char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
|
||||
A_h = reinterpret_cast<unsigned char *> (malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
memset(A_h, 0, sizeElements);
|
||||
|
||||
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
if (bAsync) {
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
} else {
|
||||
HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent));
|
||||
}
|
||||
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
|
||||
myparms.srcPtr = devPitchedPtr;
|
||||
myparms.extent = extent;
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
INFO("testMemsetMaxValue: index:" << i << ",computed:"
|
||||
<< std::hex << static_cast<int>(A_h[i]) << ",memsetval:"
|
||||
<< std::hex << memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipFree(devPitchedPtr.ptr));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* Function seeks device ptr to random slice and performs Memset operation
|
||||
* on the slice selected.
|
||||
*/
|
||||
static void seekAndSet3DArraySlice(bool bAsync) {
|
||||
char array3D[ZSIZE_S][YSIZE_S][XSIZE_S]{};
|
||||
dim3 arr_dimensions = dim3(ZSIZE_S, YSIZE_S, XSIZE_S);
|
||||
hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x,
|
||||
arr_dimensions.y, arr_dimensions.z);
|
||||
hipPitchedPtr devicePitchedPointer;
|
||||
int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL;
|
||||
|
||||
HIP_CHECK(hipMalloc3D(&devicePitchedPointer, extent));
|
||||
HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent));
|
||||
|
||||
// select random slice for memset
|
||||
unsigned int seed = time(nullptr);
|
||||
int slice_index = rand_r(&seed) % ZSIZE_S;
|
||||
|
||||
INFO("memset3d for sliceindex " << slice_index);
|
||||
|
||||
// Get attributes from device pitched pointer
|
||||
size_t pitch = devicePitchedPointer.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
|
||||
// Point devptr to selected slice
|
||||
char *devPtrSlice = (reinterpret_cast<char *>(devicePitchedPointer.ptr))
|
||||
+ slice_index * slicePitch;
|
||||
hipExtent extentSlice = make_hipExtent(sizeof(char) * arr_dimensions.x,
|
||||
arr_dimensions.y, 1);
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrSlice, pitch,
|
||||
arr_dimensions.x, arr_dimensions.y);
|
||||
|
||||
if (bAsync) {
|
||||
// Memset selected slice (Async)
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked,
|
||||
extentSlice, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
} else {
|
||||
// Memset selected slice
|
||||
HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, extentSlice));
|
||||
}
|
||||
|
||||
// Copy result back to host buffer
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x,
|
||||
arr_dimensions.x, arr_dimensions.y);
|
||||
myparms.srcPtr = devicePitchedPointer;
|
||||
myparms.extent = extent;
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (int z = 0; z < ZSIZE_S; z++) {
|
||||
for (int y = 0; y < YSIZE_S; y++) {
|
||||
for (int x = 0; x < XSIZE_S; x++) {
|
||||
if (z == slice_index) {
|
||||
if (array3D[z][y][x] != memsetval4seeked) {
|
||||
INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" << z
|
||||
<< "," << y << "," << x << ") " << "computed:" << std::hex
|
||||
<< array3D[z][y][x] << ", memsetval:" << std::hex
|
||||
<< memsetval4seeked);
|
||||
REQUIRE(false);
|
||||
}
|
||||
} else {
|
||||
if (array3D[z][y][x] != memsetval) {
|
||||
INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" << z
|
||||
<< "," << y << "," << x << ") " << "computed:" << std::hex
|
||||
<< array3D[z][y][x] << ", memsetval:" << std::hex
|
||||
<< memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(devicePitchedPointer.ptr));
|
||||
}
|
||||
|
||||
/**
|
||||
* Function seeks device ptr to selected portion of 3d array
|
||||
* and performs Memset operation on the portion.
|
||||
*/
|
||||
static void seekAndSet3DArrayPortion(bool bAsync) {
|
||||
char array3D[ZSIZE_P][YSIZE_P][XSIZE_P]{};
|
||||
dim3 arr_dimensions = dim3(ZSIZE_P, YSIZE_P, XSIZE_P);
|
||||
hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x,
|
||||
arr_dimensions.y, arr_dimensions.z);
|
||||
hipPitchedPtr devicePitchedPointer;
|
||||
int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL;
|
||||
|
||||
HIP_CHECK(hipMalloc3D(&devicePitchedPointer, extent));
|
||||
HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent));
|
||||
|
||||
// For memsetting extent/size(10,10,10) in the mid portion of cube(30,30,30),
|
||||
// seek device ptr to (10,10,10) and then memset 10 bytes across x,y,z axis.
|
||||
size_t pitch = devicePitchedPointer.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
int slice_index = ZPOS_START, y = YPOS_START, x = XPOS_START;
|
||||
|
||||
// Select 10th slice
|
||||
char *devPtrSlice = (reinterpret_cast<char *>(devicePitchedPointer.ptr))
|
||||
+ slice_index * slicePitch;
|
||||
|
||||
// Now select row at height as 10
|
||||
char *current_row = reinterpret_cast<char *>(devPtrSlice + y * pitch);
|
||||
|
||||
// Now select index of selected row as 10
|
||||
char *devPtrIndexed = ¤t_row[x];
|
||||
|
||||
// Make dev Pitchedptr, extent
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrIndexed, pitch,
|
||||
arr_dimensions.x, arr_dimensions.y);
|
||||
hipExtent setExtent = make_hipExtent(sizeof(char) * XSET_LEN, YSET_LEN,
|
||||
ZSET_LEN);
|
||||
|
||||
if (bAsync) {
|
||||
// Memset selected portion (Async)
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked,
|
||||
setExtent, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
} else {
|
||||
// Memset selected portion
|
||||
HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, setExtent));
|
||||
}
|
||||
|
||||
// Copy result back to host buffer
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x,
|
||||
arr_dimensions.x, arr_dimensions.y);
|
||||
myparms.srcPtr = devicePitchedPointer;
|
||||
myparms.extent = extent;
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (int z = 0; z < ZSIZE_P; z++) {
|
||||
for (int y = 0; y < YSIZE_P; y++) {
|
||||
for (int x = 0; x < XSIZE_P; x++) {
|
||||
if ((z >= ZPOS_START && z <= ZPOS_END) &&
|
||||
(y >= YPOS_START && y <= YPOS_END) &&
|
||||
(x >= XPOS_START && x <= XPOS_END)) {
|
||||
if (array3D[z][y][x] != memsetval4seeked) {
|
||||
INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" << z
|
||||
<< "," << y << "," << x << ") " << "computed:" << std::hex
|
||||
<< array3D[z][y][x] << ", memsetval:" << std::hex
|
||||
<< memsetval4seeked);
|
||||
REQUIRE(false);
|
||||
}
|
||||
} else {
|
||||
if (array3D[z][y][x] != memsetval) {
|
||||
INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" << z
|
||||
<< "," << y << "," << x << ") " << "computed:" << std::hex
|
||||
<< array3D[z][y][x] << ", memsetval:" << std::hex
|
||||
<< memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(devicePitchedPointer.ptr));
|
||||
}
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* Test Memset3D with different combinations of extent
|
||||
* taking zero and non-zero fields.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3D_MemsetWithExtent") {
|
||||
hipExtent testExtent;
|
||||
size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT;
|
||||
|
||||
SECTION("Memset with extent width(0)") {
|
||||
// Memset with extent width(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(0, numH, depth);
|
||||
testMemsetWithExtent(0, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent height(0)") {
|
||||
// Memset with extent height(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(numW, 0, depth);
|
||||
testMemsetWithExtent(0, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent depth(0)") {
|
||||
// Memset with extent depth(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(numW, numH, 0);
|
||||
testMemsetWithExtent(0, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent width,height,depth as 0") {
|
||||
// Memset with extent width,height,depth as 0 and verify data to be intact
|
||||
testExtent = make_hipExtent(0, 0, 0);
|
||||
testMemsetWithExtent(0, testExtent);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Test Memset3DAsync with different combinations of extent
|
||||
* taking zero and non-zero fields.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_MemsetWithExtent") {
|
||||
hipExtent testExtent;
|
||||
size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT;
|
||||
|
||||
SECTION("Memset with extent width(0)") {
|
||||
// Memset with extent width(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(0, numH, depth);
|
||||
testMemsetWithExtent(1, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent height(0)") {
|
||||
// Memset with extent height(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(numW, 0, depth);
|
||||
testMemsetWithExtent(1, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent depth(0)") {
|
||||
// Memset with extent depth(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(numW, numH, 0);
|
||||
testMemsetWithExtent(1, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent width,height,depth as 0") {
|
||||
// Memset with extent width,height,depth as 0 and verify data to be intact
|
||||
testExtent = make_hipExtent(0, 0, 0);
|
||||
testMemsetWithExtent(1, testExtent);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Memset3D with max unsigned char and verify memset operation is success
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3D_MemsetMaxValue") {
|
||||
testMemsetMaxValue(0);
|
||||
}
|
||||
|
||||
/**
|
||||
* Memset3DAsync with max unsigned char and verify memset operation is success
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_MemsetMaxValue") {
|
||||
testMemsetMaxValue(1);
|
||||
}
|
||||
|
||||
/**
|
||||
* Seek and set random slice of 3d array, verify memset is success
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3D_SeekSetSlice") {
|
||||
seekAndSet3DArraySlice(0);
|
||||
}
|
||||
|
||||
/**
|
||||
* Seek and set random slice of 3d array with async, verify memset is success
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_SeekSetSlice") {
|
||||
seekAndSet3DArraySlice(1);
|
||||
}
|
||||
|
||||
/**
|
||||
* Memset3D selected portion of 3d array
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3D_SeekSetArrayPortion") {
|
||||
seekAndSet3DArrayPortion(0);
|
||||
}
|
||||
|
||||
/**
|
||||
* Memset3DAsync selected portion of 3d array
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_SeekSetArrayPortion") {
|
||||
seekAndSet3DArrayPortion(1);
|
||||
}
|
||||
@@ -0,0 +1,236 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Test hipMemset3D() with uninitialized devPitchedPtr.
|
||||
2) Test hipMemset3DAsync() with uninitialized devPitchedPtr.
|
||||
|
||||
3) Reset devPitchedPtr to zero and check return value for hipMemset3D().
|
||||
4) Reset devPitchedPtr to zero and check return value for hipMemset3DAsync().
|
||||
|
||||
5) Test hipMemset3D() with extent.width as max size_t and keeping height,
|
||||
depth as valid values.
|
||||
6) Test hipMemset3DAsync() with extent.width as max size_t and keeping height,
|
||||
depth as valid values.
|
||||
7) Test hipMemset3D() with extent.height as max size_t and keeping width,
|
||||
depth as valid values.
|
||||
8) Test hipMemset3DAsync() with extent.height as max size_t and keeping width,
|
||||
depth as valid values.
|
||||
9) Test hipMemset3D() with extent.depth as max size_t and keeping height,
|
||||
width as valid values.
|
||||
10) Test hipMemset3DAsync() with extent.depth as max size_t and keeping height,
|
||||
width as valid values.
|
||||
|
||||
11) Device Ptr out bound and extent(0) passed for hipMemset3D().
|
||||
12) Device Ptr out bound and extent(0) passed for hipMemset3DAsync().
|
||||
|
||||
13) Device Ptr out bound and valid extent passed for hipMemset3D().
|
||||
14) Device Ptr out bound and valid extent passed for hipMemset3DAsync().
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
TEST_CASE("Unit_hipMemset3D_Negative") {
|
||||
hipError_t ret;
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
constexpr int memsetval = 1;
|
||||
constexpr size_t numH = 256, numW = 256;
|
||||
constexpr size_t depth = 10;
|
||||
constexpr size_t width = numW * sizeof(char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
|
||||
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
|
||||
SECTION("Using uninitialized devpitched ptr") {
|
||||
hipPitchedPtr devPitchedUnPtr;
|
||||
|
||||
ret = hipMemset3D(devPitchedUnPtr, memsetval, extent);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Reset devPitchedPtr to zero") {
|
||||
hipPitchedPtr rdevPitchedPtr{};
|
||||
|
||||
ret = hipMemset3D(rdevPitchedPtr, memsetval, extent);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Pass extent fields as max size_t") {
|
||||
hipExtent extMW = make_hipExtent(std::numeric_limits<std::size_t>::max(),
|
||||
numH,
|
||||
depth);
|
||||
hipExtent extMH = make_hipExtent(width,
|
||||
std::numeric_limits<std::size_t>::max(),
|
||||
depth);
|
||||
hipExtent extMD = make_hipExtent(width,
|
||||
numH,
|
||||
std::numeric_limits<std::size_t>::max());
|
||||
|
||||
ret = hipMemset3D(devPitchedPtr, memsetval, extMW);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
|
||||
ret = hipMemset3D(devPitchedPtr, memsetval, extMH);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
|
||||
if ((TestContext::get()).isAmd()) {
|
||||
ret = hipMemset3D(devPitchedPtr, memsetval, extMD);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
} else {
|
||||
WARN("Test is skipped for max depth."
|
||||
<< "Cuda doesn't check the maximum depth of extent field");
|
||||
}
|
||||
}
|
||||
|
||||
SECTION("Device Ptr out bound and extent(0) passed for memset") {
|
||||
size_t pitch = devPitchedPtr.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
constexpr auto advanceOffset = 10;
|
||||
|
||||
// Point devptr to end of allocated memory
|
||||
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
|
||||
+ depth * slicePitch;
|
||||
|
||||
// Advance devptr further to go out of boundary
|
||||
devPtrMod = devPtrMod + advanceOffset;
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
|
||||
numW * sizeof(char), numH);
|
||||
hipExtent extent0{};
|
||||
ret = hipMemset3D(modDevPitchedPtr, memsetval, extent0);
|
||||
|
||||
// api expected to check extent0 and return success before going for ptr.
|
||||
REQUIRE(ret == hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Device Ptr out bound and valid extent passed for memset") {
|
||||
size_t pitch = devPitchedPtr.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
constexpr auto advanceOffset = 10;
|
||||
|
||||
// Point devptr to end of allocated memory
|
||||
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
|
||||
+ depth * slicePitch;
|
||||
|
||||
// Advance devptr further to go out of boundary
|
||||
devPtrMod = devPtrMod + advanceOffset;
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
|
||||
numW * sizeof(char), numH);
|
||||
ret = hipMemset3D(modDevPitchedPtr, memsetval, extent);
|
||||
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(devPitchedPtr.ptr));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemset3DAsync_Negative") {
|
||||
hipError_t ret;
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
hipStream_t stream;
|
||||
constexpr int memsetval = 1;
|
||||
constexpr size_t numH = 256;
|
||||
constexpr size_t numW = 256;
|
||||
constexpr size_t depth = 10;
|
||||
constexpr size_t width = numW * sizeof(char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
|
||||
SECTION("Using uninitialized devpitched ptr") {
|
||||
hipPitchedPtr devPitchedUnPtr;
|
||||
|
||||
ret = hipMemset3DAsync(devPitchedUnPtr, memsetval, extent, stream);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Reset devPitchedPtr to zero") {
|
||||
hipPitchedPtr rdevPitchedPtr{};
|
||||
|
||||
ret = hipMemset3DAsync(rdevPitchedPtr, memsetval, extent, stream);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Pass extent fields as max size_t") {
|
||||
hipExtent extMW = make_hipExtent(std::numeric_limits<std::size_t>::max(),
|
||||
numH,
|
||||
depth);
|
||||
hipExtent extMH = make_hipExtent(width,
|
||||
std::numeric_limits<std::size_t>::max(),
|
||||
depth);
|
||||
hipExtent extMD = make_hipExtent(width,
|
||||
numH,
|
||||
std::numeric_limits<std::size_t>::max());
|
||||
|
||||
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMW, stream);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
|
||||
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMH, stream);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
|
||||
if ((TestContext::get()).isAmd()) {
|
||||
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMD, stream);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
} else {
|
||||
WARN("Test is skipped for max depth."
|
||||
<< "Cuda doesn't check the maximum depth of extent field");
|
||||
}
|
||||
}
|
||||
|
||||
SECTION("Device Ptr out bound and extent(0) passed for memset") {
|
||||
size_t pitch = devPitchedPtr.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
constexpr auto advanceOffset = 10;
|
||||
|
||||
// Point devptr to end of allocated memory
|
||||
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
|
||||
+ depth * slicePitch;
|
||||
|
||||
// Advance devptr further to go out of boundary
|
||||
devPtrMod = devPtrMod + advanceOffset;
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
|
||||
numW * sizeof(char), numH);
|
||||
hipExtent extent0{};
|
||||
ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent0, stream);
|
||||
|
||||
// api expected to check extent0 and return success before going for ptr.
|
||||
REQUIRE(ret == hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Device Ptr out bound and valid extent passed for memset") {
|
||||
size_t pitch = devPitchedPtr.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
constexpr auto advanceOffset = 10;
|
||||
|
||||
// Point devptr to end of allocated memory
|
||||
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
|
||||
+ depth * slicePitch;
|
||||
|
||||
// Advance devptr further to go out of boundary
|
||||
devPtrMod = devPtrMod + advanceOffset;
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
|
||||
numW * sizeof(char), numH);
|
||||
ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent, stream);
|
||||
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK(hipFree(devPitchedPtr.ptr));
|
||||
}
|
||||
@@ -0,0 +1,266 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Validate Async behavior of hipMemset3DAsync with commands queued
|
||||
concurrently from multiple threads.
|
||||
2) Validate hipMemset3DAsync behavior when api is queued along with kernel
|
||||
function operating on same memory.
|
||||
3) Perform regression of hipMemset3D api in loop with device memory allocated
|
||||
on different gpus.
|
||||
4) Perform regression of hipMemset3DAsync api in loop with device memory
|
||||
allocated on different gpus.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
|
||||
/*
|
||||
* Defines
|
||||
*/
|
||||
#define MAX_REGRESS_ITERS 2
|
||||
#define MAX_THREADS 10
|
||||
|
||||
/**
|
||||
* kernel function sets device memory with value passed
|
||||
*/
|
||||
static __global__ void func_set_value(hipPitchedPtr devicePitchedPointer,
|
||||
hipExtent extent,
|
||||
unsigned char val) {
|
||||
// Index Calculation
|
||||
size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
size_t y = threadIdx.y + blockDim.y * blockIdx.y;
|
||||
size_t z = threadIdx.z + blockDim.z * blockIdx.z;
|
||||
|
||||
// Get attributes from device pitched pointer
|
||||
char *devicePointer = reinterpret_cast<char *>(devicePitchedPointer.ptr);
|
||||
size_t pitch = devicePitchedPointer.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
|
||||
// Loop over the device buffer
|
||||
if (z < extent.depth) {
|
||||
char *current_slice_index = devicePointer + z * slicePitch;
|
||||
if (y < extent.height) {
|
||||
// Get data array containing all elements from the current row
|
||||
char *current_row = reinterpret_cast<char *>(current_slice_index
|
||||
+ y * pitch);
|
||||
if (x < extent.width) {
|
||||
current_row[x] = val;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Thread function queues kernel function and memset cmds
|
||||
*/
|
||||
static void threadFunc(hipStream_t stream, hipPitchedPtr devpPtr,
|
||||
int memsetval, int testval, hipExtent extent, hipMemcpy3DParms myparms) {
|
||||
// Kernel Launch Configuration
|
||||
constexpr auto size = 8;
|
||||
dim3 threadsPerBlock = dim3(size, size, size);
|
||||
dim3 blocks;
|
||||
blocks = dim3((extent.width + threadsPerBlock.x - 1) / threadsPerBlock.x,
|
||||
(extent.height + threadsPerBlock.y - 1) / threadsPerBlock.y,
|
||||
(extent.depth + threadsPerBlock.z - 1) / threadsPerBlock.z);
|
||||
|
||||
hipLaunchKernelGGL(func_set_value, dim3(blocks), dim3(threadsPerBlock), 0,
|
||||
stream, devpPtr, extent, memsetval);
|
||||
HIPCHECK(hipMemset3DAsync(devpPtr, testval, extent, stream));
|
||||
HIPCHECK(hipMemcpy3DAsync(&myparms, stream));
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Performs api regression in loop
|
||||
*/
|
||||
bool loopRegression(bool bAsync) {
|
||||
bool testPassed = true;
|
||||
char *A_h;
|
||||
constexpr int memsetval = 1;
|
||||
constexpr size_t numH = 256, numW = 100, depth = 10;
|
||||
int numGpu = 0, hasPeerAccess = 0;
|
||||
size_t width = numW * sizeof(char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
std::vector<hipPitchedPtr> devPitchedPtrlist;
|
||||
hipPitchedPtr pitchedPtr, devpPtr;
|
||||
|
||||
A_h = reinterpret_cast<char *>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
memset(A_h, 0, sizeElements);
|
||||
|
||||
// Populate hipMemcpy3D parameters
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
|
||||
myparms.extent = extent;
|
||||
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipGetDeviceCount(&numGpu));
|
||||
REQUIRE(numGpu > 0);
|
||||
|
||||
// Alloc 3D arrays in all GPUs
|
||||
for (int j = 0; j < numGpu; j++) {
|
||||
HIP_CHECK(hipSetDevice(j));
|
||||
HIP_CHECK(hipMalloc3D(&pitchedPtr, extent));
|
||||
devPitchedPtrlist.push_back(pitchedPtr);
|
||||
}
|
||||
|
||||
for (int itern = 0; itern < MAX_REGRESS_ITERS; itern++) {
|
||||
// Validate hipMemset3D data consistency in multiple iters
|
||||
for (int i = 0; i < numGpu; i++) {
|
||||
for (int j = 0; j < numGpu; j++) {
|
||||
HIP_CHECK(hipDeviceCanAccessPeer(&hasPeerAccess, i, j));
|
||||
if (!hasPeerAccess) {
|
||||
// Skip and continue if no peer access
|
||||
continue;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipSetDevice(i));
|
||||
devpPtr = devPitchedPtrlist[j];
|
||||
HIP_CHECK(hipDeviceEnablePeerAccess(j, 0));
|
||||
HIP_CHECK(hipMemset3D(devpPtr, 0, extent));
|
||||
|
||||
if (bAsync) {
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMemset3DAsync(devpPtr, memsetval, extent, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
} else {
|
||||
HIP_CHECK(hipMemset3D(devpPtr, memsetval, extent));
|
||||
}
|
||||
|
||||
myparms.srcPtr = devpPtr;
|
||||
memset(A_h, 0, sizeElements);
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (size_t indx = 0; indx < elements; indx++) {
|
||||
if (A_h[indx] != memsetval) {
|
||||
testPassed = false;
|
||||
printf("RegressIter : mismatch at index:%d computed:%02x, "
|
||||
"memsetval:%02x\n", static_cast<int>(indx),
|
||||
static_cast<int>(A_h[indx]), static_cast<int>(memsetval));
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
for (int j = 0; j < numGpu; j++) {
|
||||
HIP_CHECK(hipFree(devPitchedPtrlist[j].ptr));
|
||||
}
|
||||
|
||||
free(A_h);
|
||||
return testPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Perform regression of hipMemset3D api with device memory allocated
|
||||
* on different gpus.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3D_RegressInLoop") {
|
||||
bool TestPassed = false;
|
||||
|
||||
TestPassed = loopRegression(0);
|
||||
REQUIRE(TestPassed == true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Perform regression of hipMemset3DAsync api with device memory allocated
|
||||
* on different gpus.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_RegressInLoop") {
|
||||
bool TestPassed = false;
|
||||
|
||||
TestPassed = loopRegression(1);
|
||||
REQUIRE(TestPassed == true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Async commands queued concurrently and executed
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_ConcurrencyMthread") {
|
||||
char *A_h;
|
||||
constexpr int memsetval = 1, testval = 2;
|
||||
constexpr size_t numH = 256, numW = 100, depth = 10;
|
||||
size_t width = numW * sizeof(char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
hipPitchedPtr devpPtr;
|
||||
hipStream_t stream;
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMalloc3D(&devpPtr, extent));
|
||||
|
||||
A_h = reinterpret_cast<char *>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
memset(A_h, 0, sizeElements);
|
||||
|
||||
// Populate hipMemcpy3D parameters
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.srcPtr = devpPtr;
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
|
||||
myparms.extent = extent;
|
||||
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
std::vector<std::thread> threadlist;
|
||||
|
||||
// Queue cmds concurrently from multiple threads on same stream
|
||||
for (int i = 0; i < MAX_THREADS; i++) {
|
||||
threadlist.push_back(std::thread(threadFunc, stream, devpPtr, memsetval,
|
||||
testval, extent, myparms));
|
||||
}
|
||||
|
||||
for (auto &t : threadlist) {
|
||||
t.join();
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
|
||||
for (size_t k = 0 ; k < elements ; k++) {
|
||||
if (A_h[k] != testval) {
|
||||
CAPTURE(A_h[k], testval, k);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
free(A_h);
|
||||
HIP_CHECK(hipFree(devpPtr.ptr));
|
||||
}
|
||||
Reference in New Issue
Block a user