SWDEV-299127 - Merge 'develop' into 'amd-staging'
Change-Id: Ib604cf603ab63b998130ef1c769c53d42328bace
Этот коммит содержится в:
Jenkins
@@ -172,5 +172,8 @@ cmake ..
|
||||
make
|
||||
./test_*.out
|
||||
```
|
||||
It is recommended to use Visual Studio's command prompt for this sample due to requirement of MS Librarian tool - LIB.exe on windows platform.
|
||||
Override CMAKE_C_COMPILER and CMAKE_CXX_COMPILER to hipcc as Visual Studio's compiler would use cl.exe as default compiler.
|
||||
i.e. cmake.exe -GNinja -DCMAKE_CXX_COMPILER_ID=ROCMClang -DCMAKE_C_COMPILER_ID=ROCMClang -DCMAKE_PREFIX_PATH=%HIP_PATH% -DCMAKE_C_COMPILER=%HIP_PATH%/bin/hipcc.bat -DCMAKE_CXX_COMPILER=%HIP_PATH%/bin/hipcc.bat ..
|
||||
|
||||
## For More Infomation, please refer to the HIP FAQ.
|
||||
|
||||
@@ -12,6 +12,12 @@ list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
# Find hip
|
||||
find_package(hip REQUIRED)
|
||||
|
||||
# For windows, AR is MS Librarian and that is picked by Visual Studio's command prompt.
|
||||
if (WIN32)
|
||||
find_program(libpath NAMES lib.exe)
|
||||
set (CMAKE_AR ${libpath})
|
||||
endif()
|
||||
|
||||
# Set compiler and linker
|
||||
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
|
||||
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
|
||||
@@ -22,6 +28,11 @@ option(BUILD_SHARED_LIBS "Build as a shared library" OFF)
|
||||
|
||||
set(CPP_SOURCES hipDevice.cpp)
|
||||
|
||||
# For windows, We need to tell cmake how to create static library.
|
||||
if (WIN32)
|
||||
set (CMAKE_CXX_CREATE_STATIC_LIBRARY "<CMAKE_AR> /out:<TARGET> <LINK_FLAGS> <OBJECTS>")
|
||||
endif()
|
||||
|
||||
# Generate static lib libHipDevice.a
|
||||
add_library(HipDevice STATIC ${CPP_SOURCES})
|
||||
|
||||
@@ -35,6 +46,13 @@ set(TEST_SOURCES ${CMAKE_SOURCE_DIR}/hipMain2.cpp)
|
||||
add_executable(test_device_static ${TEST_SOURCES})
|
||||
add_dependencies(test_device_static HipDevice)
|
||||
target_compile_options(test_device_static PRIVATE -fgpu-rdc)
|
||||
target_link_libraries(test_device_static PRIVATE HipDevice)
|
||||
target_link_libraries(test_device_static PRIVATE -fgpu-rdc hip::host)
|
||||
|
||||
# For windows, Change in a way to pass lib details
|
||||
if (WIN32)
|
||||
target_link_libraries(test_device_static PRIVATE -lHipDevice -L${CMAKE_BINARY_DIR})
|
||||
else()
|
||||
target_link_libraries(test_device_static PRIVATE HipDevice)
|
||||
endif()
|
||||
|
||||
target_link_libraries(test_device_static PRIVATE -fgpu-rdc amdhip64 amd_comgr)
|
||||
|
||||
|
||||
@@ -12,10 +12,15 @@ list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
# Find hip
|
||||
find_package(hip REQUIRED)
|
||||
|
||||
# For windows, AR is MS Librarian and that is pickedby Visual Studio's command prompt.
|
||||
if (WIN32)
|
||||
find_program(libpath NAMES lib.exe)
|
||||
set (CMAKE_AR ${libpath})
|
||||
endif()
|
||||
|
||||
# Set compiler and linker
|
||||
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
|
||||
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
|
||||
set(CMAKE_AR ${HIP_HIPCC_EXECUTABLE})
|
||||
set(CMAKE_BUILD_TYPE Release)
|
||||
|
||||
# Turn static library generation ON
|
||||
@@ -23,15 +28,17 @@ option(BUILD_SHARED_LIBS "Build as a shared library" OFF)
|
||||
|
||||
set(CPP_SOURCES hipOptLibrary.cpp)
|
||||
|
||||
# For windows, We need to tell cmake how to create static library.
|
||||
if (WIN32)
|
||||
set (CMAKE_CXX_CREATE_STATIC_LIBRARY "<CMAKE_AR> /out:<TARGET> <LINK_FLAGS> <OBJECTS>")
|
||||
endif()
|
||||
|
||||
# Generate static lib libHipOptLibrary.a.
|
||||
add_library(HipOptLibrary STATIC ${CPP_SOURCES})
|
||||
|
||||
# Set-up the correct flags to generate the static library.
|
||||
target_link_libraries(HipOptLibrary PRIVATE --emit-static-lib)
|
||||
target_include_directories(HipOptLibrary PRIVATE /opt/rocm/hsa/include)
|
||||
get_property(link_libraries TARGET HipOptLibrary PROPERTY LINK_LIBRARIES)
|
||||
string (REPLACE ";" " " LINK_PROPS "${link_libraries}")
|
||||
set(CMAKE_CXX_ARCHIVE_CREATE "<CMAKE_AR> -o <TARGET> ${LINK_PROPS} <LINK_FLAGS> <OBJECTS>")
|
||||
|
||||
# Create test executable that uses libHipOptLibrary.a
|
||||
set(TEST_SOURCES ${CMAKE_SOURCE_DIR}/hipMain1.cpp)
|
||||
@@ -39,5 +46,10 @@ set(TEST_SOURCES ${CMAKE_SOURCE_DIR}/hipMain1.cpp)
|
||||
add_executable(test_opt_static ${TEST_SOURCES})
|
||||
add_dependencies(test_opt_static HipOptLibrary)
|
||||
target_link_libraries(test_opt_static PRIVATE -lHipOptLibrary -L${CMAKE_BINARY_DIR})
|
||||
target_link_libraries(test_opt_static PRIVATE amdhip64 amd_comgr hsa-runtime64::hsa-runtime64)
|
||||
|
||||
if (WIN32)
|
||||
target_link_libraries(test_opt_static PRIVATE amdhip64 amd_comgr)
|
||||
else()
|
||||
target_link_libraries(test_opt_static PRIVATE amdhip64 amd_comgr hsa-runtime64::hsa-runtime64)
|
||||
endif()
|
||||
|
||||
|
||||
@@ -14,6 +14,7 @@ set(LINUX_TEST_SRC
|
||||
hipMemCoherencyTstMProc.cc
|
||||
hipIpcEventHandle.cc
|
||||
hipIpcMemAccessTest.cc
|
||||
deviceAllocationMproc.cc
|
||||
)
|
||||
|
||||
# the last argument linker libraries is required for this test but optional to the function
|
||||
|
||||
@@ -0,0 +1,319 @@
|
||||
/*
|
||||
Copyright (c) 2022 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_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
|
||||
#ifdef __linux__
|
||||
#include <unistd.h>
|
||||
#include <stdlib.h>
|
||||
#include <sys/wait.h>
|
||||
#include <dlfcn.h>
|
||||
#endif
|
||||
|
||||
#define SIZE 2097152
|
||||
// GPU threads
|
||||
#define BLOCKSIZE 512
|
||||
#define GRIDSIZE 256
|
||||
|
||||
__device__ static char* dev_common_ptr = nullptr;
|
||||
|
||||
/**
|
||||
* This kernel allocates a memory chunk using malloc().
|
||||
*/
|
||||
static __global__ void kerTestDeviceMalloc(size_t size) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Allocate
|
||||
if (myId == 0) {
|
||||
dev_common_ptr = reinterpret_cast<char*> (malloc(size));
|
||||
if (dev_common_ptr == nullptr) {
|
||||
printf("Device Allocation Failed! \n");
|
||||
return;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* This kernel writes to the memory location allocated in kernel
|
||||
* kerTestDeviceMalloc or kerTestDeviceNew.
|
||||
*/
|
||||
static __global__ void kerTestDeviceWrite() {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Allocate
|
||||
if (dev_common_ptr == nullptr) {
|
||||
printf("Device Allocation Failed! \n");
|
||||
return;
|
||||
}
|
||||
*(dev_common_ptr + myId) = SCHAR_MAX;
|
||||
}
|
||||
|
||||
/**
|
||||
* This kernel frees the memory chunk allocated in kernel
|
||||
* kerTestDeviceMalloc using free().
|
||||
*/
|
||||
static __global__ void kerTestDeviceFree(int *result) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Allocate
|
||||
if (myId == 0) {
|
||||
if (dev_common_ptr != nullptr) {
|
||||
*result = 1;
|
||||
for (int idx = 0; idx < (BLOCKSIZE*GRIDSIZE); idx++) {
|
||||
if (*(dev_common_ptr + myId) != SCHAR_MAX) {
|
||||
*result = 0;
|
||||
break;
|
||||
}
|
||||
}
|
||||
free(dev_common_ptr);
|
||||
} else {
|
||||
*result = 0;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* This kernel allocates a memory chunk using new operator.
|
||||
*/
|
||||
static __global__ void kerTestDeviceNew(size_t size) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Allocate
|
||||
if (myId == 0) {
|
||||
dev_common_ptr = new char[size];
|
||||
if (dev_common_ptr == nullptr) {
|
||||
printf("Device Allocation Failed! \n");
|
||||
return;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* This kernel frees the memory chunk allocated in kernel
|
||||
* kerTestDeviceNew using delete operator.
|
||||
*/
|
||||
static __global__ void kerTestDeviceDelete(int *result) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Allocate
|
||||
if (myId == 0) {
|
||||
if (dev_common_ptr != nullptr) {
|
||||
*result = 1;
|
||||
for (int idx = 0; idx < (BLOCKSIZE*GRIDSIZE); idx++) {
|
||||
if (*(dev_common_ptr + myId) != SCHAR_MAX) {
|
||||
*result = 0;
|
||||
break;
|
||||
}
|
||||
}
|
||||
delete[] dev_common_ptr;
|
||||
} else {
|
||||
*result = 0;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test device malloc()/new in both Parent and Child Process.
|
||||
* Allocate SIZE bytes in both parent and child process. Verify
|
||||
* the allocated size in both parent and child process.
|
||||
*/
|
||||
static bool testDeviceAllocMulProc(bool testmalloc) {
|
||||
int fd[2];
|
||||
pid_t childpid;
|
||||
bool testResult = false;
|
||||
size_t avail = 0, tot = 0;
|
||||
// create pipe descriptors
|
||||
pipe(fd);
|
||||
// fork process
|
||||
childpid = fork();
|
||||
if (childpid > 0) { // Parent
|
||||
close(fd[1]);
|
||||
// Allocate in parent
|
||||
if (testmalloc) {
|
||||
kerTestDeviceMalloc<<<1, 1>>>(SIZE);
|
||||
} else {
|
||||
kerTestDeviceNew<<<1, 1>>>(SIZE);
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
// Check allocated memory size
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
if ((tot - avail) < SIZE) {
|
||||
return false;
|
||||
}
|
||||
// parent will wait to read the device cnt
|
||||
read(fd[0], &testResult, sizeof(testResult));
|
||||
// close the read-descriptor
|
||||
close(fd[0]);
|
||||
// wait for child exit
|
||||
wait(NULL);
|
||||
// At this point the child process exits.
|
||||
// Ensure that device memory allocated from child is freed.
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
if ((tot - avail) < SIZE) {
|
||||
testResult = false;
|
||||
}
|
||||
} else if (!childpid) { // Child
|
||||
// Wait for hipDeviceSetLimit() completion in parent.
|
||||
close(fd[0]);
|
||||
// Allocate in child
|
||||
if (testmalloc) {
|
||||
kerTestDeviceMalloc<<<1, 1>>>(SIZE);
|
||||
} else {
|
||||
kerTestDeviceNew<<<1, 1>>>(SIZE);
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
// Check allocated memory size
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
if ((tot - avail) < SIZE) {
|
||||
testResult = false;
|
||||
} else {
|
||||
testResult = true;
|
||||
}
|
||||
// send the value on the write-descriptor:
|
||||
write(fd[1], &testResult, sizeof(testResult));
|
||||
// close the write descriptor:
|
||||
close(fd[1]);
|
||||
exit(0);
|
||||
}
|
||||
return testResult;
|
||||
}
|
||||
|
||||
/**
|
||||
* Test device malloc()/new, write and free()/delete[]
|
||||
* from both Parent and Child Process. From both Parent and
|
||||
* Child Process invoke the kernel to allocate memory, the
|
||||
* kernel to write to the allocated memory and a third kernel
|
||||
* to verify the memory contents and free it.
|
||||
*/
|
||||
static bool testDeviceMemMulProc(bool testmalloc) {
|
||||
int fd[2];
|
||||
bool testResult = false;
|
||||
pid_t childpid;
|
||||
int testResultChild = 0;
|
||||
size_t size = BLOCKSIZE*GRIDSIZE;
|
||||
// create pipe descriptors
|
||||
pipe(fd);
|
||||
// fork process
|
||||
childpid = fork();
|
||||
if (childpid > 0) { // Parent
|
||||
close(fd[1]);
|
||||
int *result_d{nullptr}, *result_h{nullptr};
|
||||
HIP_CHECK(hipMalloc(&result_d, sizeof(int)));
|
||||
result_h = reinterpret_cast<int*> (malloc(sizeof(int)));
|
||||
REQUIRE(result_h != nullptr);
|
||||
// Allocate in parent
|
||||
if (testmalloc) {
|
||||
kerTestDeviceMalloc<<<1, 1>>>(size);
|
||||
} else {
|
||||
kerTestDeviceNew<<<1, 1>>>(size);
|
||||
}
|
||||
// Write
|
||||
kerTestDeviceWrite<<<GRIDSIZE, BLOCKSIZE>>>();
|
||||
// Free
|
||||
if (testmalloc) {
|
||||
kerTestDeviceFree<<<1, 1>>>(result_d);
|
||||
} else {
|
||||
kerTestDeviceDelete<<<1, 1>>>(result_d);
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
*result_h = 0;
|
||||
HIP_CHECK(hipMemcpy(result_h, result_d, sizeof(int),
|
||||
hipMemcpyDefault));
|
||||
if (*result_h == 0) {
|
||||
testResult = false;
|
||||
} else {
|
||||
testResult = true;
|
||||
}
|
||||
// parent will wait to read the device cnt
|
||||
read(fd[0], &testResultChild, sizeof(int));
|
||||
if (testResultChild == 0) {
|
||||
testResult &= false;
|
||||
} else {
|
||||
testResult &= true;
|
||||
}
|
||||
// close the read-descriptor
|
||||
close(fd[0]);
|
||||
hipFree(result_d);
|
||||
free(result_h);
|
||||
// wait for child exit
|
||||
wait(NULL);
|
||||
} else if (!childpid) { // Child
|
||||
// Wait for hipDeviceSetLimit() completion in parent.
|
||||
close(fd[0]);
|
||||
int *result_d{nullptr}, *result_h{nullptr};
|
||||
HIP_CHECK(hipMalloc(&result_d, sizeof(int)));
|
||||
result_h = reinterpret_cast<int*> (malloc(sizeof(int)));
|
||||
REQUIRE(result_h != nullptr);
|
||||
// Allocate in child
|
||||
if (testmalloc) {
|
||||
kerTestDeviceMalloc<<<1, 1>>>(size);
|
||||
} else {
|
||||
kerTestDeviceNew<<<1, 1>>>(size);
|
||||
}
|
||||
// Write
|
||||
kerTestDeviceWrite<<<GRIDSIZE, BLOCKSIZE>>>();
|
||||
// Free
|
||||
if (testmalloc) {
|
||||
kerTestDeviceFree<<<1, 1>>>(result_d);
|
||||
} else {
|
||||
kerTestDeviceDelete<<<1, 1>>>(result_d);
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
*result_h = 0;
|
||||
HIP_CHECK(hipMemcpy(result_h, result_d, sizeof(int),
|
||||
hipMemcpyDefault));
|
||||
// send the value on the write-descriptor:
|
||||
write(fd[1], result_h, sizeof(int));
|
||||
// close the write descriptor:
|
||||
close(fd[1]);
|
||||
hipFree(result_d);
|
||||
free(result_h);
|
||||
exit(0);
|
||||
}
|
||||
return testResult;
|
||||
}
|
||||
|
||||
/**
|
||||
* Multiprocess device side malloc test.
|
||||
*/
|
||||
TEST_CASE("Unit_deviceAllocation_Malloc_MultProcess") {
|
||||
auto res = testDeviceAllocMulProc(true);
|
||||
REQUIRE(res == true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Multiprocess device side new test.
|
||||
*/
|
||||
TEST_CASE("Unit_deviceAllocation_New_MultProcess") {
|
||||
auto res = testDeviceAllocMulProc(false);
|
||||
REQUIRE(res == true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Multiprocess device side malloc, write and free test.
|
||||
*/
|
||||
TEST_CASE("Unit_deviceAllocation_MallocFree_MultProcess") {
|
||||
auto res = testDeviceMemMulProc(true);
|
||||
REQUIRE(res == true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Multiprocess device side new, write and delete test.
|
||||
*/
|
||||
TEST_CASE("Unit_deviceAllocation_NewDelete_MultProcess") {
|
||||
auto res = testDeviceMemMulProc(false);
|
||||
REQUIRE(res == true);
|
||||
}
|
||||
@@ -6,3 +6,4 @@ if(HIP_PLATFORM MATCHES "amd")
|
||||
add_subdirectory(printf)
|
||||
add_subdirectory(stream)
|
||||
endif()
|
||||
add_subdirectory(deviceallocation)
|
||||
|
||||
@@ -0,0 +1,8 @@
|
||||
# Common Tests - Test independent of all platforms
|
||||
set(TEST_SRC
|
||||
Stress_deviceAllocationStress.cc
|
||||
)
|
||||
|
||||
hip_add_exe_to_target(NAME devalloc
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME stress_test)
|
||||
@@ -0,0 +1,487 @@
|
||||
/*
|
||||
Copyright (c) 2022 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_test_common.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <unistd.h>
|
||||
// Size Macros
|
||||
#define MEMORY_CHUNK_SIZE (1024*1024)
|
||||
#define MEMORY_CHUNK_SIZE_ODD (1025*1025)
|
||||
#define MAXIMUM_CHUNKS (256*1024)
|
||||
// Subtest Macros
|
||||
#define NO_ALLOCATION_ONHOST 0
|
||||
#define ALLOCATE_ONHOST_HIPMALLOCMANAGED 1
|
||||
#define ALLOCATE_ONHOST_HIPMALLOC 2
|
||||
// Test Type Macros
|
||||
#define TEST_MALLOC_FREE 1
|
||||
#define TEST_NEW_DELETE 2
|
||||
// GPU threads
|
||||
#define BLOCKSIZE 512
|
||||
#define GRIDSIZE 512
|
||||
// Test parameters
|
||||
// Two different loops
|
||||
#define NUM_OF_LOOP_SINGLE_KER 100000
|
||||
#define NUM_OF_LOOP_MULTIPLE_KER 20000
|
||||
|
||||
// The following flag is defined for platforms (nvidia)
|
||||
// which honors device memory limit. For AMD this flag
|
||||
// is disabled and defect is raised.
|
||||
#if HT_NVIDIA
|
||||
#define HT_HONORS_DEVICEMEMORY_LIMIT
|
||||
#endif
|
||||
|
||||
#ifdef HT_HONORS_DEVICEMEMORY_LIMIT
|
||||
__device__ static char* dev_mem_glob[MAXIMUM_CHUNKS];
|
||||
#endif
|
||||
__device__ static int* dev_mem[GRIDSIZE];
|
||||
__device__ static int* dev_common_ptr;
|
||||
|
||||
#ifdef HT_HONORS_DEVICEMEMORY_LIMIT
|
||||
/**
|
||||
* This kernel checks kernel allocation of size more than available
|
||||
* memory.
|
||||
*/
|
||||
static __global__ void kerTestDynamicAllocNeg(int test_type,
|
||||
size_t perThreadSize,
|
||||
int *ret) {
|
||||
// Allocate
|
||||
char* ptr = nullptr;
|
||||
printf("Memory to allocate in GPU = %zu \n", perThreadSize);
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
ptr = reinterpret_cast<char*> (malloc(perThreadSize));
|
||||
} else {
|
||||
ptr = new char[perThreadSize];
|
||||
}
|
||||
printf("Allocation Done \n");
|
||||
if (ptr == nullptr) {
|
||||
printf("Allocation Failed. PASSED! \n");
|
||||
*ret = 0;
|
||||
return;
|
||||
} else {
|
||||
// Free memory
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
free(ptr);
|
||||
} else {
|
||||
delete[] ptr;
|
||||
}
|
||||
*ret = -1;
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* This kernel allocates memory till nullptr is returned.
|
||||
*/
|
||||
static __global__ void kerAllocTillExhaust(int test_type,
|
||||
size_t *total_allocated_mem,
|
||||
size_t mem_chunk_size) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Allocate memory in thread 0 of block 0
|
||||
if (0 == myId) {
|
||||
for (int idx = 0; idx < MAXIMUM_CHUNKS; idx++) {
|
||||
dev_mem_glob[idx] = nullptr;
|
||||
}
|
||||
int idx = 0;
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
do {
|
||||
dev_mem_glob[idx] =
|
||||
reinterpret_cast<char*> (malloc(mem_chunk_size));
|
||||
if (idx >= MAXIMUM_CHUNKS) {
|
||||
break;
|
||||
}
|
||||
} while (dev_mem_glob[idx++] != nullptr);
|
||||
} else {
|
||||
do {
|
||||
dev_mem_glob[idx] =
|
||||
reinterpret_cast<char*> (new char[mem_chunk_size]);
|
||||
if (idx >= MAXIMUM_CHUNKS) {
|
||||
break;
|
||||
}
|
||||
} while (dev_mem_glob[idx++] != nullptr);
|
||||
}
|
||||
idx = 0;
|
||||
*total_allocated_mem = 0;
|
||||
while ((dev_mem_glob[idx] != nullptr) &&
|
||||
(idx < MAXIMUM_CHUNKS)) {
|
||||
*total_allocated_mem = *total_allocated_mem + mem_chunk_size;
|
||||
idx++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* This kernel deletes the memory.
|
||||
*/
|
||||
static __global__ void kerFreeAll(int test_type) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if (0 == myId) {
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
int idx = 0;
|
||||
while (dev_mem_glob[idx] != nullptr) {
|
||||
free(dev_mem_glob[idx++]);
|
||||
if (idx >= MAXIMUM_CHUNKS) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
int idx = 0;
|
||||
while (dev_mem_glob[idx] != nullptr) {
|
||||
delete[] (dev_mem_glob[idx++]);
|
||||
if (idx >= MAXIMUM_CHUNKS) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
/**
|
||||
* This kernel allocates memory once in thread 0 of each block and
|
||||
* access this memory in all threads of the block. The memory is
|
||||
* finally deleted in last thread of each block.
|
||||
*/
|
||||
static __global__ void kerBlockLevelMemoryAllocation(int *outputBuf,
|
||||
int test_type) {
|
||||
int myThreadId = threadIdx.x, lastThreadId = (blockDim.x - 1);
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Allocate memory in thread 0
|
||||
if (0 == myThreadId) {
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
dev_mem[blockIdx.x] =
|
||||
reinterpret_cast<int*> (malloc(blockDim.x*sizeof(int)));
|
||||
} else {
|
||||
dev_mem[blockIdx.x] =
|
||||
reinterpret_cast<int*> (new int[blockDim.x]);
|
||||
}
|
||||
}
|
||||
// All threads wait at this barrier
|
||||
__syncthreads();
|
||||
// Check allocated memory in all threads in block before access
|
||||
if (dev_mem[blockIdx.x] == nullptr) {
|
||||
printf("Device Allocation Failed in thread = %d \n", myId);
|
||||
return;
|
||||
}
|
||||
int *ptr = reinterpret_cast<int*> (dev_mem[blockIdx.x]);
|
||||
// Copy to buffer
|
||||
ptr[myThreadId] = myId;
|
||||
// All threads wait
|
||||
__syncthreads();
|
||||
// Copy memory to host and free the memory in thread <blockDim.x - 1>
|
||||
if (lastThreadId == myThreadId) {
|
||||
for (size_t idx = 0; idx < blockDim.x; idx++) {
|
||||
outputBuf[idx + blockDim.x * blockIdx.x] = ptr[idx];
|
||||
}
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
free(ptr);
|
||||
} else {
|
||||
delete[] ptr;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* This kernel allocates memory in one thread.
|
||||
*/
|
||||
static __global__ void kerAlloc(int test_type) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Allocate memory in thread 0 of block 0
|
||||
if (0 == myId) {
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
dev_common_ptr =
|
||||
reinterpret_cast<int*> (malloc(blockDim.x*gridDim.x*sizeof(int)));
|
||||
} else {
|
||||
dev_common_ptr =
|
||||
reinterpret_cast<int*> (new int[blockDim.x*gridDim.x]);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* This kernel writes to memory allocated in <kerAlloc>.
|
||||
*/
|
||||
static __global__ void kerWrite() {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Check allocated memory in all threads in block before access
|
||||
if (dev_common_ptr == nullptr) {
|
||||
printf("Device Allocation Failed in thread = %d \n", myId);
|
||||
return;
|
||||
}
|
||||
// Copy to buffer
|
||||
dev_common_ptr[myId] = myId;
|
||||
}
|
||||
|
||||
/**
|
||||
* This kernel copies the contents of memory allocated in <kerAlloc>
|
||||
* to host and deletes the memory from thread 0.
|
||||
*/
|
||||
static __global__ void kerFree(int *outputBuf, int test_type) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Check allocated memory in all threads in block before access
|
||||
if (dev_common_ptr == nullptr) {
|
||||
printf("Device Allocation Failed in thread = %d \n", myId);
|
||||
return;
|
||||
}
|
||||
if (0 == myId) {
|
||||
for (size_t idx = 0; idx < (blockDim.x*gridDim.x); idx++) {
|
||||
outputBuf[idx] = dev_common_ptr[idx];
|
||||
}
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
free(dev_common_ptr);
|
||||
} else {
|
||||
delete[] dev_common_ptr;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#ifdef HT_HONORS_DEVICEMEMORY_LIMIT
|
||||
/**
|
||||
* Local function: Launch kerAllocTillExhaust<<<>>> and
|
||||
* kerFreeAll<<<>>> to test memory allocation till all device
|
||||
* memory is exhausted.
|
||||
*/
|
||||
static bool TestAllocationOfAllAvailableMemory(int test_type,
|
||||
int category, size_t mem_chunk_size) {
|
||||
size_t avail1 = 0, avail2 = 0, tot = 0;
|
||||
constexpr size_t host_alloc = 2147483648; // 2 GB
|
||||
HIP_CHECK(hipMemGetInfo(&avail1, &tot));
|
||||
#if HT_NVIDIA
|
||||
HIP_CHECK(hipDeviceSetLimit(hipLimitMallocHeapSize, avail1));
|
||||
#endif
|
||||
size_t *tot_alloc_mem_d = nullptr, *tot_alloc_mem_h = nullptr;
|
||||
tot_alloc_mem_h =
|
||||
reinterpret_cast<size_t*> (malloc(sizeof(size_t)));
|
||||
REQUIRE(nullptr != tot_alloc_mem_h);
|
||||
HIP_CHECK(hipMalloc(&tot_alloc_mem_d, sizeof(size_t)));
|
||||
REQUIRE(nullptr != tot_alloc_mem_d);
|
||||
char *devptrHost = nullptr;
|
||||
if (category == ALLOCATE_ONHOST_HIPMALLOCMANAGED) {
|
||||
HIP_CHECK(hipMallocManaged(&devptrHost, host_alloc));
|
||||
} else if (category == ALLOCATE_ONHOST_HIPMALLOC) {
|
||||
HIP_CHECK(hipMalloc(&devptrHost, host_alloc));
|
||||
}
|
||||
HIP_CHECK(hipMemGetInfo(&avail1, &tot));
|
||||
INFO("Total available memory " << tot);
|
||||
INFO("Available memory before allocation " << avail1);
|
||||
// Launch Test Kernel
|
||||
kerAllocTillExhaust<<<1, 1>>>(test_type, tot_alloc_mem_d,
|
||||
mem_chunk_size);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
// Copy to host buffer
|
||||
HIP_CHECK(hipMemcpy(tot_alloc_mem_h, tot_alloc_mem_d,
|
||||
sizeof(size_t), hipMemcpyDefault));
|
||||
HIP_CHECK(hipMemGetInfo(&avail2, &tot));
|
||||
kerFreeAll<<<1, 1>>>(test_type);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
// Copy to host buffer
|
||||
bool bPassed = false;
|
||||
INFO("Available memory after allocation " << avail2);
|
||||
if (category == NO_ALLOCATION_ONHOST) {
|
||||
size_t allocated_dev_mem = (tot - avail2);
|
||||
if (allocated_dev_mem >= *tot_alloc_mem_h) {
|
||||
bPassed = true;
|
||||
}
|
||||
} else if ((category == ALLOCATE_ONHOST_HIPMALLOCMANAGED) ||
|
||||
(category == ALLOCATE_ONHOST_HIPMALLOC)) {
|
||||
size_t allocated_dev_mem = (tot - avail2 - host_alloc);
|
||||
if (allocated_dev_mem >= *tot_alloc_mem_h) {
|
||||
bPassed = true;
|
||||
}
|
||||
hipFree(devptrHost);
|
||||
}
|
||||
hipFree(tot_alloc_mem_d);
|
||||
free(tot_alloc_mem_h);
|
||||
return bPassed;
|
||||
}
|
||||
#endif
|
||||
/**
|
||||
* Local function: Launch kerBlockLevelMemoryAllocation<<<>>>
|
||||
* in a loop to stress test allocation and deallocation.
|
||||
*/
|
||||
static bool TestMemoryAllocationInLoop(int test_type,
|
||||
bool isMultikernel = false) {
|
||||
int *outputVec_d{nullptr}, *outputVec_h{nullptr};
|
||||
int arraysize = (BLOCKSIZE * GRIDSIZE);
|
||||
outputVec_h = reinterpret_cast<int*> (malloc(sizeof(int) * arraysize));
|
||||
REQUIRE(outputVec_h != nullptr);
|
||||
HIP_CHECK(hipMalloc(&outputVec_d, (sizeof(int) * arraysize)));
|
||||
bool bPassed = true;
|
||||
// Launch Test Kernel
|
||||
int max_index = 0;
|
||||
if (isMultikernel) {
|
||||
max_index = NUM_OF_LOOP_MULTIPLE_KER;
|
||||
} else {
|
||||
max_index = NUM_OF_LOOP_SINGLE_KER;
|
||||
}
|
||||
for (int idx = 0; idx < max_index; idx++) {
|
||||
if (isMultikernel) {
|
||||
kerAlloc<<<GRIDSIZE, BLOCKSIZE>>>(test_type);
|
||||
kerWrite<<<GRIDSIZE, BLOCKSIZE>>>();
|
||||
kerFree<<<GRIDSIZE, BLOCKSIZE>>>(outputVec_d, test_type);
|
||||
} else {
|
||||
kerBlockLevelMemoryAllocation<<<GRIDSIZE, BLOCKSIZE>>>(outputVec_d,
|
||||
test_type);
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
// Copy to host buffer
|
||||
HIP_CHECK(hipMemcpy(outputVec_h, outputVec_d, sizeof(int) * arraysize,
|
||||
hipMemcpyDefault));
|
||||
bPassed = true;
|
||||
for (int idx = 0; idx < arraysize; idx++) {
|
||||
if (outputVec_h[idx] != idx) {
|
||||
bPassed = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!bPassed) break;
|
||||
}
|
||||
hipFree(outputVec_d);
|
||||
free(outputVec_h);
|
||||
return bPassed;
|
||||
}
|
||||
|
||||
#ifdef HT_HONORS_DEVICEMEMORY_LIMIT
|
||||
/**
|
||||
* Scenario: Test malloc till nullptr is returned using even chunksize.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_malloc_Even") {
|
||||
REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_MALLOC_FREE,
|
||||
NO_ALLOCATION_ONHOST, MEMORY_CHUNK_SIZE));
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario: Test malloc till nullptr is returned using odd chunksize.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_malloc_Odd") {
|
||||
REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_MALLOC_FREE,
|
||||
NO_ALLOCATION_ONHOST, MEMORY_CHUNK_SIZE_ODD));
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario: Test new till nullptr is returned using even chunksize.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_new_Even") {
|
||||
REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_NEW_DELETE,
|
||||
NO_ALLOCATION_ONHOST, MEMORY_CHUNK_SIZE));
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario: Test new till nullptr is returned using odd chunksize.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_new_Odd") {
|
||||
REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_NEW_DELETE,
|
||||
NO_ALLOCATION_ONHOST, MEMORY_CHUNK_SIZE_ODD));
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario: This test checks device allocation using malloc till nullptr
|
||||
* is returned. Device memory is also allocated using hipmallocmanaged
|
||||
* from host.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_malloc_hipmallocmanaged") {
|
||||
REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_MALLOC_FREE,
|
||||
ALLOCATE_ONHOST_HIPMALLOCMANAGED, MEMORY_CHUNK_SIZE));
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario: This test checks device allocation using new till nullptr
|
||||
* is returned. Device memory is also allocated using hipmallocmanaged
|
||||
* from host.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_new_hipmallocmanaged") {
|
||||
REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_NEW_DELETE,
|
||||
ALLOCATE_ONHOST_HIPMALLOCMANAGED, MEMORY_CHUNK_SIZE));
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario: This test checks device allocation using malloc till nullptr
|
||||
* is returned. Device memory is also allocated using hipmalloc from host.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_malloc_hipmalloc") {
|
||||
REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_MALLOC_FREE,
|
||||
ALLOCATE_ONHOST_HIPMALLOC, MEMORY_CHUNK_SIZE));
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario: This test checks device allocation using new till nullptr
|
||||
* is returned. Device memory is also allocated using hipmalloc from host.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_new_hipmalloc") {
|
||||
REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_NEW_DELETE,
|
||||
ALLOCATE_ONHOST_HIPMALLOC, MEMORY_CHUNK_SIZE));
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario: This test validates device allocation negative scenario
|
||||
* when size > available memory.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_Negative") {
|
||||
int *ret_d{nullptr}, *ret_h{nullptr};
|
||||
size_t avail = 0, tot = 0;
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
printf("Available Memory in GPU = %zu \n", avail);
|
||||
ret_h = reinterpret_cast<int*> (malloc(sizeof(int)));
|
||||
REQUIRE(ret_h != nullptr);
|
||||
HIP_CHECK(hipMalloc(&ret_d, (sizeof(int))));
|
||||
SECTION("Test allocation with malloc") {
|
||||
kerTestDynamicAllocNeg<<<1, 1>>>(TEST_MALLOC_FREE, (avail + 1), ret_d);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipMemcpy(ret_h, ret_d, sizeof(int), hipMemcpyDefault));
|
||||
REQUIRE(0 == *ret_h);
|
||||
}
|
||||
|
||||
SECTION("Test allocation with new") {
|
||||
kerTestDynamicAllocNeg<<<1, 1>>>(TEST_NEW_DELETE, (avail + 1), ret_d);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipMemcpy(ret_h, ret_d, sizeof(int), hipMemcpyDefault));
|
||||
REQUIRE(0 == *ret_h);
|
||||
}
|
||||
hipFree(ret_d);
|
||||
free(ret_h);
|
||||
}
|
||||
#endif
|
||||
/**
|
||||
* Scenario: This test performs stress test of malloc/free in a loop
|
||||
* using single kernel.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_malloc_loop_singlekernel") {
|
||||
REQUIRE(true == TestMemoryAllocationInLoop(TEST_MALLOC_FREE, false));
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario: This test performs stress test of new/delete in a loop
|
||||
* using single kernel.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_new_loop_singlekernel") {
|
||||
REQUIRE(true == TestMemoryAllocationInLoop(TEST_NEW_DELETE, false));
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario: This test performs stress test of malloc/free in a loop
|
||||
* using multiple kernel.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_malloc_loop_multkernel") {
|
||||
REQUIRE(true == TestMemoryAllocationInLoop(TEST_MALLOC_FREE, true));
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario: This test performs stress test of new/delete in a loop
|
||||
* using multiple kernel.
|
||||
*/
|
||||
TEST_CASE("Stress_deviceAllocation_new_loop_multkernel") {
|
||||
REQUIRE(true == TestMemoryAllocationInLoop(TEST_NEW_DELETE, true));
|
||||
}
|
||||
@@ -15,6 +15,11 @@ set(TEST_SRC
|
||||
syncthreadsor.cc
|
||||
)
|
||||
|
||||
if(UNIX)
|
||||
set(TEST_SRC ${TEST_SRC}
|
||||
deviceAllocation.cc)
|
||||
endif()
|
||||
|
||||
# AMD only tests
|
||||
set(AMD_TEST_SRC
|
||||
unsafeAtomicAdd.cc
|
||||
@@ -43,6 +48,14 @@ set(AMD_ARCH_SPEC_TEST_SRC
|
||||
unsafeAtomicAdd_NonCoherent_withunsafeflag.cc
|
||||
)
|
||||
|
||||
|
||||
if(UNIX)
|
||||
add_custom_target(kerDevAllocMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
|
||||
add_custom_target(kerDevWriteMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevWriteMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevWriteMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
|
||||
add_custom_target(kerDevFreeMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevFreeMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevFreeMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
|
||||
add_custom_target(kerDevAllocSingleKer.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocSingleKer.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocSingleKer.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
|
||||
endif()
|
||||
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
if (DEFINED OFFLOAD_ARCH_STR)
|
||||
string(FIND ${OFFLOAD_ARCH_STR} "gfx90a" ARCH_CHECK)
|
||||
@@ -78,3 +91,7 @@ elseif(HIP_PLATFORM MATCHES "nvidia")
|
||||
TEST_TARGET_NAME build_tests
|
||||
COMPILE_OPTIONS --Wno-deprecated-declarations)
|
||||
endif()
|
||||
|
||||
if(UNIX)
|
||||
add_dependencies(build_tests kerDevAllocMultCO.code kerDevWriteMultCO.code kerDevFreeMultCO.code kerDevAllocSingleKer.code)
|
||||
endif()
|
||||
|
||||
@@ -0,0 +1,36 @@
|
||||
/*
|
||||
Copyright (c) 2022 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.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#define INTERNAL_BUFFER_SIZE 8
|
||||
// Test Type
|
||||
#define TEST_MALLOC_FREE 1
|
||||
#define TEST_NEW_DELETE 2
|
||||
// Kernel Params
|
||||
#define BLOCKSIZE 64
|
||||
#define GRIDSIZE 32
|
||||
// Code Obj
|
||||
#define DEV_ALLOC_SINGKER_COBJ "kerDevAllocSingleKer.code"
|
||||
#define DEV_ALLOC_SINGKER_COBJ_FUNC "ker_TestDynamicAllocInAllThreads_CodeObj"
|
||||
#define DEV_ALLOC_MULCOBJ "kerDevAllocMultCO.code"
|
||||
#define DEV_WRITE_MULCOBJ "kerDevWriteMultCO.code"
|
||||
#define DEV_FREE_MULCOBJ "kerDevFreeMultCO.code"
|
||||
#define DEV_ALLOC_MULCODEOBJ_ALLOC "ker_Alloc_MultCodeObj"
|
||||
#define DEV_ALLOC_MULCODEOBJ_WRITE "ker_Write_MultCodeObj"
|
||||
#define DEV_ALLOC_MULCODEOBJ_FREE "ker_Free_MultCodeObj"
|
||||
@@ -0,0 +1,132 @@
|
||||
/*
|
||||
Copyright (c) 2022 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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <cfloat>
|
||||
#include <atomic>
|
||||
#include "./defs.h"
|
||||
|
||||
static __device__ int* deviceAlloc(int test_type);
|
||||
static __device__ void deviceWrite(int myId, int *devmem);
|
||||
static __device__ void deviceFree(int *outputBuf, int *devmem,
|
||||
int test_type, int myId);
|
||||
|
||||
/**
|
||||
* Allocation base and derived class to test dynamic allocation.
|
||||
*/
|
||||
class baseAlloc{
|
||||
public:
|
||||
virtual __device__ int* alloc(size_t size) = 0;
|
||||
virtual __device__ void free(int* ptr) = 0;
|
||||
};
|
||||
|
||||
class derivedAlloc: public baseAlloc{
|
||||
public:
|
||||
virtual __device__ int* alloc(size_t size) {
|
||||
return new int[size];
|
||||
}
|
||||
virtual __device__ void free(int* ptr) {
|
||||
delete ptr;
|
||||
}
|
||||
};
|
||||
|
||||
/**
|
||||
* Allocation Structure to test dynamic allocation.
|
||||
*/
|
||||
struct deviceAllocFunc{
|
||||
int* (*alloc)(int);
|
||||
void (*write)(int, int*);
|
||||
void (*free)(int*, int*, int, int);
|
||||
};
|
||||
|
||||
/**
|
||||
* Simple Structure to test dynamic allocation.
|
||||
*/
|
||||
struct simpleStruct{
|
||||
int32_t i;
|
||||
double d;
|
||||
float f;
|
||||
int16_t s;
|
||||
char c;
|
||||
int32_t iarr[INTERNAL_BUFFER_SIZE];
|
||||
bool operator!=(const struct simpleStruct &inpStr) {
|
||||
if ((i != inpStr.i) || (d != inpStr.d) ||
|
||||
(f != inpStr.f) || (s != inpStr.s) || (c != inpStr.c)) {
|
||||
return true;
|
||||
}
|
||||
for (int32_t idx = 0; idx < INTERNAL_BUFFER_SIZE; idx++) {
|
||||
if (iarr[idx] != inpStr.iarr[idx]) {
|
||||
return true;
|
||||
}
|
||||
}
|
||||
return false;
|
||||
}
|
||||
};
|
||||
|
||||
/**
|
||||
* Simple Structure containing thread information
|
||||
*/
|
||||
struct threadInfo{
|
||||
int threadid;
|
||||
int blockid;
|
||||
int32_t ival;
|
||||
double dval;
|
||||
float fval;
|
||||
int16_t sval;
|
||||
char cval;
|
||||
};
|
||||
|
||||
/**
|
||||
* C/C++ Union
|
||||
*/
|
||||
union testInfoUnion{
|
||||
int32_t ival;
|
||||
double dval;
|
||||
float fval;
|
||||
int16_t sval;
|
||||
char cval;
|
||||
};
|
||||
|
||||
/**
|
||||
* Complex (nested) Structure to test dynamic allocation using malloc.
|
||||
*/
|
||||
struct complexStructure{
|
||||
struct threadInfo *sthreadInfo;
|
||||
__device__ void alloc_internal_members(int test_type, size_t size) {
|
||||
sthreadInfo = nullptr;
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
sthreadInfo = reinterpret_cast<struct threadInfo*>(
|
||||
malloc(size*sizeof(struct threadInfo)));
|
||||
} else {
|
||||
sthreadInfo = new struct threadInfo[size];
|
||||
}
|
||||
}
|
||||
|
||||
__device__ void free_internal_members(int test_type) {
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
free(sthreadInfo);
|
||||
} else {
|
||||
delete[] sthreadInfo;
|
||||
}
|
||||
sthreadInfo = nullptr;
|
||||
}
|
||||
};
|
||||
Разница между файлами не показана из-за своего большого размера
Загрузить разницу
@@ -0,0 +1,39 @@
|
||||
/*
|
||||
Copyright (c) 2022 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"
|
||||
#include "./defs.h"
|
||||
|
||||
/**
|
||||
* This kernel allocates memory in thread 0.
|
||||
*/
|
||||
extern "C" __global__ void ker_Alloc_MultCodeObj(int **dev_mem,
|
||||
int test_type) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Allocate memory in thread 0 of block 0
|
||||
if (0 == myId) {
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
*dev_mem =
|
||||
reinterpret_cast<int*> (malloc(blockDim.x*gridDim.x*sizeof(int)));
|
||||
} else {
|
||||
*dev_mem =
|
||||
reinterpret_cast<int*> (new int[blockDim.x*gridDim.x]);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,57 @@
|
||||
/*
|
||||
Copyright (c) 2022 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"
|
||||
#include "./defs.h"
|
||||
/**
|
||||
* This kernel allocates and deallocates memory in every thread.
|
||||
*/
|
||||
extern "C" __global__ void ker_TestDynamicAllocInAllThreads_CodeObj(
|
||||
int *outputBuf, int test_type, int value,
|
||||
size_t perThreadSize) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Allocate
|
||||
size_t size = 0;
|
||||
int* ptr = nullptr;
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
size = perThreadSize * sizeof(int);
|
||||
ptr = reinterpret_cast<int*> (malloc(size));
|
||||
} else {
|
||||
size = perThreadSize;
|
||||
ptr = new int[perThreadSize];
|
||||
}
|
||||
if (ptr == nullptr) {
|
||||
printf("Device Allocation in thread %d Failed! \n", myId);
|
||||
return;
|
||||
}
|
||||
// Set memory
|
||||
for (size_t idx = 0; idx < perThreadSize; idx++) {
|
||||
ptr[idx] = value;
|
||||
}
|
||||
// Copy to output buffer
|
||||
for (size_t idx = 0; idx < perThreadSize; idx++) {
|
||||
outputBuf[myId*perThreadSize + idx] = ptr[idx];
|
||||
}
|
||||
// Free memory
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
free(ptr);
|
||||
} else {
|
||||
delete[] ptr;
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,47 @@
|
||||
/*
|
||||
Copyright (c) 2022 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"
|
||||
#include "./defs.h"
|
||||
|
||||
/**
|
||||
* This kernel copies the contents of memory allocated in
|
||||
* ker_Alloc_MultCodeObj<<<>>> to host and deletes the memory
|
||||
* from thread 0.
|
||||
*/
|
||||
extern "C" __global__ void ker_Free_MultCodeObj(int *outputBuf,
|
||||
int **dev_mem, int test_type) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Check allocated memory in all threads in block before access
|
||||
if (*dev_mem == nullptr) {
|
||||
printf("Device Allocation Failed in thread = %d \n", myId);
|
||||
return;
|
||||
}
|
||||
|
||||
if (0 == myId) {
|
||||
for (size_t idx = 0; idx < (blockDim.x*gridDim.x); idx++) {
|
||||
outputBuf[idx] = (*dev_mem)[idx];
|
||||
}
|
||||
if (test_type == TEST_MALLOC_FREE) {
|
||||
free(*dev_mem);
|
||||
} else {
|
||||
delete[] (*dev_mem);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,36 @@
|
||||
/*
|
||||
Copyright (c) 2022 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"
|
||||
#include "./defs.h"
|
||||
|
||||
/**
|
||||
* This kernel writes to memory allocated in ker_Alloc_MultCodeObj<<<>>>.
|
||||
*/
|
||||
extern "C" __global__ void ker_Write_MultCodeObj(int **dev_mem,
|
||||
int value) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
// Check allocated memory in all threads in block before access
|
||||
if (*dev_mem == nullptr) {
|
||||
printf("Device Allocation Failed in thread = %d \n", myId);
|
||||
return;
|
||||
}
|
||||
// Copy to buffer
|
||||
(*dev_mem)[myId] = value;
|
||||
}
|
||||
@@ -96,6 +96,7 @@ set(TEST_SRC
|
||||
hipMemcpySync.cc
|
||||
hipMemsetSync.cc
|
||||
hipMemsetAsync.cc
|
||||
hipMemAdvise.cc
|
||||
)
|
||||
else()
|
||||
set(TEST_SRC
|
||||
@@ -167,6 +168,7 @@ set(TEST_SRC
|
||||
hipMemcpySync.cc
|
||||
hipMemsetSync.cc
|
||||
hipMemsetAsync.cc
|
||||
hipMemAdvise.cc
|
||||
)
|
||||
endif()
|
||||
|
||||
|
||||
@@ -0,0 +1,962 @@
|
||||
/*
|
||||
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, INNCLUDING 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 ANNY 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.
|
||||
*/
|
||||
|
||||
/* Test Case Description:
|
||||
Scenario-1: The following Function Tests the working of flags which can be
|
||||
assigned to HMM memory using hipMemAdvise() api
|
||||
Scenario-2: Negative tests on hipMemAdvise() api
|
||||
Scenario-3: The following function tests various scenarios around the flag
|
||||
'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api
|
||||
Scenario-4: The following function tests various scenarios around the flag
|
||||
'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api
|
||||
Scenario-5: The following function verifies if assigning of a flag
|
||||
invalidates the earlier flag which was assigned to the same memory region
|
||||
using hipMemAdvise()
|
||||
Scenario-6: The following function tests if peers can set
|
||||
hipMemAdviseSetAccessedBy flag
|
||||
on HMM memory prefetched on each of the other gpus
|
||||
Scenario-7: Set AccessedBy flag and check value returned by
|
||||
hipMemRangeGetAttribute() It should be -2(same is observed on cuda)
|
||||
Scenario-8: Set AccessedBy flag to device 0 on Hmm memory and prefetch the
|
||||
memory to device 1, then probe for AccessedBy flag using
|
||||
hipMemRangeGetAttribute() we should still see the said flag is set for
|
||||
device 0
|
||||
Scenario-9: 1) Set AccessedBy to device 0 followed by PreferredLocation to
|
||||
device 1 check for AccessedBy flag using hipMemRangeGetAttribute() it should
|
||||
return 0
|
||||
2) Unset AccessedBy to 0 and set it to device 1 followed by
|
||||
PreferredLocation to device 1, check for AccessedBy flag using
|
||||
hipMemRangeGetAttribute() it should return 1
|
||||
Scenario-10: Set AccessedBy flag to HMM memory launch a kernel and then unset
|
||||
AccessedBy, launch kernel. We should not have any access issues
|
||||
Scenario-11: Allocate memory using aligned_alloc(), assign PreferredLocation
|
||||
flag to the allocated memory and launch a kernel. Kernel should get executed
|
||||
successfully without hang or segfault
|
||||
Scenario-12: Allocate Hmm memory, set advise to PreferredLocation and then
|
||||
get attribute using the api hipMemRangeGetAttribute() for
|
||||
hipMemRangeAttributeLastPrefetchLocation the value returned should be -2
|
||||
Scenario-13: Allocate HMM memory, set PreferredLocation to device 0, Prfetch
|
||||
the mem to device1, probe for hipMemRangeAttributeLastPrefetchLocation using
|
||||
hipMemRangeGetAttribute(), we should get 1
|
||||
Scenario-14: Allocate HMM memory, set ReadMostly followed by
|
||||
PreferredLocation, probe for hipMemRangeAttributeReadMostly and
|
||||
hipMemRangeAttributePreferredLocation
|
||||
using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly.
|
||||
In other words setting of hipMemRangeAttributePreferredLocation should not
|
||||
impact hipMemRangeAttributeReadMostly advise to the memory
|
||||
Scenario-15: Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and
|
||||
launch kernel on all other gpus except 0. This test case may discover any
|
||||
effect or access denial case arising due to setting ReadMostly only to a
|
||||
particular gpu
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#if __linux__
|
||||
#include <unistd.h>
|
||||
#include <sys/mman.h>
|
||||
#include <sys/wait.h>
|
||||
#endif
|
||||
|
||||
// Kernel function
|
||||
__global__ void MemAdvseKernel(int n, int *x) {
|
||||
int index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (index < n)
|
||||
x[index] = x[index] * x[index];
|
||||
}
|
||||
|
||||
// Kernel
|
||||
__global__ void MemAdvise2(int *Hmm, int n) {
|
||||
int index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int stride = blockDim.x * gridDim.x;
|
||||
for (int i = index; i < n; i += stride) {
|
||||
Hmm[i] = Hmm[i] + 10;
|
||||
}
|
||||
}
|
||||
|
||||
// Kernel
|
||||
__global__ void MemAdvise3(int *Hmm, int *Hmm1, int n) {
|
||||
int index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int stride = blockDim.x * gridDim.x;
|
||||
for (int i = index; i < n; i += stride) {
|
||||
Hmm1[i] = Hmm[i] + 10;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
static bool CheckError(hipError_t err, int LineNo) {
|
||||
if (err == hipSuccess) {
|
||||
WARN("Error expected but received hipSuccess at line no.:" << LineNo);
|
||||
return false;
|
||||
} else {
|
||||
return true;
|
||||
}
|
||||
}
|
||||
|
||||
static int HmmAttrPrint() {
|
||||
int managed = 0;
|
||||
WARN("The following are the attribute values related to HMM for"
|
||||
" device 0:\n");
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
|
||||
WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeConcurrentManagedAccess, 0));
|
||||
WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccess, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" << managed);
|
||||
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
WARN("hipDeviceAttributeManagedMemory: " << managed);
|
||||
return managed;
|
||||
}
|
||||
|
||||
|
||||
// The following Function Tests the working of flags which can be assigned
|
||||
// to HMM memory using hipMemAdvise() api
|
||||
TEST_CASE("Unit_hipMemAdvise_TstFlags") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
bool IfTestPassed = true;
|
||||
int NumDevs = 0, *Outpt = nullptr;
|
||||
int MEM_SIZE = 4*1024, A_CONST = 9999;
|
||||
float *Hmm = nullptr;
|
||||
int AttrVal = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
Outpt = new int(NumDevs);
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal));
|
||||
// With the following for loop we iterate through each of the Gpus in the
|
||||
// system set and unset the flags and check the behavior.
|
||||
for (int i = 0; i < NumDevs; ++i) {
|
||||
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetReadMostly, i));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal),
|
||||
hipMemRangeAttributeReadMostly, Hmm,
|
||||
MEM_SIZE * 2));
|
||||
if (AttrVal != 1) {
|
||||
WARN("Attempt to set hipMemAdviseSetReadMostly flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetReadMostly,
|
||||
i));
|
||||
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal),
|
||||
hipMemRangeAttributeReadMostly, Hmm,
|
||||
(MEM_SIZE * 2)));
|
||||
if (AttrVal != 0) {
|
||||
WARN("Attempt to Unset hipMemAdviseSetReadMostly flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
AttrVal = A_CONST;
|
||||
// Currently hipMemAdviseSetPreferredLocation and
|
||||
// hipMemAdviseSetAccessedBy
|
||||
// flags are resulting in issues: SWDEV-267357
|
||||
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2,
|
||||
hipMemAdviseSetPreferredLocation, i));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal),
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
Hmm, (MEM_SIZE * 2)));
|
||||
if (AttrVal != i) {
|
||||
WARN("Attempt to set hipMemAdviseSetPreferredLocation flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
AttrVal = A_CONST;
|
||||
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2,
|
||||
hipMemAdviseUnsetPreferredLocation, i));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal),
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
Hmm, (MEM_SIZE * 2)));
|
||||
if (AttrVal == i) {
|
||||
WARN("Attempt to Unset hipMemAdviseUnsetPreferredLocation ");
|
||||
WARN("flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
for (int m = 0; m < NumDevs; ++m) {
|
||||
Outpt[m] = A_CONST;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetAccessedBy, i));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt),
|
||||
hipMemRangeAttributeAccessedBy, Hmm,
|
||||
(MEM_SIZE * 2)));
|
||||
if ((Outpt[0]) != i) {
|
||||
WARN("Attempt to set hipMemAdviseSetAccessedBy flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
for (int m = 0; m < NumDevs; ++m) {
|
||||
Outpt[m] = A_CONST;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetAccessedBy,
|
||||
i));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt),
|
||||
hipMemRangeAttributeAccessedBy, Hmm,
|
||||
(MEM_SIZE * 2)));
|
||||
if ((Outpt[0]) >= 0) {
|
||||
WARN("Attempt to Unset hipMemAdviseUnsetAccessedBy flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
delete [] Outpt;
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemAdvise_NegtveTsts") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
bool IfTestPassed = true;
|
||||
int NumDevs = 0, MEM_SIZE = 4*1024;
|
||||
float *Hmm = nullptr;
|
||||
std::string str;
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal));
|
||||
#if HT_AMD
|
||||
// Passing invalid value(99) device param
|
||||
IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2,
|
||||
hipMemAdviseSetReadMostly, 99), __LINE__);
|
||||
|
||||
// Passing invalid value(-12) device param
|
||||
IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2,
|
||||
hipMemAdviseSetReadMostly, -12), __LINE__);
|
||||
#endif
|
||||
// Passing NULL as first parameter instead of valid pointer to a memory
|
||||
IfTestPassed &= CheckError(hipMemAdvise(NULL, MEM_SIZE * 2,
|
||||
hipMemAdviseSetReadMostly, 0), __LINE__);
|
||||
|
||||
// Passing 0 for count(2nd param) parameter
|
||||
IfTestPassed &= CheckError(hipMemAdvise(Hmm, 0, hipMemAdviseSetReadMostly,
|
||||
0), __LINE__);
|
||||
|
||||
// Passing count much more than actually allocated value
|
||||
IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 6,
|
||||
hipMemAdviseSetReadMostly, 0), __LINE__);
|
||||
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
// The following function tests various scenarios around the flag
|
||||
// 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api
|
||||
TEST_CASE("Unit_hipMemAdvise_PrefrdLoc") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
// Check that when a page fault occurs for the memory region set to devPtr,
|
||||
// the data is migrated to the destn processor
|
||||
int MEM_SIZE = 4096, A_CONST = 9999;
|
||||
int *Hmm = nullptr, NumDevs = 0, dev = A_CONST;
|
||||
bool IfTestPassed = true;
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 3, hipMemAttachGlobal));
|
||||
for (int i = 0; i < ((MEM_SIZE * 3)/4); ++i) {
|
||||
Hmm[i] = 4;
|
||||
}
|
||||
for (int devId = 0; devId < NumDevs; ++devId) {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3,
|
||||
hipMemAdviseSetPreferredLocation, devId));
|
||||
int NumElms = ((MEM_SIZE * 3)/4);
|
||||
MemAdvseKernel<<<NumElms/32, 32>>>(NumElms, Hmm);
|
||||
int dev = A_CONST;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev),
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
Hmm, MEM_SIZE * 3));
|
||||
if (dev != devId) {
|
||||
WARN("Memory observed to be not available on expected location\n");
|
||||
WARN("line no: " << __LINE__);
|
||||
WARN("dev: " << dev);
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
|
||||
// Check that when preferred location is set for a memory region,
|
||||
// data can still be prefetched using hipMemPrefetchAsync
|
||||
hipStream_t strm;
|
||||
dev = A_CONST;
|
||||
for (int devId = 0; devId < NumDevs; ++devId) {
|
||||
HIP_CHECK(hipSetDevice(devId));
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3,
|
||||
hipMemAdviseSetPreferredLocation, devId));
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE * 3, devId, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev),
|
||||
hipMemRangeAttributeLastPrefetchLocation,
|
||||
Hmm, MEM_SIZE * 3));
|
||||
if (dev != devId) {
|
||||
WARN("Memory reported to be not available at the Prefetched ");
|
||||
WARN("location with device id: " << devId);
|
||||
WARN("line no: " << __LINE__);
|
||||
WARN("dev: " << dev);
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
// The following function tests various scenarios around the flag
|
||||
// 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api
|
||||
|
||||
TEST_CASE("Unit_hipMemAdvise_ReadMostly") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
bool IfTestPassed = true;
|
||||
int MEM_SIZE = 4096, A_CONST = 9999;
|
||||
float *Hmm = nullptr;
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE));
|
||||
for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) {
|
||||
Hmm[i] = A_CONST;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, 0));
|
||||
// Checking if the data can be read after setting hipMemAdviseSetReadMostly
|
||||
for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) {
|
||||
if (Hmm[i] != A_CONST) {
|
||||
WARN("Didn't find expected value in Hmm memory after setting");
|
||||
WARN(" hipMemAdviseSetReadMostly flag line no.: " << __LINE__);
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
|
||||
// Checking if the memory region can be modified
|
||||
for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) {
|
||||
Hmm[i] = A_CONST;
|
||||
}
|
||||
|
||||
for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) {
|
||||
if (Hmm[i] != A_CONST) {
|
||||
WARN("Didn't find expected value in Hmm memory after Modification\n");
|
||||
WARN("line no.: " << __LINE__);
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
|
||||
int out = A_CONST;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&out, 4, hipMemRangeAttributeReadMostly,
|
||||
Hmm, MEM_SIZE));
|
||||
if (out != 1) {
|
||||
WARN("out value: " << out);
|
||||
IfTestPassed = false;
|
||||
}
|
||||
// Checking the advise attribute after prefetch
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, 0, 0));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&out, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly, Hmm,
|
||||
MEM_SIZE));
|
||||
if (out != 1) {
|
||||
WARN("Attribute assigned to memory changed after calling ");
|
||||
WARN("hipMemPrefetchAsync(). line no.: " << __LINE__);
|
||||
WARN("out value: " << out);
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
// The following function verifies if assigning of a flag invalidates the
|
||||
// earlier flag which was assigned to the same memory region using
|
||||
// hipMemAdvise()
|
||||
TEST_CASE("Unit_hipMemAdvise_TstFlgOverrideEffect") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
bool IfTestPassed = true;
|
||||
int MEM_SIZE = 4*4096, A_CONST = 9999;
|
||||
float *Hmm = nullptr;
|
||||
int NumDevs = 0, dev = A_CONST;
|
||||
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal));
|
||||
for (int i = 0; i < NumDevs; ++i) {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, i));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly, Hmm,
|
||||
MEM_SIZE));
|
||||
if (dev != 1) {
|
||||
WARN("hipMemAdviseSetReadMostly flag did not take affect despite ");
|
||||
WARN("setting it using hipMemAdvise(). line no.: " << __LINE__);
|
||||
IfTestPassed = false;
|
||||
break;
|
||||
}
|
||||
dev = A_CONST;
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetPreferredLocation,
|
||||
i));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int),
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
Hmm, MEM_SIZE));
|
||||
if (dev != i) {
|
||||
WARN("hipMemAdviseSetPreferredLocation flag did not take affect ");
|
||||
WARN("despite setting it using hipMemAdvise()\n");
|
||||
WARN("line no.: " << __LINE__);
|
||||
IfTestPassed = false;
|
||||
break;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, i));
|
||||
dev = A_CONST;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int),
|
||||
hipMemRangeAttributeAccessedBy, Hmm,
|
||||
MEM_SIZE));
|
||||
if (dev != i) {
|
||||
WARN("hipMemAdviseSetAccessedBy flag did not take affect despite ");
|
||||
WARN("setting it using hipMemAdvise(). line no.: " << __LINE__);
|
||||
IfTestPassed = false;
|
||||
break;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetAccessedBy, i));
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
// The following function tests if peers can set hipMemAdviseSetAccessedBy flag
|
||||
// on HMM memory prefetched on each of the other gpus
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_hipMemAdvise_TstAccessedByPeer") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
bool IfTestPassed = true;
|
||||
int *Hmm = nullptr, MEM_SIZE = 4*4096, A_CONST = 9999;;
|
||||
int NumDevs = 0, CanAccessPeer = A_CONST, flag = 0;
|
||||
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
if (NumDevs < 2) {
|
||||
SUCCEED("Test TestSetAccessedByPeer() need atleast two Gpus to test"
|
||||
" the scenario. This system has GPUs less than 2");
|
||||
}
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal));
|
||||
for (int i = 0; i < NumDevs; ++i) {
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, i, 0));
|
||||
for (int j = 0; j < NumDevs; ++j) {
|
||||
if (i == j)
|
||||
continue;
|
||||
HIP_CHECK(hipSetDevice(j));
|
||||
HIP_CHECK(hipDeviceCanAccessPeer(&CanAccessPeer, j, i));
|
||||
if (CanAccessPeer) {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, j));
|
||||
for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) {
|
||||
Hmm[m] = 4;
|
||||
}
|
||||
HIP_CHECK(hipDeviceEnablePeerAccess(i, 0));
|
||||
MemAdvseKernel<<<(MEM_SIZE/sizeof(int)/32), 32>>>(
|
||||
(MEM_SIZE/sizeof(int)), Hmm);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
// Verifying the result
|
||||
for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) {
|
||||
if (Hmm[m] != 16) {
|
||||
flag = 1;
|
||||
}
|
||||
}
|
||||
if (flag) {
|
||||
WARN("Didnt get Expected results with device: " << j);
|
||||
WARN("line no.: " << __LINE__);
|
||||
IfTestPassed = false;
|
||||
flag = 0;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
/* Set AccessedBy flag and check value returned by hipMemRangeGetAttribute()
|
||||
It should be -2(same is observed on cuda)*/
|
||||
TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int *Hmm = NULL, data = 999;
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, 2*4096));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeLastPrefetchLocation,
|
||||
Hmm, 2*4096));
|
||||
if (data != -2) {
|
||||
WARN("Didnt get expected value!!\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
/* Set AccessedBy flag to device 0 on Hmm memory and prefetch the memory to
|
||||
device 1, then probe for AccessedBy flag using hipMemRangeGetAttribute()
|
||||
we should still see the said flag is set for device 0*/
|
||||
TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg2") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int *Hmm = NULL, data = 999, Ngpus = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&Ngpus));
|
||||
if (Ngpus >= 2) {
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, 2*4096));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0));
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, 2*4096, 1, strm));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeAccessedBy, Hmm, 2*4096));
|
||||
if (data != 0) {
|
||||
WARN("Didnt get expected behavior at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
/* 1) Set AccessedBy to device 0 followed by PreferredLocation to device 1
|
||||
check for AccessedBy flag using hipMemRangeGetAttribute() it should
|
||||
return 0
|
||||
2) Unset AccessedBy to 0 and set it to device 1 followed by
|
||||
PreferredLocation to device 1, check for AccessedBy flag using
|
||||
hipMemRangeGetAttribute() it should return 1*/
|
||||
|
||||
TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg3") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int *Hmm = NULL, data = 999, Ngpus = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&Ngpus));
|
||||
if (Ngpus >= 2) {
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, 2*4096));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 1));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeAccessedBy, Hmm, 2*4096));
|
||||
if (data != 0) {
|
||||
WARN("Didnt get expected behavior at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 1));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeAccessedBy, Hmm, 2*4096));
|
||||
if (data != 1) {
|
||||
WARN("Didnt get expected behavior at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/* Set AccessedBy flag to HMM memory launch a kernel and then unset
|
||||
AccessedBy, launch kernel. We should not have any access issues*/
|
||||
|
||||
TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg4") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64;
|
||||
int DataMismatch = 0;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int))));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)),
|
||||
hipMemAdviseSetAccessedBy, 0));
|
||||
// Initializing memory
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
Hmm[i] = InitVal;
|
||||
}
|
||||
dim3 dimBlock(blockSize, 1, 1);
|
||||
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
|
||||
// launching kernel from each one of the gpus
|
||||
MemAdvise2<<<dimGrid, dimBlock, 0, strm>>>(Hmm, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
|
||||
// verifying the final result
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
if (Hmm[i] != (InitVal + 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
|
||||
if (DataMismatch != 0) {
|
||||
WARN("DataMismatch is observed at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)),
|
||||
hipMemAdviseUnsetAccessedBy, 0));
|
||||
MemAdvise2<<<dimGrid, dimBlock, 0, strm>>>(Hmm, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// verifying the final result
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
if (Hmm[i] != (InitVal + (2*10))) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
|
||||
if (DataMismatch != 0) {
|
||||
WARN("DataMismatch is observed at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/* Allocate memory using aligned_alloc(), assign PreferredLocation flag to
|
||||
the allocated memory and launch a kernel. Kernel should get executed
|
||||
successfully without hang or segfault*/
|
||||
#if __linux__ && HT_AMD
|
||||
TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") {
|
||||
if ((setenv("HSA_XNACK", "1", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
// The following code block is used to check for gfx906/8 so as to skip if
|
||||
// any of the gpus available
|
||||
int fd1[2]; // Used to store two ends of first pipe
|
||||
pid_t p;
|
||||
if (pipe(fd1) == -1) {
|
||||
fprintf(stderr, "Pipe Failed");
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
/* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
|
||||
GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
|
||||
int GpuId[2] = {0, 0};
|
||||
p = fork();
|
||||
|
||||
if (p < 0) {
|
||||
fprintf(stderr, "fork Failed");
|
||||
REQUIRE(false);
|
||||
} else if (p > 0) { // parent process
|
||||
close(fd1[1]); // Close writing end of first pipe
|
||||
// Wait for child to send a string
|
||||
wait(NULL);
|
||||
// Read string from child and close reading end.
|
||||
read(fd1[0], GpuId, 2 * sizeof(int));
|
||||
close(fd1[0]);
|
||||
if ((GpuId[0] == 1) || (GpuId[0] == 1)) {
|
||||
WARN("This test is not applicable on MI60 & MI100."
|
||||
"Skipping the test!!");
|
||||
exit(0);
|
||||
}
|
||||
} else { // child process
|
||||
close(fd1[0]); // Close read end of first pipe
|
||||
hipDeviceProp_t prop;
|
||||
HIPCHECK(hipGetDeviceProperties(&prop, 0));
|
||||
char *p = NULL;
|
||||
p = strstr(prop.gcnArchName, "gfx906");
|
||||
if (p) {
|
||||
WARN("gfx906 gpu found on this system!!");
|
||||
GpuId[0] = 1;
|
||||
}
|
||||
p = strstr(prop.gcnArchName, "gfx908");
|
||||
if (p) {
|
||||
WARN("gfx908 gpu found on this system!!");
|
||||
GpuId[1] = 1;
|
||||
}
|
||||
// Write concatenated string and close writing end
|
||||
write(fd1[1], GpuId, 2 * sizeof(int));
|
||||
close(fd1[1]);
|
||||
exit(0);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) {
|
||||
// The below part should be inside fork
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123;
|
||||
// Mllc = reinterpret_cast<(int *)>(aligned_alloc(4096, MemSz));
|
||||
Mllc = reinterpret_cast<int*>(aligned_alloc(4096, 4096*4));
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
Mllc[i] = InitVal;
|
||||
}
|
||||
hipStream_t strm;
|
||||
int DataMismatch = 0;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
// The following hipMemAdvise() call is made to know if advise on
|
||||
// aligned_alloc() is causing any issue
|
||||
HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0));
|
||||
HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
MemAdvise2<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
if (Mllc[i] != (InitVal + 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
WARN("DataMismatch observed!!");
|
||||
exit(9); // 9 for failure
|
||||
} else {
|
||||
exit(10); // 10 for Pass result
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
} else {
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
/* Allocate Hmm memory, set advise to PreferredLocation and then get
|
||||
attribute using the api hipMemRangeGetAttribute() for
|
||||
hipMemRangeAttributeLastPrefetchLocation the value returned should be -2*/
|
||||
|
||||
TEST_CASE("Unit_hipMemAdvise_TstMemAdvisePrefrdLoc") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int *Hmm = NULL, data = 999;
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, 4096));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeLastPrefetchLocation,
|
||||
Hmm, 4096));
|
||||
if (data != -2) {
|
||||
WARN("Didnt receive expected value.");
|
||||
REQUIRE(false);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/* Allocate HMM memory, set PreferredLocation to device 0, Prfetch the mem
|
||||
to device1, probe for hipMemRangeAttributeLastPrefetchLocation using
|
||||
hipMemRangeGetAttribute(), we should get 1*/
|
||||
|
||||
TEST_CASE("Unit_hipMemAdvise_TstMemAdviseLstPreftchLoc") {
|
||||
int NumDevs = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
if (NumDevs >= 2) {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int *Hmm = NULL, data = 999;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipSetDevice(1));
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, 4096));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0));
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, 4096, 1, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeLastPrefetchLocation,
|
||||
Hmm, 4096));
|
||||
if (data != 1) {
|
||||
WARN("Didnt receive expected value!!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
} else {
|
||||
SUCCEED("This system has less than 2 gpus hence skipping the test.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/* Allocate HMM memory, set ReadMostly followed by PreferredLocation, probe
|
||||
for hipMemRangeAttributeReadMostly and hipMemRangeAttributePreferredLocation
|
||||
using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly.
|
||||
In other words setting of hipMemRangeAttributePreferredLocation should not
|
||||
impact hipMemRangeAttributeReadMostly advise to the memory*/
|
||||
|
||||
TEST_CASE("Unit_hipMemAdvise_TstMemAdviseMultiFlag") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int *Hmm = NULL, data = 999;
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, 4096));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetReadMostly, 0));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly, Hmm,
|
||||
4096));
|
||||
if (data != 1) {
|
||||
WARN("Didnt receive expected value at line: " << data);
|
||||
REQUIRE(false);
|
||||
}
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributePreferredLocation, Hmm,
|
||||
4096));
|
||||
if (data != 0) {
|
||||
WARN("Didnt receive expected value at line: " << data);
|
||||
REQUIRE(false);
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
/*Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and launch kernel
|
||||
on all other gpus except 0. This test case may discover any effect or
|
||||
access denial case arising due to setting ReadMostly only to a particular
|
||||
gpu*/
|
||||
|
||||
TEST_CASE("Unit_hipMemAdvise_ReadMosltyMgpuTst") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int Ngpus = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&Ngpus));
|
||||
if (Ngpus < 2) {
|
||||
SUCCEED("This test needs atleast two gpus to run."
|
||||
"Hence skipping the test.\n");
|
||||
}
|
||||
int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64;
|
||||
int *Hmm1 = NULL, DataMismatch = 0;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int))));
|
||||
// Initializing memory
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
Hmm[i] = InitVal;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)),
|
||||
hipMemAdviseSetReadMostly, 0));
|
||||
dim3 dimBlock(blockSize, 1, 1);
|
||||
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
|
||||
#if HT_AMD
|
||||
SECTION("Launch Kernel on all other gpus") {
|
||||
// launching kernel from each one of the gpus
|
||||
for (int i = 1; i < Ngpus; ++i) {
|
||||
DataMismatch = 0;
|
||||
HIP_CHECK(hipSetDevice(i));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int))));
|
||||
MemAdvise3<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Hmm1, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// verifying the results
|
||||
for (int j = 0; j < NumElms; ++j) {
|
||||
if (Hmm1[j] != (InitVal + 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
WARN("DataMismatch is observed with the gpu: " << i);
|
||||
REQUIRE(false);
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm1));
|
||||
}
|
||||
}
|
||||
|
||||
SECTION("Launch Kernel on all other gpus and manipulate the content") {
|
||||
for (int i = 0; i < Ngpus; ++i) {
|
||||
DataMismatch = 0;
|
||||
HIP_CHECK(hipSetDevice(i));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)),
|
||||
hipMemAdviseSetReadMostly, i));
|
||||
MemAdvise2<<<dimGrid, dimBlock, 0, strm>>>(Hmm, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
}
|
||||
// verifying the final result
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
if (Hmm[i] != (InitVal + Ngpus * 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
|
||||
if (DataMismatch != 0) {
|
||||
WARN("DataMismatch is observed at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
#endif
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
TEST_CASE("Unit_hipMemAdvise_TstSetUnsetPrfrdLoc") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int *Hmm = NULL, data = 999;
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, 2*4096));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributePreferredLocation, Hmm, 2*4096));
|
||||
if (data != 0) {
|
||||
WARN("Didnt receive expected value!!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetPreferredLocation, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributePreferredLocation, Hmm, 2*4096));
|
||||
if (data != -2) {
|
||||
WARN("Didnt receive expected value!!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
@@ -14,6 +14,8 @@ set(TEST_SRC
|
||||
hipStreamValue.cc
|
||||
hipStreamWithCUMask.cc
|
||||
hipStreamACb_MultiThread.cc
|
||||
hipStreamSynchronize.cc
|
||||
hipStreamQuery.cc
|
||||
hipStreamWaitEvent.cc
|
||||
)
|
||||
else()
|
||||
@@ -32,6 +34,8 @@ set(TEST_SRC
|
||||
# Fixing would break ABI, to be re-enabled when the fix is made.
|
||||
streamCommon.cc
|
||||
hipStreamValue.cc
|
||||
hipStreamSynchronize.cc
|
||||
hipStreamQuery.cc
|
||||
)
|
||||
|
||||
# set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17)
|
||||
@@ -39,4 +43,5 @@ endif()
|
||||
|
||||
hip_add_exe_to_target(NAME StreamTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests)
|
||||
TEST_TARGET_NAME build_tests
|
||||
COMPILE_OPTIONS -std=c++17)
|
||||
|
||||
@@ -509,8 +509,8 @@ bool validateStreamPrioritiesWithEvents() {
|
||||
#define OP(x) \
|
||||
free(src_h_##x); \
|
||||
free(dst_h_##x); \
|
||||
hipFree(src_d_##x); \
|
||||
hipFree(dst_d_##x);
|
||||
HIP_CHECK(hipFree(src_d_##x)); \
|
||||
HIP_CHECK(hipFree(dst_d_##x));
|
||||
OP(low)
|
||||
OP(normal)
|
||||
OP(high)
|
||||
|
||||
@@ -0,0 +1,125 @@
|
||||
/*
|
||||
Copyright (c) 2022 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, INNCLUDING 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 ANNY 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_test_common.hh>
|
||||
#include "streamCommon.hh"
|
||||
|
||||
/**
|
||||
* @brief Check that querying a stream with no work returns hipSuccess
|
||||
*
|
||||
**/
|
||||
TEST_CASE("Unit_hipStreamQuery_WithNoWork") {
|
||||
hipStream_t stream{nullptr};
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipStreamQuery(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Check that querying a stream with finished work returns hipSuccess
|
||||
*
|
||||
**/
|
||||
TEST_CASE("Unit_hipStreamQuery_WithFinishedWork") {
|
||||
hipStream_t stream{nullptr};
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
hip::stream::empty_kernel<<<dim3(1), dim3(1), 0, stream>>>();
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
|
||||
HIP_CHECK(hipStreamQuery(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
#if !HT_NVIDIA
|
||||
/**
|
||||
* @brief Check that submitting work to a destroyed stream sets its status as
|
||||
* hipErrorContextIsDestroyed
|
||||
*
|
||||
* Test removed for Nvidia devices because it returns unexpected error
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamQuery_WithDestroyedStream") {
|
||||
hipStream_t stream{nullptr};
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorContextIsDestroyed);
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Check that submitting work to an uninitialized stream sets its status as
|
||||
* hipErrorContextIsDestroyed
|
||||
*
|
||||
* Test removed for Nvidia devices because it returns unexpected error
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamQuery_WithUninitializedStream") {
|
||||
hipStream_t stream{reinterpret_cast<hipStream_t>(0xFFFF)};
|
||||
HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorContextIsDestroyed);
|
||||
}
|
||||
#endif
|
||||
|
||||
#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */
|
||||
|
||||
/**
|
||||
* @brief Check that submitting work to a stream sets the status of the nullStream to
|
||||
* hipErrorNotReady
|
||||
*
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamQuery_SubmitWorkOnStreamAndQueryNullStream") {
|
||||
{
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIP_CHECK(hipStreamQuery(hip::nullStream));
|
||||
HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream);
|
||||
HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Check that submitting work to the nullStream properly sets its status as
|
||||
* hipErrorNotReady.
|
||||
*
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamQuery_NullStreamQuery") {
|
||||
HIP_CHECK(hipStreamQuery(hip::nullStream));
|
||||
HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::nullStream);
|
||||
HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(hip::nullStream));
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Check that querying a stream with pending work returns hipErrorNotReady
|
||||
*
|
||||
**/
|
||||
TEST_CASE("Unit_hipStreamQuery_WithPendingWork") {
|
||||
hipStream_t waitingStream{nullptr};
|
||||
HIP_CHECK(hipStreamCreate(&waitingStream));
|
||||
|
||||
HipTest::runKernelForDuration(std::chrono::milliseconds(500), waitingStream);
|
||||
|
||||
HIP_CHECK_ERROR(hipStreamQuery(waitingStream), hipErrorNotReady);
|
||||
HIP_CHECK(hipStreamSynchronize(waitingStream));
|
||||
HIP_CHECK(hipStreamQuery(waitingStream));
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(waitingStream));
|
||||
}
|
||||
#endif
|
||||
@@ -0,0 +1,156 @@
|
||||
/*
|
||||
Copyright (c) 2022 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, INNCLUDING 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 ANNY 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_test_common.hh>
|
||||
#include "streamCommon.hh"
|
||||
|
||||
namespace hipStreamSynchronizeTest {
|
||||
|
||||
/**
|
||||
* @brief Check that hipStreamSynchronize handles empty streams properly.
|
||||
*
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamSynchronize_EmptyStream") {
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */
|
||||
|
||||
/**
|
||||
* @brief Check that all work executing in a stream is finished after a call to
|
||||
* hipStreamSynchronize.
|
||||
*
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamSynchronize_FinishWork") {
|
||||
const hipStream_t explicitStream = reinterpret_cast<hipStream_t>(-1);
|
||||
hipStream_t stream = GENERATE_COPY(explicitStream, hip::nullStream, hip::streamPerThread);
|
||||
|
||||
const bool isExplicitStream = stream == explicitStream;
|
||||
if (isExplicitStream) {
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
}
|
||||
|
||||
HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream);
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamQuery(stream));
|
||||
|
||||
if (isExplicitStream) {
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Check that synchronizing the nullStream implicitly synchronizes all executing streams.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamSynchronize_NullStreamSynchronization") {
|
||||
int totalStreams = 10;
|
||||
|
||||
std::vector<hipStream_t> streams{};
|
||||
|
||||
for (int i = 0; i < totalStreams; ++i) {
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
streams.push_back(stream);
|
||||
}
|
||||
|
||||
for (int i = 0; i < totalStreams; ++i) {
|
||||
HipTest::runKernelForDuration(std::chrono::milliseconds(1000), streams[i]);
|
||||
}
|
||||
|
||||
for (int i = 0; i < totalStreams; ++i) {
|
||||
HIP_CHECK_ERROR(hipStreamQuery(streams[i]), hipErrorNotReady);
|
||||
}
|
||||
|
||||
HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(hip::nullStream));
|
||||
HIP_CHECK(hipStreamQuery(hip::nullStream));
|
||||
|
||||
for (int i = 0; i < totalStreams; ++i) {
|
||||
HIP_CHECK(hipStreamQuery(streams[i]));
|
||||
}
|
||||
|
||||
for (int i = 0; i < totalStreams; ++i) {
|
||||
HIP_CHECK(hipStreamDestroy(streams[i]));
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Check that synchronizing one stream does implicitly synchronize other streams.
|
||||
* Check that submiting work to the nullStream does not affect synchronization of other
|
||||
* streams. Check that querying the nullStream does not affect synchronization of other streams.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamSynchronize_SynchronizeStreamAndQueryNullStream") {
|
||||
#if HT_AMD
|
||||
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-22");
|
||||
#else
|
||||
|
||||
hipStream_t stream1;
|
||||
hipStream_t stream2;
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream1));
|
||||
HIP_CHECK(hipStreamCreate(&stream2));
|
||||
|
||||
HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream1);
|
||||
HipTest::runKernelForDuration(std::chrono::milliseconds(2000), stream2);
|
||||
|
||||
SECTION("Do not use NullStream") {}
|
||||
SECTION("Submit Kernel to NullStream") {
|
||||
hip::stream::empty_kernel<<<1, 1, 0, hip::nullStream> > >();
|
||||
}
|
||||
SECTION("Query NullStream") {
|
||||
HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
|
||||
}
|
||||
|
||||
HIP_CHECK_ERROR(hipStreamQuery(stream1), hipErrorNotReady);
|
||||
HIP_CHECK_ERROR(hipStreamQuery(stream2), hipErrorNotReady);
|
||||
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream1));
|
||||
HIP_CHECK(hipStreamQuery(stream1));
|
||||
HIP_CHECK_ERROR(hipStreamQuery(stream2), hipErrorNotReady);
|
||||
HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream2));
|
||||
HIP_CHECK(hipStreamQuery(stream2));
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream1));
|
||||
HIP_CHECK(hipStreamDestroy(stream2));
|
||||
#endif
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief Check that synchronizing the nullStream also synchronizes the hipStreamPerThread
|
||||
* special stream.
|
||||
*
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamSynchronize_NullStreamAndStreamPerThread") {
|
||||
HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::streamPerThread);
|
||||
HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
|
||||
HIP_CHECK_ERROR(hipStreamQuery(hip::streamPerThread), hipErrorNotReady);
|
||||
HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::nullStream);
|
||||
HIP_CHECK(hipStreamSynchronize(hip::nullStream))
|
||||
HIP_CHECK_ERROR(hipStreamQuery(hip::streamPerThread), hipSuccess);
|
||||
HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipSuccess);
|
||||
}
|
||||
#endif
|
||||
} // namespace hipStreamSynchronizeTest
|
||||
@@ -17,276 +17,377 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#include <hip_test_common.hh>
|
||||
#include <memory>
|
||||
#include <type_traits>
|
||||
|
||||
constexpr unsigned int writeFlag = 0;
|
||||
|
||||
#define DEFINE_HIP_STREAM_VALUE(TYPE, BITS, ...) hipStream##TYPE##Value##BITS(__VA_ARGS__)
|
||||
template <typename UIntT, typename... Args> auto waitFunc(Args... args) {
|
||||
if constexpr (std::is_same<UIntT, uint32_t>::value) {
|
||||
return hipStreamWaitValue32(args...);
|
||||
} else {
|
||||
return hipStreamWaitValue64(args...);
|
||||
}
|
||||
};
|
||||
|
||||
#define CHECK_HIP_STREAM_VALUE(TYPE, BITS, ...) \
|
||||
HIP_CHECK(DEFINE_HIP_STREAM_VALUE(TYPE, BITS, __VA_ARGS__));
|
||||
template <typename UIntT, typename... Args> auto writeFunc(Args... args) {
|
||||
if constexpr (std::is_same<UIntT, uint32_t>::value) {
|
||||
return hipStreamWriteValue32(args...);
|
||||
} else {
|
||||
return hipStreamWriteValue64(args...);
|
||||
}
|
||||
};
|
||||
|
||||
#define NEG_TEST_ERROR_CHECK(TYPE, BITS, errorCode, ...) \
|
||||
HIP_CHECK_ERROR(DEFINE_HIP_STREAM_VALUE(TYPE, BITS, __VA_ARGS__), errorCode);
|
||||
// Random predefined 32 and 64 bit values
|
||||
using value32_t = std::integral_constant<uint32_t, 0x70F0F0FF>;
|
||||
using value64_t = std::integral_constant<uint64_t, 0x7FFF0000FFFF0000>;
|
||||
template <typename UIntT>
|
||||
using testValue =
|
||||
typename std::conditional<std::is_same<UIntT, uint32_t>::value, value32_t, value64_t>::type;
|
||||
|
||||
#if HT_AMD
|
||||
// Random predefiend 32 and 64 bit values
|
||||
constexpr uint32_t value32 = 0x70F0F0FF;
|
||||
constexpr uint64_t value64 = 0x7FFF0000FFFF0000;
|
||||
constexpr uint32_t DATA_INIT = 0x1234;
|
||||
constexpr uint32_t DATA_UPDATE = 0X4321;
|
||||
|
||||
template <typename intT> struct TEST_WAIT {
|
||||
using uintT = typename std::make_unsigned<intT>::type;
|
||||
int compareOp;
|
||||
uintT mask;
|
||||
uintT waitValue;
|
||||
intT signalValueFail;
|
||||
intT signalValuePass;
|
||||
template <typename UIntT> struct TEST_WAIT {
|
||||
static_assert(std::is_same<UIntT, uint32_t>::value or std::is_same<UIntT, uint64_t>::value,
|
||||
"only implemented for 32 bit and 64 bit unsigned integers");
|
||||
unsigned int compareOp;
|
||||
UIntT mask = ~static_cast<UIntT>(0);
|
||||
UIntT waitValue;
|
||||
UIntT signalValueFail;
|
||||
UIntT signalValuePass;
|
||||
|
||||
TEST_WAIT(int compareOp, uintT waitValue, intT signalValueFail, intT signalValuePass)
|
||||
TEST_WAIT(unsigned int compareOp, UIntT waitValue, UIntT signalValueFail, UIntT signalValuePass)
|
||||
: compareOp{compareOp},
|
||||
waitValue{waitValue},
|
||||
signalValueFail{signalValueFail},
|
||||
signalValuePass{signalValuePass} {
|
||||
mask = static_cast<uintT>(0xFFFFFFFFFFFFFFFF);
|
||||
}
|
||||
signalValuePass{signalValuePass} {}
|
||||
|
||||
TEST_WAIT(int compareOp, uintT mask, uintT waitValue, intT signalValueFail, intT signalValuePass)
|
||||
TEST_WAIT(unsigned int compareOp, UIntT mask, UIntT waitValue, UIntT signalValueFail,
|
||||
UIntT signalValuePass)
|
||||
: compareOp{compareOp},
|
||||
mask{mask},
|
||||
waitValue{waitValue},
|
||||
signalValueFail{signalValueFail},
|
||||
signalValuePass{signalValuePass} {}
|
||||
};
|
||||
typedef TEST_WAIT<int32_t> TEST_WAIT32;
|
||||
typedef TEST_WAIT<int64_t> TEST_WAIT64;
|
||||
|
||||
using TEST_WAIT32 = TEST_WAIT<uint32_t>;
|
||||
using TEST_WAIT64 = TEST_WAIT<uint64_t>;
|
||||
|
||||
bool streamWaitValueSupported() {
|
||||
int device_num = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&device_num));
|
||||
int waitValueSupport;
|
||||
for (int device_id = 0; device_id < device_num; ++device_id) {
|
||||
HIP_CHECK(hipSetDevice(device_id));
|
||||
waitValueSupport = 0;
|
||||
HIP_CHECK(hipDeviceGetAttribute(&waitValueSupport, hipDeviceAttributeCanUseStreamWaitValue,
|
||||
device_id));
|
||||
int waitValueSupport = 0;
|
||||
auto getAttributeError = hipDeviceGetAttribute(
|
||||
&waitValueSupport, hipDeviceAttributeCanUseStreamWaitValue, device_id);
|
||||
if (getAttributeError != hipSuccess) {
|
||||
HipTest::HIP_SKIP_TEST("attribute not supported");
|
||||
return false;
|
||||
}
|
||||
if (waitValueSupport == 1) return true;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
// hipStreamWriteValue Tests
|
||||
TEST_CASE("Unit_hipStreamValue_Write") {
|
||||
int64_t* signalPtr;
|
||||
// The different types of memory that can be used with hipStream[Wait|Write]
|
||||
enum class PtrType { HostPtr, DevicePtr, DevicePtrToHost, Signal };
|
||||
|
||||
hipStream_t stream{nullptr};
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
// Helper class to expose the pointer that is used with hipStream[Write|Wait]Value and also store a
|
||||
// unique pointer with the deleter to simplify cleanup
|
||||
// Also includes functions to update and get the value directly
|
||||
template <PtrType type, typename UIntT, typename UniquePtrWithDeleter> class TestPtr {
|
||||
// This stores the memory that must be deleted, as well as the deleter
|
||||
UniquePtrWithDeleter ptrToDelete;
|
||||
|
||||
// Allocate Host Memory
|
||||
auto hostPtr64 = std::unique_ptr<uint64_t>(new uint64_t(1));
|
||||
auto hostPtr32 = std::unique_ptr<uint32_t>(new uint32_t(1));
|
||||
public:
|
||||
// The pointer that should be used with hipStream[Write|Wait]Value
|
||||
UIntT* ptr;
|
||||
|
||||
// Register Host Memory
|
||||
HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0));
|
||||
HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0));
|
||||
TestPtr(UIntT* ptr, UniquePtrWithDeleter ptrToDelete)
|
||||
: ptrToDelete(std::move(ptrToDelete)), ptr(ptr) {}
|
||||
|
||||
// Register Signal Memory
|
||||
HIP_CHECK(hipExtMallocWithFlags((void**)&signalPtr, 8, hipMallocSignalMemory));
|
||||
|
||||
// Initialise Data
|
||||
*signalPtr = 0x0;
|
||||
*hostPtr64 = 0x0;
|
||||
*hostPtr32 = 0x0;
|
||||
|
||||
SECTION("Registered host memory hipStreamWriteValue32") {
|
||||
INFO("Test writting to registered host pointer using hipStreamWriteValue32");
|
||||
HIP_CHECK(hipStreamWriteValue32(stream, hostPtr32.get(), value32, writeFlag));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_ASSERT(*hostPtr32 == value32);
|
||||
// directly retrieve the value from wherever it was allocated
|
||||
UIntT getValue(size_t offset = 0) {
|
||||
if constexpr (type == PtrType::Signal || type == PtrType::HostPtr ||
|
||||
type == PtrType::DevicePtrToHost) {
|
||||
return ptrToDelete.get()[offset];
|
||||
} else {
|
||||
static_assert(type == PtrType::DevicePtr, "Expected DevicePtr");
|
||||
UIntT value;
|
||||
HIP_CHECK(hipMemcpy(&value, ptr + offset, sizeof(UIntT), hipMemcpyDeviceToHost));
|
||||
return value;
|
||||
}
|
||||
}
|
||||
|
||||
SECTION("Registered host memory hipStreamWriteValue64") {
|
||||
INFO("Test writting to registered host pointer using hipStreamWriteValue32");
|
||||
HIP_CHECK(hipStreamWriteValue64(stream, hostPtr64.get(), value64, writeFlag));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_ASSERT(*hostPtr64 == value64);
|
||||
// directly set the value wherever it was allocated
|
||||
void setValue(UIntT value, size_t offset = 0) {
|
||||
if constexpr (type == PtrType::Signal || type == PtrType::DevicePtrToHost ||
|
||||
type == PtrType::HostPtr) {
|
||||
ptrToDelete.get()[offset] = value;
|
||||
} else {
|
||||
// hipMemcpy causes deadlock, so use hipStreamWriteValue
|
||||
static_assert(type == PtrType::DevicePtr, "Expected DevicePtr");
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
|
||||
HIP_CHECK(writeFunc<UIntT>(stream, ptr + offset, value, writeFlag));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
// Test writting device pointer
|
||||
void* devicePtr64;
|
||||
void* devicePtr32;
|
||||
HIP_CHECK(hipHostGetDevicePointer((void**)&devicePtr64, hostPtr64.get(), 0));
|
||||
HIP_CHECK(hipHostGetDevicePointer((void**)&devicePtr32, hostPtr32.get(), 0));
|
||||
// Reset values
|
||||
*hostPtr64 = 0x0;
|
||||
*hostPtr32 = 0x0;
|
||||
// required for the static assert
|
||||
template <PtrType> inline constexpr bool AMD_ACTIVE = HT_AMD == 1;
|
||||
|
||||
SECTION("Device Memory hipStreamWriteValue32") {
|
||||
INFO("Test writting to device pointer using hipStreamWriteValue32");
|
||||
HIP_CHECK(hipStreamWriteValue32(stream, devicePtr32, value32, writeFlag));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_ASSERT(*hostPtr32 == value32);
|
||||
template <PtrType type, typename UIntT> auto allocMem() {
|
||||
constexpr std::size_t arraySize = 1024;
|
||||
if constexpr (type == PtrType::Signal) {
|
||||
static_assert(std::is_same<UIntT, uint64_t>::value,
|
||||
"signal memory should only be used with 64bit memory");
|
||||
|
||||
// Allocate Signal Memory
|
||||
uint64_t* signalPtr{};
|
||||
|
||||
static_assert(AMD_ACTIVE<type>,
|
||||
"nvidia backend compiler doesn't like hipExtMallocWithFlags, even in this "
|
||||
"constexpr branch");
|
||||
#if HT_AMD
|
||||
// 8 is the only acceptable size
|
||||
HIP_CHECK(
|
||||
hipExtMallocWithFlags(reinterpret_cast<void**>(&signalPtr), 8, hipMallocSignalMemory));
|
||||
#endif
|
||||
|
||||
// Init Memory
|
||||
*signalPtr = 0;
|
||||
|
||||
auto freeStuff = [](uint64_t* sPtr) { HIP_CHECK(hipFree(sPtr)); };
|
||||
return TestPtr<type, UIntT, std::unique_ptr<uint64_t, decltype(freeStuff)>>{
|
||||
signalPtr, std::unique_ptr<uint64_t, decltype(freeStuff)>(signalPtr, freeStuff)};
|
||||
} else if constexpr (type == PtrType::DevicePtrToHost) {
|
||||
auto hostPtr = new UIntT[arraySize];
|
||||
|
||||
// Register Host Memory
|
||||
HIP_CHECK(hipHostRegister(hostPtr, sizeof(UIntT) * arraySize, 0));
|
||||
|
||||
// Init memory
|
||||
std::fill(hostPtr, hostPtr + arraySize, 0);
|
||||
|
||||
UIntT* devicePtr;
|
||||
// Test writing device pointer
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&devicePtr), hostPtr, 0));
|
||||
auto freeStuff = [](UIntT* ptr) {
|
||||
HIP_CHECK(hipHostUnregister(ptr));
|
||||
delete[] ptr;
|
||||
};
|
||||
|
||||
return TestPtr<type, UIntT, std::unique_ptr<UIntT[], decltype(freeStuff)>>{
|
||||
devicePtr, std::unique_ptr<UIntT[], decltype(freeStuff)>(hostPtr, freeStuff)};
|
||||
} else if constexpr (type == PtrType::HostPtr) {
|
||||
auto hostPtr = new UIntT[arraySize];
|
||||
|
||||
// Register Host Memory
|
||||
HIP_CHECK(hipHostRegister(hostPtr, sizeof(UIntT) * arraySize, 0));
|
||||
|
||||
// Init memory
|
||||
std::fill(hostPtr, hostPtr + arraySize, 0);
|
||||
|
||||
auto freeStuff = [](UIntT* ptr) {
|
||||
HIP_CHECK(hipHostUnregister(ptr));
|
||||
delete[] ptr;
|
||||
};
|
||||
|
||||
return TestPtr<type, UIntT, std::unique_ptr<UIntT[], decltype(freeStuff)>>{
|
||||
hostPtr, std::unique_ptr<UIntT[], decltype(freeStuff)>(hostPtr, freeStuff)};
|
||||
} else {
|
||||
static_assert(type == PtrType::DevicePtr, "Expected DevicePtr");
|
||||
UIntT* devicePtr;
|
||||
HIP_CHECK(hipMalloc(&devicePtr, sizeof(UIntT) * arraySize));
|
||||
HIP_CHECK(hipMemset(devicePtr, 0, sizeof(UIntT) * arraySize));
|
||||
auto freeStuff = [](UIntT* ptr) { HIP_CHECK(hipFree(ptr)); };
|
||||
return TestPtr<type, UIntT, std::unique_ptr<UIntT, decltype(freeStuff)>>{
|
||||
devicePtr, std::unique_ptr<UIntT, decltype(freeStuff)>(devicePtr, freeStuff)};
|
||||
}
|
||||
|
||||
SECTION("Device Memory hipStreamWriteValue64") {
|
||||
INFO("Test writting to device pointer using hipStreamWriteValue64");
|
||||
HIP_CHECK(hipStreamWriteValue64(stream, devicePtr64, value64, writeFlag));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_ASSERT(*hostPtr64 == value64);
|
||||
}
|
||||
|
||||
// Test Writing to Signal Memory
|
||||
SECTION("Signal Memory hipStreamWriteValue64") {
|
||||
INFO("Test writting to signal memory using hipStreamWriteValue64");
|
||||
HIP_CHECK(hipStreamWriteValue64(stream, signalPtr, value64, writeFlag));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_ASSERT(*signalPtr == value64);
|
||||
}
|
||||
|
||||
// Cleanup
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK(hipHostUnregister(hostPtr64.get()));
|
||||
HIP_CHECK(hipHostUnregister(hostPtr32.get()));
|
||||
HIP_CHECK(hipFree(signalPtr));
|
||||
}
|
||||
|
||||
// hipStreamWaitValue Tests
|
||||
template <bool isBlocking, typename intT, typename TEST_T>
|
||||
void initData(intT* dataPtr, int64_t* signalPtr, TEST_T tc, std::vector<hipEvent_t>& events) {
|
||||
// Initialize memory to be waited on
|
||||
*signalPtr = isBlocking ? tc.signalValueFail : tc.signalValuePass;
|
||||
|
||||
|
||||
// Initialize host pointers
|
||||
dataPtr[0] = DATA_INIT;
|
||||
dataPtr[1] = DATA_INIT;
|
||||
|
||||
|
||||
hipEvent_t firstWriteEvent{nullptr};
|
||||
hipEvent_t secondWriteEvent{nullptr};
|
||||
HIP_CHECK(hipEventCreate(&firstWriteEvent));
|
||||
HIP_CHECK(hipEventCreate(&secondWriteEvent));
|
||||
events.push_back(firstWriteEvent);
|
||||
events.push_back(secondWriteEvent);
|
||||
}
|
||||
|
||||
template <bool isBlocking, typename intT, typename TEST_T>
|
||||
void syncAndCheckData(hipStream_t stream, intT* dataPtr, int64_t* signalPtr, TEST_T tc,
|
||||
std::vector<hipEvent_t>& events) {
|
||||
// Ensure first part of host memory is updated
|
||||
HIP_CHECK(hipStreamWaitEvent(stream, events[0], 0));
|
||||
HIP_ASSERT(dataPtr[0] == DATA_UPDATE);
|
||||
if (isBlocking) {
|
||||
// Ensure second part of host memory isn't updated yet
|
||||
HIP_ASSERT(hipEventQuery(events[1]) == hipErrorNotReady);
|
||||
HIP_ASSERT(dataPtr[1] == DATA_INIT);
|
||||
// Update value to release stream
|
||||
*signalPtr = tc.signalValuePass;
|
||||
// allows the creation of a list of offsets while avoiding it for signal memory
|
||||
template <PtrType type> constexpr auto get_offsets() {
|
||||
if constexpr (type == PtrType::Signal) {
|
||||
return std::array<size_t, 1>{0};
|
||||
} else {
|
||||
return std::array<size_t, 6>{0, 1, 2, 3, 31, 1023};
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_ASSERT(hipEventQuery(events[1]) == hipSuccess);
|
||||
// Finally ensure that second part of host memory is updated
|
||||
HIP_ASSERT(dataPtr[1] == DATA_UPDATE);
|
||||
}
|
||||
|
||||
template <typename intT> void cleanup(hipStream_t& stream, intT* dataPtr, int64_t* signalPtr) {
|
||||
// Cleanup
|
||||
HIP_CHECK(hipFree(signalPtr));
|
||||
HIP_CHECK(hipHostUnregister(dataPtr));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
template <typename UIntT, PtrType ptrTypeValue> struct TestParams {
|
||||
using UIntType = UIntT;
|
||||
constexpr static PtrType ptrType = ptrTypeValue;
|
||||
};
|
||||
|
||||
template <typename intT, bool isBlocking, typename TEST_T> void testWait(TEST_T tc) {
|
||||
#if HT_AMD
|
||||
TEMPLATE_TEST_CASE("Unit_hipStreamValue_Write", "", (TestParams<uint32_t, PtrType::HostPtr>),
|
||||
(TestParams<uint32_t, PtrType::DevicePtr>),
|
||||
(TestParams<uint32_t, PtrType::DevicePtrToHost>),
|
||||
(TestParams<uint64_t, PtrType::HostPtr>),
|
||||
(TestParams<uint64_t, PtrType::DevicePtr>),
|
||||
(TestParams<uint64_t, PtrType::DevicePtrToHost>),
|
||||
(TestParams<uint64_t, PtrType::Signal>)) {
|
||||
#else
|
||||
TEMPLATE_TEST_CASE("Unit_hipStreamValue_Write", "", (TestParams<uint32_t, PtrType::HostPtr>),
|
||||
(TestParams<uint32_t, PtrType::DevicePtr>),
|
||||
(TestParams<uint32_t, PtrType::DevicePtrToHost>),
|
||||
(TestParams<uint64_t, PtrType::HostPtr>),
|
||||
(TestParams<uint64_t, PtrType::DevicePtr>),
|
||||
(TestParams<uint64_t, PtrType::DevicePtrToHost>)) {
|
||||
#endif
|
||||
#if HT_AMD
|
||||
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-126");
|
||||
return;
|
||||
#endif
|
||||
if (!streamWaitValueSupported()) {
|
||||
UNSCOPED_INFO(" hipStreamWaitValue: not supported on this device , skipping ...");
|
||||
HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device.");
|
||||
return;
|
||||
}
|
||||
|
||||
// Initialize stream
|
||||
using UIntT = typename TestType::UIntType;
|
||||
constexpr auto ptrType = TestType::ptrType;
|
||||
constexpr auto writeValue = testValue<UIntT>::value;
|
||||
|
||||
hipStream_t stream{nullptr};
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
const auto offsets = get_offsets<ptrType>();
|
||||
const auto offset = GENERATE_COPY(from_range(std::begin(offsets), std::end(offsets)));
|
||||
|
||||
CAPTURE(offset);
|
||||
|
||||
// Allocate Host Memory
|
||||
std::unique_ptr<intT> dataPtr(new intT(2));
|
||||
auto ptr = allocMem<ptrType, UIntT>();
|
||||
UIntT* target = ptr.ptr + offset;
|
||||
HIP_CHECK(writeFunc<UIntT>(stream, target, writeValue, writeFlag));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
REQUIRE(ptr.getValue(offset) == writeValue);
|
||||
|
||||
// Register Host Memory
|
||||
HIP_CHECK(hipHostRegister(&(dataPtr.get()[0]), sizeof(intT), 0));
|
||||
HIP_CHECK(hipHostRegister(&(dataPtr.get()[1]), sizeof(intT), 0));
|
||||
// Cleanup
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
// Allocate Signal Memory
|
||||
int64_t* signalPtr;
|
||||
HIP_CHECK(hipExtMallocWithFlags((void**)&signalPtr, 8, hipMallocSignalMemory));
|
||||
template <bool isBlocking, typename UIntT, typename TestPtr>
|
||||
void syncAndCheckData(hipStream_t stream, UIntT* dataPtr, TestPtr signalPtr, size_t offset,
|
||||
TEST_WAIT<UIntT> tc, std::array<hipEvent_t, 2>& events) {
|
||||
// Ensure first part of host memory is updated
|
||||
HIP_CHECK(hipEventSynchronize(events[0]));
|
||||
REQUIRE(dataPtr[0] == DATA_UPDATE);
|
||||
|
||||
std::vector<hipEvent_t> events;
|
||||
initData<isBlocking>(dataPtr.get(), signalPtr, tc, events);
|
||||
|
||||
if (std::is_same<intT, int32_t>::value) {
|
||||
CHECK_HIP_STREAM_VALUE(Write, 32, stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag)
|
||||
HIP_CHECK(hipEventRecord(events[0], stream));
|
||||
|
||||
if (static_cast<uint32_t>(tc.mask) != 0xFFFFFFFF) {
|
||||
CHECK_HIP_STREAM_VALUE(Wait, 32, stream, signalPtr, static_cast<uint32_t>(tc.waitValue),
|
||||
tc.compareOp, static_cast<uint32_t>(tc.mask));
|
||||
} else {
|
||||
CHECK_HIP_STREAM_VALUE(Wait, 32, stream, signalPtr, tc.waitValue, tc.compareOp);
|
||||
}
|
||||
|
||||
CHECK_HIP_STREAM_VALUE(Write, 32, stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag)
|
||||
} else {
|
||||
CHECK_HIP_STREAM_VALUE(Write, 64, stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag)
|
||||
HIP_CHECK(hipEventRecord(events[0], stream));
|
||||
|
||||
if (tc.mask != 0xFFFFFFFFFFFFFFFF) {
|
||||
CHECK_HIP_STREAM_VALUE(Wait, 64, stream, signalPtr, tc.waitValue, tc.compareOp, tc.mask);
|
||||
} else {
|
||||
CHECK_HIP_STREAM_VALUE(Wait, 64, stream, signalPtr, tc.waitValue, tc.compareOp);
|
||||
}
|
||||
|
||||
CHECK_HIP_STREAM_VALUE(Write, 64, stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag)
|
||||
if constexpr (isBlocking) {
|
||||
// Ensure second part of host memory isn't updated yet
|
||||
HIP_CHECK_ERROR(hipEventQuery(events[1]), hipErrorNotReady);
|
||||
REQUIRE(dataPtr[1] == DATA_INIT);
|
||||
// Update value to release stream
|
||||
signalPtr.setValue(tc.signalValuePass, offset);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipEventQuery(events[1]));
|
||||
// Finally ensure that second part of host memory is updated
|
||||
REQUIRE(dataPtr[1] == DATA_UPDATE);
|
||||
}
|
||||
|
||||
|
||||
template <typename TestType, bool isBlocking>
|
||||
void testWait(TEST_WAIT<typename TestType::UIntType> tc) {
|
||||
if (!streamWaitValueSupported()) {
|
||||
HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device.");
|
||||
return;
|
||||
}
|
||||
#if HT_AMD
|
||||
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-128");
|
||||
return;
|
||||
#endif
|
||||
|
||||
using UIntT = typename TestType::UIntType;
|
||||
constexpr auto ptrType = TestType::ptrType;
|
||||
constexpr UIntT defaultMask = ~static_cast<UIntT>(0);
|
||||
|
||||
// Initialize stream
|
||||
hipStream_t stream{};
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
// Allocate Host Memory
|
||||
auto dataPtr = std::make_unique<UIntT[]>(2);
|
||||
// Register Host Memory
|
||||
HIP_CHECK(hipHostRegister(dataPtr.get(), sizeof(UIntT), 0));
|
||||
HIP_CHECK(hipHostRegister(dataPtr.get() + 1, sizeof(UIntT), 0));
|
||||
std::fill(dataPtr.get(), dataPtr.get() + 2, DATA_INIT);
|
||||
|
||||
std::array<hipEvent_t, 2> events;
|
||||
HIP_CHECK(hipEventCreate(&events[0]));
|
||||
HIP_CHECK(hipEventCreate(&events[1]));
|
||||
|
||||
|
||||
const auto offsets = get_offsets<ptrType>();
|
||||
const auto offset = GENERATE_COPY(from_range(std::begin(offsets), std::end(offsets)));
|
||||
|
||||
auto waitPtr = allocMem<ptrType, UIntT>();
|
||||
UIntT* const target = waitPtr.ptr + offset;
|
||||
waitPtr.setValue(isBlocking ? tc.signalValueFail : tc.signalValuePass, offset);
|
||||
|
||||
HIP_CHECK(writeFunc<UIntT>(stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag));
|
||||
HIP_CHECK(hipEventRecord(events[0], stream));
|
||||
|
||||
if (tc.mask != defaultMask) {
|
||||
HIP_CHECK(waitFunc<UIntT>(stream, target, tc.waitValue, tc.compareOp, tc.mask));
|
||||
} else {
|
||||
HIP_CHECK(waitFunc<UIntT>(stream, target, tc.waitValue, tc.compareOp));
|
||||
}
|
||||
|
||||
HIP_CHECK(writeFunc<UIntT>(stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag));
|
||||
|
||||
HIP_CHECK(hipEventRecord(events[1], stream));
|
||||
|
||||
syncAndCheckData<isBlocking>(stream, dataPtr.get(), signalPtr, tc, events);
|
||||
cleanup(stream, dataPtr.get(), signalPtr);
|
||||
}
|
||||
#undef CHECK_HIP_STREAM_VALUE
|
||||
syncAndCheckData<isBlocking>(stream, dataPtr.get(), std::move(waitPtr), offset, tc, events);
|
||||
|
||||
// Cleanup
|
||||
HIP_CHECK(hipEventDestroy(events[0]));
|
||||
HIP_CHECK(hipEventDestroy(events[1]));
|
||||
HIP_CHECK(hipHostUnregister(dataPtr.get()));
|
||||
HIP_CHECK(hipHostUnregister(dataPtr.get() + 1));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
// TEMPLATE_TEST_CASE wasn't working within a macro, so sections were used instead
|
||||
#define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32(suffix, test_t) \
|
||||
TEST_CASE("Unit_hipStreamValue_Wait32_Blocking_" + std::string(suffix)) { \
|
||||
testWait<int32_t, true>(test_t); \
|
||||
SECTION("HostPtr") { testWait<TestParams<uint32_t, PtrType::HostPtr>, true>(test_t); } \
|
||||
SECTION("DevicePtr") { testWait<TestParams<uint32_t, PtrType::DevicePtr>, true>(test_t); } \
|
||||
SECTION("DevicePtrToHost") { \
|
||||
testWait<TestParams<uint32_t, PtrType::DevicePtrToHost>, true>(test_t); \
|
||||
} \
|
||||
} \
|
||||
TEST_CASE("Unit_hipStreamValue_Wait32_NonBlocking_" + std::string(suffix)) { \
|
||||
testWait<int32_t, false>(test_t); \
|
||||
SECTION("HostPtr") { testWait<TestParams<uint32_t, PtrType::HostPtr>, false>(test_t); } \
|
||||
SECTION("DevicePtr") { testWait<TestParams<uint32_t, PtrType::DevicePtr>, false>(test_t); } \
|
||||
SECTION("DevicePtrToHost") { \
|
||||
testWait<TestParams<uint32_t, PtrType::DevicePtrToHost>, false>(test_t); \
|
||||
} \
|
||||
}
|
||||
|
||||
|
||||
// Using Mask
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte_1",
|
||||
TEST_WAIT64( // mask will ignore few MSB bits
|
||||
hipStreamWaitValueGte, 0x0000FFFFFFFFFFFF,
|
||||
0x000000007FFF0001, 0x7FFF00007FFF0000,
|
||||
0x000000007FFF0001))
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte_2",
|
||||
TEST_WAIT64(hipStreamWaitValueGte, 0xF, 0x4, 0x3, 0x6))
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte",
|
||||
TEST_WAIT32(hipStreamWaitValueGte, 0xF, 0x4, 0x3, 0x6))
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Eq_1",
|
||||
TEST_WAIT64( // mask will ignore few MSB bits
|
||||
hipStreamWaitValueEq, 0x0000FFFFFFFFFFFF,
|
||||
0x000000000FFF0001, 0x7FFF00000FFF0000,
|
||||
0x7F0000000FFF0001))
|
||||
TEST_WAIT32( // mask will ignore few MSB bits
|
||||
hipStreamWaitValueEq, 0x0000FFFF, 0x00000001,
|
||||
0x0FFF0000, 0x0FFF0001))
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Eq_2",
|
||||
TEST_WAIT64(hipStreamWaitValueEq, 0xFF, 0x11, 0x25, 0x11))
|
||||
TEST_WAIT32(hipStreamWaitValueEq, 0xFF, 0x11, 0x25, 0x11))
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_And",
|
||||
TEST_WAIT64( // mask will discard bits 8 to 11
|
||||
TEST_WAIT32( // mask will discard bits 8 to 11
|
||||
hipStreamWaitValueAnd, 0xFF, 0xF4A, 0xF35, 0X02))
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Nor_1",
|
||||
TEST_WAIT64( // mask is set to ignore the sign bit.
|
||||
hipStreamWaitValueNor, 0x7FFFFFFFFFFFFFFF,
|
||||
0x7FFFFFFFFFFFF247, 0x7FFFFFFFFFFFFdbd,
|
||||
0x7FFFFFFFFFFFFdb5))
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Nor_2",
|
||||
TEST_WAIT64( // mask is set to apply NOR for bits 0 to 3.
|
||||
hipStreamWaitValueNor, 0xF, 0x7E, 0x7D, 0x76))
|
||||
|
||||
// Not Using Mask
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_Eq",
|
||||
@@ -299,19 +400,47 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_And",
|
||||
TEST_WAIT32(hipStreamWaitValueAnd, 0x70F0F0F0, 0x0F0F0F0F,
|
||||
0X1F0F0F0F))
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_Nor",
|
||||
TEST_WAIT32(hipStreamWaitValueNor, 0x7AAAAAAA,
|
||||
static_cast<int32_t>(0x85555555),
|
||||
static_cast<int32_t>(0x9AAAAAAA)))
|
||||
TEST_WAIT32(hipStreamWaitValueNor, 0x7AAAAAAA, 0x85555555,
|
||||
0x9AAAAAAA))
|
||||
|
||||
#undef DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32
|
||||
|
||||
#if HT_AMD
|
||||
// TEMPLATE_TEST_CASE wasn't working within a macro, so sections were used instead
|
||||
#define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64(suffix, test_t) \
|
||||
TEST_CASE("Unit_hipStreamValue_Wait64_Blocking_" + std::string(suffix)) { \
|
||||
testWait<int64_t, true>(test_t); \
|
||||
SECTION("HostPtr") { testWait<TestParams<uint64_t, PtrType::HostPtr>, true>(test_t); } \
|
||||
SECTION("DevicePtr") { testWait<TestParams<uint64_t, PtrType::DevicePtr>, true>(test_t); } \
|
||||
SECTION("DevicePtrToHost") { \
|
||||
testWait<TestParams<uint64_t, PtrType::DevicePtrToHost>, true>(test_t); \
|
||||
} \
|
||||
SECTION("Signal") { testWait<TestParams<uint64_t, PtrType::Signal>, true>(test_t); } \
|
||||
} \
|
||||
TEST_CASE("Unit_hipStreamValue_Wait64_NonBlocking_" + std::string(suffix)) { \
|
||||
testWait<int64_t, false>(test_t); \
|
||||
SECTION("HostPtr") { testWait<TestParams<uint64_t, PtrType::HostPtr>, false>(test_t); } \
|
||||
SECTION("DevicePtr") { testWait<TestParams<uint64_t, PtrType::DevicePtr>, false>(test_t); } \
|
||||
SECTION("DevicePtrToHost") { \
|
||||
testWait<TestParams<uint64_t, PtrType::DevicePtrToHost>, false>(test_t); \
|
||||
} \
|
||||
SECTION("Signal") { testWait<TestParams<uint64_t, PtrType::Signal>, false>(test_t); } \
|
||||
}
|
||||
#else
|
||||
#define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64(suffix, test_t) \
|
||||
TEST_CASE("Unit_hipStreamValue_Wait64_Blocking_" + std::string(suffix)) { \
|
||||
SECTION("HostPtr") { testWait<TestParams<uint64_t, PtrType::HostPtr>, true>(test_t); } \
|
||||
SECTION("DevicePtr") { testWait<TestParams<uint64_t, PtrType::DevicePtr>, true>(test_t); } \
|
||||
SECTION("DevicePtrToHost") { \
|
||||
testWait<TestParams<uint64_t, PtrType::DevicePtrToHost>, true>(test_t); \
|
||||
} \
|
||||
} \
|
||||
TEST_CASE("Unit_hipStreamValue_Wait64_NonBlocking_" + std::string(suffix)) { \
|
||||
SECTION("HostPtr") { testWait<TestParams<uint64_t, PtrType::HostPtr>, false>(test_t); } \
|
||||
SECTION("DevicePtr") { testWait<TestParams<uint64_t, PtrType::DevicePtr>, false>(test_t); } \
|
||||
SECTION("DevicePtrToHost") { \
|
||||
testWait<TestParams<uint64_t, PtrType::DevicePtrToHost>, false>(test_t); \
|
||||
} \
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
// Using Mask
|
||||
@@ -332,14 +461,6 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Eq_2",
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_And",
|
||||
TEST_WAIT64( // mask will discard bits 8 to 11
|
||||
hipStreamWaitValueAnd, 0xFF, 0xF4A, 0xF35, 0X02))
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Nor_1",
|
||||
TEST_WAIT64( // mask is set to ignore the sign bit.
|
||||
hipStreamWaitValueNor, 0x7FFFFFFFFFFFFFFF,
|
||||
0x7FFFFFFFFFFFF247, 0x7FFFFFFFFFFFFdbd,
|
||||
0x7FFFFFFFFFFFFdb5))
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Nor_2",
|
||||
TEST_WAIT64( // mask is set to apply NOR for bits 0 to 3.
|
||||
hipStreamWaitValueNor, 0xF, 0x7E, 0x7D, 0x76))
|
||||
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_Gte",
|
||||
TEST_WAIT64(hipStreamWaitValueGte, 0x7FFFFFFFFFFF0001,
|
||||
@@ -352,94 +473,72 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_And",
|
||||
0x0F0F0F0F0F0F0F0F, 0X1F0F0F0F0F0F0F0F))
|
||||
DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_Nor",
|
||||
TEST_WAIT64(hipStreamWaitValueNor, 0x4724724747247247,
|
||||
static_cast<int64_t>(0xbddbddbdbddbddbd),
|
||||
static_cast<int64_t>(0xbddbddbdbddbddb3)))
|
||||
0xbddbddbdbddbddbd, 0xbddbddbdbddbddb3))
|
||||
#undef DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64
|
||||
|
||||
#endif
|
||||
|
||||
// Negative Tests
|
||||
TEST_CASE("Unit_hipStreamValue_Negative_InvalidMemory") {
|
||||
|
||||
#if HT_AMD
|
||||
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-96");
|
||||
return;
|
||||
#endif
|
||||
if (!streamWaitValueSupported()) {
|
||||
HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device.");
|
||||
return;
|
||||
}
|
||||
|
||||
hipStream_t stream{nullptr};
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
REQUIRE(stream != nullptr);
|
||||
|
||||
// Allocate Host Memory
|
||||
auto hostPtr32 = std::unique_ptr<uint32_t>(new uint32_t(1));
|
||||
auto hostPtr64 = std::unique_ptr<uint64_t>(new uint64_t(1));
|
||||
|
||||
// Register Host Memory
|
||||
HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0));
|
||||
HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0));
|
||||
|
||||
// Set dummy data
|
||||
*hostPtr64 = 0x0;
|
||||
*hostPtr32 = 0x0;
|
||||
|
||||
auto compareOp = hipStreamWaitValueGte;
|
||||
const auto compareOp = hipStreamWaitValueGte;
|
||||
const auto expectedError = hipErrorInvalidValue;
|
||||
|
||||
// Memory pointer negative tests
|
||||
|
||||
INFO("Testing Invalid Memory Pointer for hipStreamWriteValue32");
|
||||
NEG_TEST_ERROR_CHECK(Write, 32, hipErrorNotSupported, stream, nullptr, 0, writeFlag)
|
||||
|
||||
INFO("Testing Invalid Memory Pointer for hipStreamWriteValue64");
|
||||
NEG_TEST_ERROR_CHECK(Write, 64, hipErrorNotSupported, stream, nullptr, 0, writeFlag)
|
||||
|
||||
INFO("Testing Invalid Memory Pointer for hipStreamWaitValue32");
|
||||
NEG_TEST_ERROR_CHECK(Wait, 32, hipErrorNotSupported, stream, nullptr, 0, compareOp)
|
||||
|
||||
INFO("Testing Invalid Memory Pointer for hipStreamWaitValue64");
|
||||
NEG_TEST_ERROR_CHECK(Wait, 64, hipErrorNotSupported, stream, nullptr, 0, compareOp)
|
||||
SECTION("Invalid Memory Pointer for hipStreamWriteValue32") {
|
||||
HIP_CHECK_ERROR(hipStreamWriteValue32(stream, nullptr, 0, writeFlag), expectedError);
|
||||
}
|
||||
SECTION("Invalid Memory Pointer for hipStreamWriteValue64") {
|
||||
HIP_CHECK_ERROR(hipStreamWriteValue64(stream, nullptr, 0, writeFlag), expectedError);
|
||||
}
|
||||
SECTION("Invalid Memory Pointer for hipStreamWaitValue32") {
|
||||
HIP_CHECK_ERROR(hipStreamWaitValue32(stream, nullptr, 0, compareOp), expectedError);
|
||||
}
|
||||
SECTION("Invalid Memory Pointer for hipStreamWaitValue32") {
|
||||
HIP_CHECK_ERROR(hipStreamWaitValue64(stream, nullptr, 0, compareOp), expectedError);
|
||||
}
|
||||
|
||||
// Cleanup
|
||||
HIP_CHECK(hipHostUnregister(hostPtr32.get()));
|
||||
HIP_CHECK(hipHostUnregister(hostPtr64.get()));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipStreamWaitValue_Negative_InvalidFlag") {
|
||||
TEMPLATE_TEST_CASE("Unit_hipStreamWaitValue_Negative_InvalidFlag", "", uint32_t, uint64_t) {
|
||||
#if HT_AMD
|
||||
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-96");
|
||||
return;
|
||||
#endif
|
||||
if (!streamWaitValueSupported()) {
|
||||
HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device.");
|
||||
return;
|
||||
}
|
||||
|
||||
hipStream_t stream{nullptr};
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
REQUIRE(stream != nullptr);
|
||||
|
||||
// Allocate Host Memory
|
||||
auto hostPtr32 = std::unique_ptr<uint32_t>(new uint32_t(1));
|
||||
auto hostPtr64 = std::unique_ptr<uint64_t>(new uint64_t(1));
|
||||
auto hostPtr = std::make_unique<TestType>();
|
||||
|
||||
// Register Host Memory
|
||||
HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0));
|
||||
HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0));
|
||||
HIP_CHECK(hipHostRegister(hostPtr.get(), sizeof(TestType), 0));
|
||||
|
||||
// Set dummy data
|
||||
*hostPtr64 = 0x0;
|
||||
*hostPtr32 = 0x0;
|
||||
*hostPtr = 0x0;
|
||||
|
||||
/* EXSWCPHIPT-96 */
|
||||
INFO("Testing Invalid flag for hipStreamWaitValue32");
|
||||
NEG_TEST_ERROR_CHECK(Wait, 32, hipErrorNotSupported, stream, hostPtr32.get(), 0, -1)
|
||||
INFO("Testing Invalid flag for hipStreamWaitValue64");
|
||||
NEG_TEST_ERROR_CHECK(Wait, 64, hipErrorNotSupported, stream, hostPtr64.get(), 0, -1)
|
||||
HIP_CHECK_ERROR(waitFunc<TestType>(stream, hostPtr.get(), 0, -1), hipErrorInvalidValue);
|
||||
|
||||
// Cleanup
|
||||
HIP_CHECK(hipHostUnregister(hostPtr32.get()));
|
||||
HIP_CHECK(hipHostUnregister(hostPtr64.get()));
|
||||
HIP_CHECK(hipHostUnregister(hostPtr.get()));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
#undef NEG_TEST_ERROR_CHECK
|
||||
|
||||
@@ -149,8 +149,8 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_ValidateCallbackFunc") {
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, 0));
|
||||
createDefaultCUMask(&defaultCUMask, props.multiProcessorCount);
|
||||
|
||||
hipExtStreamCreateWithCUMask(&mystream, defaultCUMask.size(),
|
||||
defaultCUMask.data());
|
||||
HIP_CHECK(hipExtStreamCreateWithCUMask(&mystream, defaultCUMask.size(),
|
||||
defaultCUMask.data()));
|
||||
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice,
|
||||
mystream));
|
||||
const unsigned blocks = GRIDSIZE;
|
||||
@@ -244,7 +244,7 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") {
|
||||
|
||||
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, streams[0], dA[0], dC[0], N);
|
||||
hipDeviceSynchronize();
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto single_end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> single_kernel_time = single_end - single_start;
|
||||
@@ -264,7 +264,7 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") {
|
||||
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, streams[np], dA[np], dC[np], N);
|
||||
}
|
||||
hipDeviceSynchronize();
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto all_end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> all_kernel_time = all_end - all_start;
|
||||
@@ -288,8 +288,8 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") {
|
||||
delete [] hA;
|
||||
delete [] hC;
|
||||
for (int np = 0; np < KNumPartition; np++) {
|
||||
hipFree(dC[np]);
|
||||
hipFree(dA[np]);
|
||||
HIP_CHECK(hipFree(dC[np]));
|
||||
HIP_CHECK(hipFree(dA[np]));
|
||||
HIP_CHECK(hipStreamDestroy(streams[np]));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -66,44 +66,8 @@ bool checkStreamFlags_(hipStream_t stream, bool checkFlags = false, unsigned fla
|
||||
|
||||
inline namespace stream {
|
||||
|
||||
__device__ int defaultSemaphore = 0;
|
||||
|
||||
__global__ void signaling_kernel(int* semaphore) {
|
||||
size_t tid{blockIdx.x * blockDim.x + threadIdx.x};
|
||||
if (tid == 0) {
|
||||
if (semaphore == nullptr) {
|
||||
atomicAdd(&defaultSemaphore, 1);
|
||||
} else {
|
||||
atomicAdd(semaphore, 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void waiting_kernel(int* semaphore) {
|
||||
size_t tid{blockIdx.x * blockDim.x + threadIdx.x};
|
||||
if (tid == 0) {
|
||||
if (semaphore == nullptr) {
|
||||
while (atomicCAS(&defaultSemaphore, 1, 2) == 0) {
|
||||
}
|
||||
} else {
|
||||
while (atomicCAS(semaphore, 1, 2) == 0) {
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
std::thread startSignalingThread(int* semaphore) {
|
||||
std::thread signalingThread([semaphore]() {
|
||||
hipStream_t signalingStream;
|
||||
HIP_CHECK_THREAD(hipStreamCreateWithFlags(&signalingStream, hipStreamNonBlocking));
|
||||
|
||||
signaling_kernel<<<1, 1, 0, signalingStream>>>(semaphore);
|
||||
HIP_CHECK_THREAD(hipStreamSynchronize(signalingStream));
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(signalingStream));
|
||||
});
|
||||
|
||||
return signalingThread;
|
||||
}
|
||||
/* Empty kernel to ensure work finishes on the stream quickly */
|
||||
__global__ void empty_kernel() {}
|
||||
|
||||
bool checkStream(hipStream_t stream) {
|
||||
{ // Check default flags
|
||||
|
||||
@@ -24,33 +24,12 @@ THE SOFTWARE.
|
||||
namespace hip {
|
||||
inline namespace stream {
|
||||
|
||||
/* Empty kernel to ensure work finishes on the stream quickly */
|
||||
__global__ void empty_kernel();
|
||||
|
||||
const hipStream_t nullStream = nullptr;
|
||||
const hipStream_t streamPerThread = hipStreamPerThread;
|
||||
|
||||
/**
|
||||
* @brief Kernel that signals a semaphore to change value from 0 to 1.
|
||||
*
|
||||
* @param semaphore the semaphore that needs to be signaled.
|
||||
*/
|
||||
__global__ void signaling_kernel(int* semaphore = nullptr);
|
||||
|
||||
/**
|
||||
* @brief Kernel that busy waits until the specified semaphore goes from 0 to 1.
|
||||
*
|
||||
* @param semaphore the semaphore to wait for.
|
||||
*/
|
||||
__global__ void waiting_kernel(int* semaphore = nullptr);
|
||||
|
||||
/**
|
||||
* @brief Creates a thread that runs a signaling_kernel on a non-blocking stream.
|
||||
* hipStreamNonBlocking is used here to avoid interfering with tests for the Null Stream.
|
||||
* You must call HIP_CHECK_THREAD_FINALIZE after joining this thread.
|
||||
*
|
||||
* @param semaphore memory location to signal
|
||||
* @return std::thread thread that has to be joined after the testing is done.
|
||||
*/
|
||||
std::thread startSignalingThread(int* semaphore = nullptr);
|
||||
|
||||
// Checks stream for valid values of flags and priority
|
||||
bool checkStream(hipStream_t stream);
|
||||
|
||||
|
||||
Ссылка в новой задаче
Block a user