SWDEV-299127 - Merge 'develop' into 'amd-staging'
Change-Id: I570bec5adc6dd2fa445cf072e2d27e21e548b15d
Šī revīzija ir iekļauta:
Jenkins
@@ -127,7 +127,7 @@ Breakpoint 1, main ()
|
||||
```
|
||||
|
||||
### Other Debugging Tools
|
||||
There are also other debugging tools availble online developers can google and choose the one best suits the debugging requirements.
|
||||
There are also other debugging tools available online developers can google and choose the one best suits the debugging requirements.
|
||||
|
||||
## Debugging HIP Applications
|
||||
|
||||
@@ -192,7 +192,7 @@ AMD_SERIALIZE_COPY, for serializing copies.
|
||||
So HIP runtime can wait for GPU idle before/after any GPU command depending on the environment setting.
|
||||
|
||||
### Making Device visible
|
||||
For system with multiple devices, it's possible to make only certain device(s) visible to HIP via setting environment varible,
|
||||
For system with multiple devices, it's possible to make only certain device(s) visible to HIP via setting environment variable,
|
||||
HIP_VISIBLE_DEVICES, only devices whose index is present in the sequence are visible to HIP.
|
||||
|
||||
For example,
|
||||
@@ -200,7 +200,7 @@ For example,
|
||||
$ HIP_VISIBLE_DEVICES=0,1
|
||||
```
|
||||
|
||||
or in the appliation,
|
||||
or in the application,
|
||||
```
|
||||
if (totalDeviceNum > 2) {
|
||||
setenv("HIP_VISIBLE_DEVICES", "0,1,2", 1);
|
||||
@@ -210,11 +210,11 @@ if (totalDeviceNum > 2) {
|
||||
```
|
||||
|
||||
### Dump code object
|
||||
Developers can dump code object to anylize compiler related issues via setting environment variable,
|
||||
Developers can dump code object to analyze compiler related issues via setting environment variable,
|
||||
GPU_DUMP_CODE_OBJECT
|
||||
|
||||
### HSA related environment variables
|
||||
HSA provides some environment varibles help to analize issues in driver or hardware, for example,
|
||||
HSA provides some environment variables help to analyze issues in driver or hardware, for example,
|
||||
|
||||
HSA_ENABLE_SDMA=0
|
||||
It causes host-to-device and device-to-host copies to use compute shader blit kernels rather than the dedicated DMA copy engines.
|
||||
@@ -225,8 +225,24 @@ HSA_ENABLE_INTERRUPT=0
|
||||
Causes completion signals to be detected with memory-based polling rather than interrupts.
|
||||
This environment variable can be useful to diagnose interrupt storm issues in the driver.
|
||||
|
||||
### Summary of environment variables in HIP
|
||||
|
||||
The following is the summary of the most useful environment variables in HIP.
|
||||
|
||||
| **Environment variable** | **Default value** | **Usage** |
|
||||
| ---------------------------------------------------------------------------------------------------------------| ----------------- | --------- |
|
||||
| AMD_LOG_LEVEL <br><sub> Enable HIP log on different Level. </sub> | 0 | 0: Disable log. <br> 1: Enable log on error level. <br> 2: Enable log on warning and below levels. <br> 0x3: Enable log on information and below levels. <br> 0x4: Decode and display AQL packets. |
|
||||
| AMD_LOG_MASK <br><sub> Enable HIP log on different Level. </sub> | 0x7FFFFFFF | 0x1: Log API calls. <br> 0x02: Kernel and Copy Commands and Barriers. <br> 0x4: Synchronization and waiting for commands to finish. <br> 0x8: Enable log on information and below levels. <br> 0x20: Queue commands and queue contents. <br> 0x40:Signal creation, allocation, pool. <br> 0x80: Locks and thread-safety code. <br> 0x100: Copy debug. <br> 0x200: Detailed copy debug. <br> 0x400: Resource allocation, performance-impacting events. <br> 0x800: Initialization and shutdown. <br> 0x1000: Misc debug, not yet classified. <br> 0x2000: Show raw bytes of AQL packet. <br> 0x4000: Show code creation debug. <br> 0x8000: More detailed command info, including barrier commands. <br> 0x10000: Log message location. <br> 0xFFFFFFFF: Log always even mask flag is zero. |
|
||||
| HIP_VISIBLE_DEVICES <br><sub> Only devices whose index is present in the sequence are visible to HIP. </sub> | | 0,1,2: Depending on the number of devices on the system. |
|
||||
| GPU_DUMP_CODE_OBJECT <br><sub> Dump code object. </sub> | 0 | 0: Disable. <br> 1: Enable. |
|
||||
| AMD_SERIALIZE_KERNEL <br><sub> Serialize kernel enqueue. </sub> | 0 | 1: Wait for completion before enqueue. <br> 2: Wait for completion after enqueue. <br> 3: Both. |
|
||||
| AMD_SERIALIZE_COPY <br><sub> Serialize copies. </sub> | 0 | 1: Wait for completion before enqueue. <br> 2: Wait for completion after enqueue. <br> 3: Both. |
|
||||
| HIP_HOST_COHERENT <br><sub> Coherent memory in hipHostMalloc. </sub> | 0 | 0: memory is not coherent between host and GPU. <br> 1: memory is coherent with host. |
|
||||
| AMD_DIRECT_DISPATCH <br><sub> Enable direct kernel dispatch. </sub> | 0 | 0: Disable. <br> 1: Enable. |
|
||||
|
||||
|
||||
## General Debugging Tips
|
||||
- 'gdb --args' can be used to conviently pass the executable and arguments to gdb.
|
||||
- 'gdb --args' can be used to conveniently pass the executable and arguments to gdb.
|
||||
- From inside GDB, you can set environment variables "set env". Note the command does not use an '=' sign:
|
||||
|
||||
```
|
||||
|
||||
@@ -66,12 +66,15 @@ ROCm defines two coherency options for host memory:
|
||||
- Non-coherent memory : Can be cached by GPU, but cannot support synchronization while the kernel is running. Non-coherent memory can be optionally synchronized only at command (end-of-kernel or copy command) boundaries. This memory is appropriate for high-performance access when fine-grain synchronization is not required.
|
||||
|
||||
HIP provides the developer with controls to select which type of memory is used via allocation flags passed to hipHostMalloc and the HIP_HOST_COHERENT environment variable. By default, the environment variable HIP_HOST_COHERENT is set to 0 in HIP.
|
||||
- hipHostMallocCoherent=0, hipHostMallocNonCoherent=0: Use HIP_HOST_COHERENT environment variable,
|
||||
The control logic in the current version of HIP is as follows:
|
||||
- No flags are passed in: the host memory allocation is coherent, the HIP_HOST_COHERENT environment variable is ignored.
|
||||
- hipHostMallocCoherent=1: The host memory allocation will be coherent, the HIP_HOST_COHERENT environment variable is ignored.
|
||||
- hipHostMallocMapped=1: The host memory allocation will be coherent, the HIP_HOST_COHERENT environment variable is ignored.
|
||||
- hipHostMallocNonCoherent=1, hipHostMallocCoherent=0, and hipHostMallocMapped=0: The host memory will be non-coherent, the HIP_HOST_COHERENT environment variable is ignored.
|
||||
- hipHostMallocCoherent=0, hipHostMallocNonCoherent=0, hipHostMallocMapped=0, but one of the other HostMalloc flags is set:
|
||||
- If HIP_HOST_COHERENT is defined as 1, the host memory allocation is coherent.
|
||||
- If HIP_HOST_COHERENT is not defined, or defined as 0, the host memory allocation is non-coherent.
|
||||
- hipHostMallocCoherent=1, hipHostMallocNonCoherent=0: The host memory allocation will be coherent. HIP_HOST_COHERENT env variable is ignored.
|
||||
- hipHostMallocCoherent=0, hipHostMallocNonCoherent=1: The host memory allocation will be non-coherent. HIP_HOST_COHERENT env variable is ignored.
|
||||
- hipHostMallocCoherent=1, hipHostMallocNonCoherent=1: Illegal.
|
||||
- hipHostMallocCoherent=1, hipHostMallocNonCoherent=1: Illegal.
|
||||
|
||||
### Visibility of Zero-Copy Host Memory
|
||||
Coherent host memory is automatically visible at synchronization points.
|
||||
|
||||
@@ -69,6 +69,7 @@ add_subdirectory(unit)
|
||||
add_subdirectory(ABM)
|
||||
add_subdirectory(hipTestMain)
|
||||
add_subdirectory(stress)
|
||||
add_subdirectory(TypeQualifiers)
|
||||
|
||||
if(UNIX)
|
||||
add_subdirectory(multiproc)
|
||||
|
||||
@@ -0,0 +1,12 @@
|
||||
# Common Tests
|
||||
set(TEST_SRC
|
||||
hipManagedKeyword.cc
|
||||
)
|
||||
|
||||
|
||||
# Create shared lib of all tests
|
||||
add_library(TypeQualifiers SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests TypeQualifiers)
|
||||
|
||||
@@ -69,3 +69,14 @@ target_link_libraries(StressTest PRIVATE memory stdc++fs)
|
||||
add_dependencies(build_stress_test StressTest)
|
||||
add_custom_target(stress_test COMMAND StressTest)
|
||||
|
||||
# Space Specifiers/Qualifiers exe
|
||||
add_executable(TypeQualifierTests EXCLUDE_FROM_ALL main.cc hip_test_context.cc)
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set_property(TARGET TypeQualifierTests PROPERTY CXX_STANDARD 17)
|
||||
else()
|
||||
target_compile_options(TypeQualifierTests PUBLIC -std=c++17)
|
||||
endif()
|
||||
target_link_libraries(TypeQualifierTests PRIVATE TypeQualifiers stdc++fs)
|
||||
|
||||
catch_discover_tests(TypeQualifierTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
|
||||
add_dependencies(build_tests TypeQualifierTests)
|
||||
|
||||
@@ -23,10 +23,46 @@ THE SOFTWARE.
|
||||
#pragma once
|
||||
#include "hip_test_common.hh"
|
||||
|
||||
#ifdef __linux__
|
||||
#include <sys/sysinfo.h>
|
||||
#endif
|
||||
|
||||
namespace HipTest {
|
||||
static inline int getGeviceCount() {
|
||||
int dev = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&dev));
|
||||
HIP_CHECK(hipGetDeviceCount(&dev));
|
||||
return dev;
|
||||
}
|
||||
|
||||
// Get Free Memory from the system
|
||||
static size_t getMemoryAmount() {
|
||||
#ifdef __linux__
|
||||
struct sysinfo info{};
|
||||
sysinfo(&info);
|
||||
return info.freeram / (1024 * 1024); // MB
|
||||
#elif defined(_WIN32)
|
||||
MEMORYSTATUSEX statex;
|
||||
statex.dwLength = sizeof(statex);
|
||||
GlobalMemoryStatusEx(&statex);
|
||||
return (statex.ullAvailPhys / (1024 * 1024)); // MB
|
||||
#endif
|
||||
}
|
||||
|
||||
static size_t getHostThreadCount(const size_t memPerThread,
|
||||
const size_t maxThreads) {
|
||||
if (memPerThread == 0) return 0;
|
||||
auto memAmount = getMemoryAmount();
|
||||
const auto processor_count = std::thread::hardware_concurrency();
|
||||
if (processor_count == 0 || memAmount == 0) return 0;
|
||||
size_t thread_count = 0;
|
||||
if ((processor_count * memPerThread) < memAmount)
|
||||
thread_count = processor_count;
|
||||
else
|
||||
thread_count = reinterpret_cast<size_t>(memAmount / memPerThread);
|
||||
if (maxThreads > 0) {
|
||||
return (thread_count > maxThreads) ? maxThreads : thread_count;
|
||||
}
|
||||
return thread_count;
|
||||
}
|
||||
|
||||
} // namespace HipTest
|
||||
|
||||
@@ -0,0 +1,88 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
/**
|
||||
* @brief Error codes retured by rocm_smi_lib functions
|
||||
*/
|
||||
typedef enum {
|
||||
RSMI_STATUS_SUCCESS = 0x0, //!< Operation was successful
|
||||
RSMI_STATUS_INVALID_ARGS, //!< Passed in arguments are not valid
|
||||
RSMI_STATUS_NOT_SUPPORTED, //!< The requested information or
|
||||
//!< action is not available for the
|
||||
//!< given input, on the given system
|
||||
RSMI_STATUS_FILE_ERROR, //!< Problem accessing a file. This
|
||||
//!< may because the operation is not
|
||||
//!< supported by the Linux kernel
|
||||
//!< version running on the executing
|
||||
//!< machine
|
||||
RSMI_STATUS_PERMISSION, //!< Permission denied/EACCESS file
|
||||
//!< error. Many functions require
|
||||
//!< root access to run.
|
||||
RSMI_STATUS_OUT_OF_RESOURCES, //!< Unable to acquire memory or other
|
||||
//!< resource
|
||||
RSMI_STATUS_INTERNAL_EXCEPTION, //!< An internal exception was caught
|
||||
RSMI_STATUS_INPUT_OUT_OF_BOUNDS, //!< The provided input is out of
|
||||
//!< allowable or safe range
|
||||
RSMI_STATUS_INIT_ERROR, //!< An error occurred when rsmi
|
||||
//!< initializing internal data
|
||||
//!< structures
|
||||
RSMI_INITIALIZATION_ERROR = RSMI_STATUS_INIT_ERROR,
|
||||
RSMI_STATUS_NOT_YET_IMPLEMENTED, //!< The requested function has not
|
||||
//!< yet been implemented in the
|
||||
//!< current system for the current
|
||||
//!< devices
|
||||
RSMI_STATUS_NOT_FOUND, //!< An item was searched for but not
|
||||
//!< found
|
||||
RSMI_STATUS_INSUFFICIENT_SIZE, //!< Not enough resources were
|
||||
//!< available for the operation
|
||||
RSMI_STATUS_INTERRUPT, //!< An interrupt occurred during
|
||||
//!< execution of function
|
||||
RSMI_STATUS_UNEXPECTED_SIZE, //!< An unexpected amount of data
|
||||
//!< was read
|
||||
RSMI_STATUS_NO_DATA, //!< No data was found for a given
|
||||
//!< input
|
||||
RSMI_STATUS_UNEXPECTED_DATA, //!< The data read or provided to
|
||||
//!< function is not what was expected
|
||||
RSMI_STATUS_BUSY, //!< A resource or mutex could not be
|
||||
//!< acquired because it is already
|
||||
//!< being used
|
||||
RSMI_STATUS_REFCOUNT_OVERFLOW, //!< An internal reference counter
|
||||
//!< exceeded INT32_MAX
|
||||
|
||||
RSMI_STATUS_UNKNOWN_ERROR = 0xFFFFFFFF, //!< An unknown error occurred
|
||||
} rsmi_status_t;
|
||||
|
||||
|
||||
/**
|
||||
* @brief Types of memory
|
||||
*/
|
||||
typedef enum {
|
||||
RSMI_MEM_TYPE_FIRST = 0,
|
||||
|
||||
RSMI_MEM_TYPE_VRAM = RSMI_MEM_TYPE_FIRST, //!< VRAM memory
|
||||
RSMI_MEM_TYPE_VIS_VRAM, //!< VRAM memory that is visible
|
||||
RSMI_MEM_TYPE_GTT, //!< GTT memory
|
||||
|
||||
RSMI_MEM_TYPE_LAST = RSMI_MEM_TYPE_GTT
|
||||
} rsmi_memory_type_t;
|
||||
@@ -1,6 +1,5 @@
|
||||
# AMD Tests
|
||||
# Common Tests
|
||||
set(LINUX_TEST_SRC
|
||||
hipMallocConcurrency.cc
|
||||
childMalloc.cc
|
||||
hipDeviceComputeCapabilityMproc.cc
|
||||
hipDeviceGetPCIBusIdMproc.cc
|
||||
@@ -10,12 +9,16 @@ set(LINUX_TEST_SRC
|
||||
hipGetDevicePropertiesMproc.cc
|
||||
hipSetGetDeviceMproc.cc
|
||||
hipIpcMemAccessTest.cc
|
||||
hipHostMallocTestsMproc.cc
|
||||
hipMallocConcurrencyMproc.cc
|
||||
)
|
||||
|
||||
if(UNIX)
|
||||
# Create shared lib of all tests
|
||||
add_library(MultiProc SHARED EXCLUDE_FROM_ALL ${LINUX_TEST_SRC})
|
||||
|
||||
target_link_libraries(MultiProc ${CMAKE_DL_LIBS})
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests MultiProc)
|
||||
endif()
|
||||
|
||||
@@ -0,0 +1,314 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Validate page table allocation of hipHostMalloc() api when
|
||||
HIP_VISIBLE_DEVICES set to single device.
|
||||
2) Validate page table allocation of hipHostMalloc() api when
|
||||
HIP_VISIBLE_DEVICES set to list of multiple devices.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_smi.hh>
|
||||
|
||||
#if defined(__linux__)
|
||||
#include <unistd.h>
|
||||
#include <stdlib.h>
|
||||
#include <sys/wait.h>
|
||||
#include <dlfcn.h>
|
||||
|
||||
#if HT_AMD
|
||||
/* AMD specific as below tests use rocm-smi */
|
||||
|
||||
/**
|
||||
* Defines
|
||||
*/
|
||||
#define LIB_ROCMSMI "librocm_smi64.so"
|
||||
#define ALLOC_SIZE (1*1024*1024)
|
||||
|
||||
/**
|
||||
* Global variables
|
||||
*/
|
||||
static rsmi_status_t (*rsmi_dev_memory_usage_get_fp)(uint32_t,
|
||||
rsmi_memory_type_t,
|
||||
uint64_t *);
|
||||
static rsmi_status_t (*rsmi_init_fp)(uint64_t);
|
||||
static rsmi_status_t (*rsmi_shut_down_fp)();
|
||||
static void *rocm_smi_h;
|
||||
|
||||
|
||||
/**
|
||||
* Fetches Gpu device count
|
||||
*/
|
||||
static void getDeviceCount(int *pdevCnt) {
|
||||
int fd[2], val = 0;
|
||||
pid_t childpid;
|
||||
|
||||
// create pipe descriptors
|
||||
pipe(fd);
|
||||
|
||||
// disable visible_devices env from shell
|
||||
unsetenv("ROCR_VISIBLE_DEVICES");
|
||||
unsetenv("HIP_VISIBLE_DEVICES");
|
||||
|
||||
childpid = fork();
|
||||
|
||||
if (childpid > 0) { // Parent
|
||||
close(fd[1]);
|
||||
// parent will wait to read the device cnt
|
||||
read(fd[0], &val, sizeof(val));
|
||||
|
||||
// close the read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
// wait for child exit
|
||||
wait(nullptr);
|
||||
|
||||
*pdevCnt = val;
|
||||
} else if (!childpid) { // Child
|
||||
int devCnt = 1;
|
||||
// writing only, no need for read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
// send the value on the write-descriptor:
|
||||
write(fd[1], &devCnt, sizeof(devCnt));
|
||||
|
||||
// close the write descriptor:
|
||||
close(fd[1]);
|
||||
exit(0);
|
||||
} else { // failure
|
||||
*pdevCnt = 0;
|
||||
return;
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Initializes rocm smi library handles
|
||||
*/
|
||||
static bool rocm_smi_init() {
|
||||
// Open ROCm SMI Library
|
||||
if (!(rocm_smi_h = dlopen(LIB_ROCMSMI, RTLD_LAZY))) {
|
||||
printf("Error opening rocm smi library!\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
void* fnsym = dlsym(rocm_smi_h, "rsmi_dev_memory_usage_get");
|
||||
if (!fnsym) {
|
||||
printf("Error getting rsmi_dev_memory_usage_get() function\n");
|
||||
dlclose(rocm_smi_h);
|
||||
return false;
|
||||
}
|
||||
rsmi_dev_memory_usage_get_fp = reinterpret_cast<rsmi_status_t (*)(uint32_t,
|
||||
rsmi_memory_type_t, uint64_t *)>(fnsym);
|
||||
|
||||
fnsym = dlsym(rocm_smi_h, "rsmi_init");
|
||||
if (!fnsym) {
|
||||
printf("Error getting rsmi_init() function\n");
|
||||
dlclose(rocm_smi_h);
|
||||
return false;
|
||||
}
|
||||
rsmi_init_fp = reinterpret_cast<rsmi_status_t (*)(uint64_t)>(fnsym);
|
||||
|
||||
fnsym = dlsym(rocm_smi_h, "rsmi_shut_down");
|
||||
if (!fnsym) {
|
||||
printf("Error getting rsmi_shut_down() function\n");
|
||||
dlclose(rocm_smi_h);
|
||||
return false;
|
||||
}
|
||||
rsmi_shut_down_fp = reinterpret_cast<rsmi_status_t (*)()>(fnsym);
|
||||
|
||||
uint64_t init_flags = 0;
|
||||
rsmi_status_t retsmi_init;
|
||||
retsmi_init = rsmi_init_fp(init_flags);
|
||||
if (RSMI_STATUS_SUCCESS != retsmi_init) {
|
||||
printf("Error when initializing rocm_smi\n");
|
||||
dlclose(rocm_smi_h);
|
||||
return false;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* Exits rocm smi library
|
||||
*/
|
||||
static void rocm_smi_exit() {
|
||||
rsmi_shut_down_fp();
|
||||
dlclose(rocm_smi_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates page table memory allocations
|
||||
* by setting visible devices selected.
|
||||
*/
|
||||
static bool validatePageTableAllocations(const char *devList, int visDevCnt) {
|
||||
int fd[2];
|
||||
bool testResult = false;
|
||||
pid_t pid;
|
||||
int numdev = 0;
|
||||
|
||||
getDeviceCount(&numdev);
|
||||
REQUIRE(numdev > 0);
|
||||
|
||||
if (pipe(fd) < 0) {
|
||||
printf("Pipe system call failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
pid = fork();
|
||||
|
||||
if (!pid) { // Child process
|
||||
rsmi_status_t ret;
|
||||
std::vector<int> prev, current;
|
||||
uint64_t used = 0;
|
||||
int tmpdev = 0, changeCnt = 0, indx = 0;
|
||||
char *ptr = nullptr;
|
||||
bool testPassed = true;
|
||||
|
||||
// Disable visible_devices env from shell
|
||||
unsetenv("ROCR_VISIBLE_DEVICES");
|
||||
unsetenv("HIP_VISIBLE_DEVICES");
|
||||
|
||||
setenv("HIP_VISIBLE_DEVICES", devList, 1);
|
||||
|
||||
// First Call to initialize hip api
|
||||
hipGetDeviceCount(&tmpdev);
|
||||
|
||||
|
||||
// Get memory snapshot before hostmalloc
|
||||
for (indx = 0; indx < numdev; indx++) {
|
||||
ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used);
|
||||
if (RSMI_STATUS_SUCCESS != ret) {
|
||||
printf("Error while running rsmi_dev_memory_usage_get func\n");
|
||||
dlclose(rocm_smi_h);
|
||||
rsmi_shut_down_fp();
|
||||
return false;
|
||||
}
|
||||
prev.push_back(used);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipHostMalloc(&ptr, ALLOC_SIZE));
|
||||
|
||||
// Get memory snapshot after hostmalloc
|
||||
for (indx = 0; indx < numdev; indx++) {
|
||||
ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used);
|
||||
if (RSMI_STATUS_SUCCESS != ret) {
|
||||
printf("Error while running rsmi_dev_memory_usage_get func\n");
|
||||
dlclose(rocm_smi_h);
|
||||
rsmi_shut_down_fp();
|
||||
hipHostFree(ptr);
|
||||
return false;
|
||||
}
|
||||
current.push_back(used);
|
||||
}
|
||||
|
||||
for (indx = 0; indx < numdev; indx++) {
|
||||
if (current[indx] - prev[indx])
|
||||
changeCnt++;
|
||||
}
|
||||
|
||||
// Check if memory allocation happened only for visible devices
|
||||
if (changeCnt == visDevCnt) {
|
||||
testPassed = true;
|
||||
} else {
|
||||
testPassed = false;
|
||||
}
|
||||
|
||||
hipHostFree(ptr);
|
||||
|
||||
// writing only, no need for read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
// send the value on the write-descriptor:
|
||||
write(fd[1], &testPassed, sizeof(testPassed));
|
||||
|
||||
// close the write descriptor:
|
||||
close(fd[1]);
|
||||
exit(0);
|
||||
} else if (pid > 0) { // parent
|
||||
close(fd[1]);
|
||||
read(fd[0], &testResult, sizeof(testResult));
|
||||
close(fd[0]);
|
||||
wait(nullptr);
|
||||
} else {
|
||||
printf("fork() failed\n");
|
||||
HIP_ASSERT(false);
|
||||
}
|
||||
|
||||
return testResult;
|
||||
}
|
||||
|
||||
/**
|
||||
* Test page table allocation when HIP_VISIBLE_DEVICES set to
|
||||
* single device
|
||||
*/
|
||||
TEST_CASE("Unit_hipHostMalloc_SingleVisibleDevicePageAlloc") {
|
||||
int devCnt;
|
||||
std::string str;
|
||||
|
||||
if (!rocm_smi_init()) {
|
||||
WARN("Testcase skipped as rocm smi not initialized/present");
|
||||
return;
|
||||
}
|
||||
|
||||
getDeviceCount(&devCnt);
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
// Select single visible device and validate memory usage
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
str = std::to_string(i);
|
||||
REQUIRE(validatePageTableAllocations(str.c_str(), 1) == true);
|
||||
}
|
||||
|
||||
rocm_smi_exit();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test page table allocation when HIP_VISIBLE_DEVICES set to
|
||||
* multiple devices
|
||||
*/
|
||||
TEST_CASE("Unit_hipHostMalloc_MultipleVisibleDevicesPageAlloc") {
|
||||
int devCnt = 0, vdCnt = 0;
|
||||
std::string str;
|
||||
|
||||
if (!rocm_smi_init()) {
|
||||
WARN("Testcase skipped as rocm smi not initialized/present");
|
||||
return;
|
||||
}
|
||||
|
||||
getDeviceCount(&devCnt);
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
// Select multiple visible devices and validate memory usage
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
if (i == 0)
|
||||
str += std::to_string(i);
|
||||
else
|
||||
str += "," + std::to_string(i);
|
||||
|
||||
vdCnt++;
|
||||
REQUIRE(validatePageTableAllocations(str.c_str(), vdCnt) == true);
|
||||
}
|
||||
|
||||
rocm_smi_exit();
|
||||
}
|
||||
#endif // HT_AMD
|
||||
#endif // __linux__
|
||||
@@ -1,171 +0,0 @@
|
||||
#include <sys/types.h>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
#ifdef __linux__
|
||||
#include <sys/wait.h>
|
||||
#include <unistd.h>
|
||||
#endif
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <limits>
|
||||
#include <atomic>
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
|
||||
static constexpr size_t N = 4 * 1024 * 1024;
|
||||
static constexpr unsigned blocksPerCU = 6; // to hide latency
|
||||
static constexpr unsigned threadsPerBlock = 256;
|
||||
/**
|
||||
* Validates data consitency on supplied gpu
|
||||
*/
|
||||
bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
size_t prevAvl, prevTot, curAvl, curTot;
|
||||
bool TestPassed = true;
|
||||
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
|
||||
printf("tgs allocating..\n");
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
|
||||
printf("Validation PASSED for gpu %d from pid %d\n", gpu, getpid());
|
||||
} else {
|
||||
printf("%s : Validation FAILED for gpu %d from pid %d\n", __func__, gpu, getpid());
|
||||
TestPassed &= false;
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
|
||||
|
||||
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
|
||||
// In concurrent calls on one GPU, we cannot verify leaking in this way
|
||||
printf(
|
||||
"%s : Memory allocation mismatch observed."
|
||||
"Possible memory leak.\n",
|
||||
__func__);
|
||||
TestPassed &= false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
|
||||
#if 1
|
||||
/**
|
||||
* Fetches Gpu device count
|
||||
*/
|
||||
void getDeviceCount1(int* pdevCnt) {
|
||||
#ifdef __linux__
|
||||
int fd[2], val = 0;
|
||||
pid_t childpid;
|
||||
|
||||
// create pipe descriptors
|
||||
pipe(fd);
|
||||
|
||||
// disable visible_devices env from shell
|
||||
unsetenv("ROCR_VISIBLE_DEVICES");
|
||||
unsetenv("HIP_VISIBLE_DEVICES");
|
||||
|
||||
childpid = fork();
|
||||
|
||||
if (childpid > 0) { // Parent
|
||||
close(fd[1]);
|
||||
// parent will wait to read the device cnt
|
||||
read(fd[0], &val, sizeof(val));
|
||||
|
||||
// close the read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
// wait for child exit
|
||||
wait(NULL);
|
||||
|
||||
*pdevCnt = val;
|
||||
} else if (!childpid) { // Child
|
||||
int devCnt = 1;
|
||||
// writing only, no need for read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
// send the value on the write-descriptor:
|
||||
write(fd[1], &devCnt, sizeof(devCnt));
|
||||
|
||||
// close the write descriptor:
|
||||
close(fd[1]);
|
||||
exit(0);
|
||||
} else { // failure
|
||||
*pdevCnt = 1;
|
||||
return;
|
||||
}
|
||||
|
||||
#else
|
||||
HIP_CHECK(hipGetDeviceCount(pdevCnt));
|
||||
#endif
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
TEST_CASE("hipMallocChild_Concurrency_MultiGpu") {
|
||||
bool TestPassed = false;
|
||||
#ifdef __linux__
|
||||
// Parallel execution on multiple gpus from different child processes
|
||||
int devCnt = 1, pid = 0;
|
||||
|
||||
// Get GPU count
|
||||
getDeviceCount1(&devCnt);
|
||||
|
||||
// Spawn child for each GPU
|
||||
for (int gpu = 0; gpu < devCnt; gpu++) {
|
||||
if ((pid = fork()) < 0) {
|
||||
INFO("Child_Concurrency_MultiGpu : fork() returned error" << pid);
|
||||
REQUIRE(false);
|
||||
|
||||
} else if (!pid) { // Child process
|
||||
bool TestPassedChild = false;
|
||||
TestPassedChild = validateMemoryOnGPU(gpu);
|
||||
|
||||
if (TestPassedChild) {
|
||||
printf("returning exit(1) for success\n");
|
||||
exit(1); // child exit with success status
|
||||
} else {
|
||||
printf("Child_Concurrency_MultiGpu : childpid %d failed\n", getpid());
|
||||
exit(2); // child exit with failure status
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Parent shall wait for child to complete
|
||||
int cnt = 0;
|
||||
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
int pidwait = 0, exitStatus;
|
||||
pidwait = wait(&exitStatus);
|
||||
|
||||
printf("exitStatus for iter:%d is %d\n", i, exitStatus);
|
||||
if (pidwait < 0) {
|
||||
break;
|
||||
}
|
||||
|
||||
if (WEXITSTATUS(exitStatus) == 1) cnt++;
|
||||
}
|
||||
|
||||
if (cnt && (cnt == devCnt)) TestPassed = true;
|
||||
|
||||
#else
|
||||
INFO("Test hipMallocChild_Concurrency_MultiGpu skipped on non-linux");
|
||||
#endif
|
||||
REQUIRE(TestPassed == true);
|
||||
}
|
||||
@@ -0,0 +1,228 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
|
||||
1) Run hipMalloc() api/kernel code on same gpu parallely from parent and child
|
||||
processes, validate the results.
|
||||
|
||||
2) Execute hipMalloc() api simultaneously on all the gpus by spawning multiple
|
||||
child processes. Validate buffers allocated after running kernel code.
|
||||
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
|
||||
#ifdef __linux__
|
||||
#include <sys/wait.h>
|
||||
#include <sys/types.h>
|
||||
#include <unistd.h>
|
||||
|
||||
/**
|
||||
* Fetches Gpu device count
|
||||
*/
|
||||
static void getDeviceCount(int* pdevCnt) {
|
||||
int fd[2], val = 0;
|
||||
pid_t childpid;
|
||||
|
||||
// create pipe descriptors
|
||||
pipe(fd);
|
||||
|
||||
// disable visible_devices env from shell
|
||||
#ifdef HT_NVIDIA
|
||||
unsetenv("CUDA_VISIBLE_DEVICES");
|
||||
#else
|
||||
unsetenv("ROCR_VISIBLE_DEVICES");
|
||||
unsetenv("HIP_VISIBLE_DEVICES");
|
||||
#endif
|
||||
|
||||
childpid = fork();
|
||||
|
||||
if (childpid > 0) { // Parent
|
||||
close(fd[1]);
|
||||
// parent will wait to read the device cnt
|
||||
read(fd[0], &val, sizeof(val));
|
||||
|
||||
// close the read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
// wait for child exit
|
||||
wait(nullptr);
|
||||
|
||||
*pdevCnt = val;
|
||||
} else if (!childpid) { // Child
|
||||
int devCnt = 1;
|
||||
// writing only, no need for read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
// send the value on the write-descriptor:
|
||||
write(fd[1], &devCnt, sizeof(devCnt));
|
||||
|
||||
// close the write descriptor:
|
||||
close(fd[1]);
|
||||
exit(0);
|
||||
} else { // failure
|
||||
*pdevCnt = 0;
|
||||
return;
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates data consistency on supplied gpu
|
||||
*/
|
||||
static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
size_t prevAvl, prevTot, curAvl, curTot;
|
||||
bool TestPassed = true;
|
||||
constexpr auto N = 4 * 1024 * 1024;
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, 0, static_cast<const int*>(A_d),
|
||||
static_cast<const int*>(B_d), C_d, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
|
||||
printf("Validation PASSED for gpu %d from pid %d\n", gpu, getpid());
|
||||
} else {
|
||||
printf("Validation FAILED for gpu %d from pid %d\n", gpu, getpid());
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
|
||||
|
||||
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
|
||||
// In concurrent calls on one GPU, we cannot verify leaking in this way
|
||||
printf(
|
||||
"%s : Memory allocation mismatch observed."
|
||||
"Possible memory leak.\n",
|
||||
__func__);
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Parallel execution of parent and child on gpu0
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc_ChildConcurrencyDefaultGpu") {
|
||||
int devCnt = 0, pid = 0;
|
||||
constexpr auto resSuccess = 1, resFailure = 2;
|
||||
bool TestPassed = true;
|
||||
|
||||
// Get GPU count
|
||||
getDeviceCount(&devCnt);
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
if ((pid = fork()) < 0) {
|
||||
INFO("Child_Concurrency_DefaultGpu : fork() returned error : " << pid);
|
||||
HIP_ASSERT(false);
|
||||
|
||||
} else if (!pid) { // Child process
|
||||
bool TestPassedChild = false;
|
||||
|
||||
// Allocates and validates memory on Gpu0 simultaneously with parent
|
||||
TestPassedChild = validateMemoryOnGPU(0, true);
|
||||
|
||||
if (TestPassedChild) {
|
||||
exit(resSuccess); // child exit with success status
|
||||
} else {
|
||||
exit(resFailure); // child exit with failure status
|
||||
}
|
||||
|
||||
} else { // Parent process
|
||||
int exitStatus;
|
||||
|
||||
// Allocates and validates memory on Gpu0 simultaneously with child
|
||||
TestPassed = validateMemoryOnGPU(0, true);
|
||||
|
||||
// Wait and get result from child
|
||||
pid = wait(&exitStatus);
|
||||
if ((WEXITSTATUS(exitStatus) == resFailure) || (pid < 0))
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
REQUIRE(TestPassed == true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Parallel execution of api on multiple gpus from
|
||||
* different child processes.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc_ChildConcurrencyMultiGpu") {
|
||||
int devCnt = 0, pid = 0;
|
||||
constexpr auto resSuccess = 1, resFailure = 2;
|
||||
|
||||
// Get GPU count
|
||||
getDeviceCount(&devCnt);
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
// Spawn child for each GPU
|
||||
for (int gpu = 0; gpu < devCnt; gpu++) {
|
||||
if ((pid = fork()) < 0) {
|
||||
INFO("Child_Concurrency_MultiGpu : fork() returned error : " << pid);
|
||||
REQUIRE(false);
|
||||
|
||||
} else if (!pid) { // Child process
|
||||
bool TestPassedChild = false;
|
||||
TestPassedChild = validateMemoryOnGPU(gpu);
|
||||
|
||||
if (TestPassedChild) {
|
||||
exit(resSuccess); // child exit with success status
|
||||
} else {
|
||||
exit(resFailure); // child exit with failure status
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Parent shall wait for child to complete
|
||||
int passCnt = 0;
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
int pidwait = 0, exitStatus;
|
||||
pidwait = wait(&exitStatus);
|
||||
|
||||
printf("exitStatus for dev:%d is %d\n", i, WEXITSTATUS(exitStatus));
|
||||
if (pidwait < 0) {
|
||||
break;
|
||||
}
|
||||
|
||||
if (WEXITSTATUS(exitStatus) == resSuccess) passCnt++;
|
||||
}
|
||||
REQUIRE(passCnt == devCnt);
|
||||
}
|
||||
#endif // __linux__
|
||||
@@ -32,7 +32,18 @@ set(TEST_SRC
|
||||
hipHostGetFlags.cc
|
||||
hipMemoryAllocateCoherent.cc
|
||||
hipMallocManaged_MultiScenario.cc
|
||||
hipManagedKeyword.cc
|
||||
hipMemsetInvalidPtr.cc
|
||||
hipMemset.cc
|
||||
hipMemsetAsyncMultiThread.cc
|
||||
hipMemsetAsyncAndKernel.cc
|
||||
hipMemset3D.cc
|
||||
hipMemset2D.cc
|
||||
hipMemset2DAsyncMultiThreadAndKernel.cc
|
||||
hipHostMallocTests.cc
|
||||
hipMallocConcurrency.cc
|
||||
hipMemset3DFunctional.cc
|
||||
hipMemset3DNegative.cc
|
||||
hipMemset3DRegressMultiThread.cc
|
||||
)
|
||||
else()
|
||||
set(TEST_SRC
|
||||
@@ -65,7 +76,18 @@ set(TEST_SRC
|
||||
hipHostGetFlags.cc
|
||||
hipMemoryAllocateCoherent.cc
|
||||
hipMallocManaged_MultiScenario.cc
|
||||
hipManagedKeyword.cc
|
||||
hipMemsetInvalidPtr.cc
|
||||
hipMemset.cc
|
||||
hipMemsetAsyncMultiThread.cc
|
||||
hipMemsetAsyncAndKernel.cc
|
||||
hipMemset3D.cc
|
||||
hipMemset2D.cc
|
||||
hipMemset2DAsyncMultiThreadAndKernel.cc
|
||||
hipHostMallocTests.cc
|
||||
hipMallocConcurrency.cc
|
||||
hipMemset3DFunctional.cc
|
||||
hipMemset3DNegative.cc
|
||||
hipMemset3DRegressMultiThread.cc
|
||||
)
|
||||
endif()
|
||||
# Create shared lib of all tests
|
||||
|
||||
@@ -0,0 +1,61 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
|
||||
1) Test hipHostMalloc() api with ptr as nullptr and check for return value.
|
||||
2) Test hipHostMalloc() api with size as max(size_t) and check for OOM error.
|
||||
3) Test hipHostMalloc() api with flags as max(unsigned int) and validate
|
||||
return value.
|
||||
4) Pass size as zero for hipHostMalloc() api and check ptr is reset with
|
||||
with return value success.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
/**
|
||||
* Performs argument validation of hipHostMalloc api.
|
||||
*/
|
||||
TEST_CASE("Unit_hipHostMalloc_ArgValidation") {
|
||||
hipError_t ret;
|
||||
constexpr size_t allocSize = 1000;
|
||||
char *ptr;
|
||||
|
||||
SECTION("Pass ptr as nullptr") {
|
||||
ret = hipHostMalloc(static_cast<void **>(nullptr), allocSize);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Size as max(size_t)") {
|
||||
ret = hipHostMalloc(&ptr, std::numeric_limits<std::size_t>::max());
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Flags as max(uint)") {
|
||||
ret = hipHostMalloc(&ptr, allocSize,
|
||||
std::numeric_limits<unsigned int>::max());
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Pass size as zero and check ptr reset") {
|
||||
HIP_CHECK(hipHostMalloc(&ptr, 0));
|
||||
REQUIRE(ptr == nullptr);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,412 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
|
||||
1) Test hipMalloc() api passing zero size and confirming *ptr returning
|
||||
nullptr. Also pass nullptr to hipFree() api.
|
||||
|
||||
2) Pass maximum value of size_t for hipMalloc() api and make sure appropriate
|
||||
error is returned.
|
||||
|
||||
3) Check for hipMalloc() error code, passing invalid/null pointer.
|
||||
|
||||
4) Regress hipMalloc()/hipFree() in loop for bigger chunk of allocation
|
||||
with adequate number of iterations and later test for kernel execution on
|
||||
default gpu.
|
||||
|
||||
5) Regress hipMalloc()/hipFree() in loop while allocating smaller chunks
|
||||
keeping maximum number of iterations and then run kernel code on default
|
||||
gpu, perfom data validation.
|
||||
|
||||
6) Check hipMalloc() api adaptability when app creates small chunks of memory
|
||||
continuously, stores it for later use and then frees it at later point
|
||||
of time.
|
||||
|
||||
7) Multithread Scenario : Exercise hipMalloc() api parellely on all gpus from
|
||||
multiple threads and regress the api.
|
||||
|
||||
8) Validate memory usage with hipMemGetInfo() while regressing hipMalloc()
|
||||
api. Check for any possible memory leaks.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
|
||||
|
||||
#include <vector>
|
||||
#include <limits>
|
||||
#include <atomic>
|
||||
|
||||
|
||||
/* Buffer size for bigger chunks in alloc/free cycles */
|
||||
static constexpr auto BuffSizeBC = 5*1024*1024;
|
||||
|
||||
/* Buffer size for smaller chunks in alloc/free cycles */
|
||||
static constexpr auto BuffSizeSC = 16;
|
||||
|
||||
/* You may change it for individual test.
|
||||
* But default 100 is for quick return in Jenkin Build */
|
||||
static constexpr auto NumDiv = 100;
|
||||
|
||||
/* Max alloc/free iterations for smaller chunks */
|
||||
static constexpr auto MaxAllocFree_SmallChunks = (5000000/NumDiv);
|
||||
|
||||
/* Max alloc/free iterations for bigger chunks */
|
||||
static constexpr auto MaxAllocFree_BigChunks = 10000;
|
||||
|
||||
/* Max alloc and pool iterations */
|
||||
static constexpr auto MaxAllocPoolIter = (2000000/NumDiv);
|
||||
|
||||
/* Test status shared across threads */
|
||||
static std::atomic<bool> g_thTestPassed{true};
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* Validates data consistency on supplied gpu
|
||||
*/
|
||||
static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
size_t prevAvl, prevTot, curAvl, curTot;
|
||||
bool TestPassed = true;
|
||||
constexpr auto N = 4 * 1024 * 1024;
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, 0, static_cast<const int*>(A_d),
|
||||
static_cast<const int*>(B_d), C_d, N);
|
||||
|
||||
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
|
||||
UNSCOPED_INFO("Validation PASSED for gpu " << gpu);
|
||||
} else {
|
||||
UNSCOPED_INFO("Validation FAILED for gpu " << gpu);
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
|
||||
|
||||
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
|
||||
// In concurrent calls on one GPU, we cannot verify leaking in this way
|
||||
UNSCOPED_INFO(
|
||||
"validateMemoryOnGPU : Memory allocation mismatch observed."
|
||||
<< "Possible memory leak.");
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Regress memory allocation and free in loop
|
||||
*/
|
||||
static bool regressAllocInLoop(int gpu) {
|
||||
bool TestPassed = true;
|
||||
size_t tot, avail, ptot, pavail, numBytes;
|
||||
int i = 0;
|
||||
int *ptr;
|
||||
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
numBytes = BuffSizeBC;
|
||||
|
||||
// Exercise allocation in loop with bigger chunks
|
||||
for (i = 0; i < MaxAllocFree_BigChunks; i++) {
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
HIP_CHECK(hipMalloc(&ptr, numBytes));
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
HIP_CHECK(hipFree(ptr));
|
||||
|
||||
if (pavail-avail < numBytes) { // We expect pavail-avail >= numBytes
|
||||
UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " <<
|
||||
numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" <<
|
||||
pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" <<
|
||||
tot << ", pavail-avail=" << pavail-avail);
|
||||
TestPassed = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
// Exercise allocation in loop with smaller chunks and maximum iters
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
numBytes = BuffSizeSC;
|
||||
|
||||
for (i = 0; i < MaxAllocFree_SmallChunks; i++) {
|
||||
HIP_CHECK(hipMalloc(&ptr, numBytes));
|
||||
|
||||
HIP_CHECK(hipFree(ptr));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
if ((pavail != avail) || (ptot != tot)) {
|
||||
UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." <<
|
||||
"Possible memory leak.");
|
||||
TestPassed &= false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates data consistency on supplied gpu
|
||||
* In Multithreaded Environment
|
||||
*/
|
||||
static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) {
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
size_t prevAvl, prevTot, curAvl, curTot;
|
||||
bool TestPassed = true;
|
||||
constexpr auto N = 4 * 1024 * 1024;
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
HIPCHECK(hipSetDevice(gpu));
|
||||
HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot));
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, 0, static_cast<const int*>(A_d),
|
||||
static_cast<const int*>(B_d), C_d, N);
|
||||
|
||||
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
|
||||
UNSCOPED_INFO("Validation PASSED for gpu " << gpu);
|
||||
} else {
|
||||
UNSCOPED_INFO("Validation FAILED for gpu " << gpu);
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipMemGetInfo(&curAvl, &curTot));
|
||||
|
||||
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
|
||||
// In concurrent calls on one GPU, we cannot verify leaking in this way
|
||||
UNSCOPED_INFO(
|
||||
"validateMemoryOnGpuMThread : Memory allocation mismatch observed."
|
||||
"Possible memory leak.");
|
||||
TestPassed = false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Regress memory allocation and free in loop
|
||||
* In Multithreaded Environment
|
||||
*/
|
||||
static bool regressAllocInLoopMthread(int gpu) {
|
||||
bool TestPassed = true;
|
||||
size_t tot, avail, ptot, pavail, numBytes;
|
||||
int i = 0;
|
||||
int *ptr;
|
||||
|
||||
HIPCHECK(hipSetDevice(gpu));
|
||||
numBytes = BuffSizeBC;
|
||||
|
||||
// Exercise allocation in loop with bigger chunks
|
||||
for (i = 0; i < MaxAllocFree_BigChunks; i++) {
|
||||
HIPCHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
HIPCHECK(hipMalloc(&ptr, numBytes));
|
||||
HIPCHECK(hipMemGetInfo(&avail, &tot));
|
||||
HIPCHECK(hipFree(ptr));
|
||||
|
||||
if (pavail-avail < numBytes) { // We expect pavail-avail >= numBytes
|
||||
UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " <<
|
||||
numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" <<
|
||||
pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" <<
|
||||
tot << ", pavail-avail=" << pavail-avail);
|
||||
TestPassed = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
// Exercise allocation in loop with smaller chunks and maximum iters
|
||||
HIPCHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
numBytes = BuffSizeSC;
|
||||
|
||||
for (i = 0; i < MaxAllocFree_SmallChunks; i++) {
|
||||
HIPCHECK(hipMalloc(&ptr, numBytes));
|
||||
|
||||
HIPCHECK(hipFree(ptr));
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
if ((pavail != avail) || (ptot != tot)) {
|
||||
UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." <<
|
||||
"Possible memory leak.");
|
||||
TestPassed &= false;
|
||||
}
|
||||
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/*
|
||||
* Thread func to regress alloc and check data consistency
|
||||
*/
|
||||
static void threadFunc(int gpu) {
|
||||
g_thTestPassed = regressAllocInLoopMthread(gpu);
|
||||
g_thTestPassed = g_thTestPassed & validateMemoryOnGpuMThread(gpu);
|
||||
|
||||
UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" <<
|
||||
g_thTestPassed.load());
|
||||
}
|
||||
|
||||
|
||||
/* Performs Argument Validation of api */
|
||||
TEST_CASE("Unit_hipMalloc_ArgumentValidation") {
|
||||
int *ptr;
|
||||
hipError_t ret;
|
||||
|
||||
SECTION("hipMalloc() when size(0)") {
|
||||
HIP_CHECK(hipMalloc(&ptr, 0));
|
||||
// ptr expected to be reset to null ptr
|
||||
REQUIRE(ptr == nullptr);
|
||||
}
|
||||
|
||||
SECTION("hipFree() when freeing nullptr ") {
|
||||
ptr = nullptr;
|
||||
// api should return success and shudnt crash
|
||||
HIP_CHECK(hipFree(ptr));
|
||||
}
|
||||
|
||||
SECTION("hipMalloc() with invalid argument") {
|
||||
constexpr auto sizeBytes = 100;
|
||||
ret = hipMalloc(nullptr, sizeBytes);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("hipMalloc() with max size_t") {
|
||||
ret = hipMalloc(&ptr, std::numeric_limits<std::size_t>::max());
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Regress hipMalloc()/hipFree() in loop for bigger chunks and
|
||||
* smaller chunks of memory allocation
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc_LoopRegressionAllocFreeCycles") {
|
||||
int devCnt = 0;
|
||||
|
||||
// Get GPU count
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
CHECK(regressAllocInLoop(0) == true);
|
||||
CHECK(validateMemoryOnGPU(0) == true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Application Behavior Modelling.
|
||||
* Check hipMalloc() api adaptability when app creates small chunks of memory
|
||||
* continuously, stores it for later use and then frees it at later point
|
||||
* of time.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
|
||||
size_t avail, tot, pavail, ptot;
|
||||
bool ret;
|
||||
hipError_t err;
|
||||
std::vector<int *> ptrlist;
|
||||
constexpr auto BuffSize = 10;
|
||||
int devCnt, *ptr;
|
||||
|
||||
// Get GPU count
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
|
||||
// Allocate small chunks of memory million times
|
||||
for (int i = 0; i < MaxAllocPoolIter ; i++) {
|
||||
if ((err = hipMalloc(&ptr, BuffSize)) != hipSuccess) {
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
INFO("Loop regression pool allocation failure. " <<
|
||||
"Total gpu memory " << tot/(1024.0*1024.0) <<", Free memory " <<
|
||||
avail/(1024.0*1024.0) << " iter " << i << " error "
|
||||
<< hipGetErrorString(err));
|
||||
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
// Store pointers allocated to emulate memory pool of app
|
||||
ptrlist.push_back(ptr);
|
||||
}
|
||||
|
||||
// Free ptrs at later point of time
|
||||
for ( auto &t : ptrlist ) {
|
||||
HIP_CHECK(hipFree(t));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
ret = validateMemoryOnGPU(0);
|
||||
REQUIRE(ret == true);
|
||||
REQUIRE(pavail == avail);
|
||||
REQUIRE(ptot == tot);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Exercise hipMalloc() api parellely on all gpus from
|
||||
* multiple threads and regress the api.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc_Multithreaded_MultiGPU") {
|
||||
std::vector<std::thread> threadlist;
|
||||
int devCnt;
|
||||
|
||||
// Get GPU count
|
||||
HIP_CHECK(hipGetDeviceCount(&devCnt));
|
||||
REQUIRE(devCnt > 0);
|
||||
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
threadlist.push_back(std::thread(threadFunc, i));
|
||||
}
|
||||
|
||||
for (auto &t : threadlist) {
|
||||
t.join();
|
||||
}
|
||||
|
||||
REQUIRE(g_thTestPassed == true);
|
||||
}
|
||||
@@ -0,0 +1,281 @@
|
||||
/*
|
||||
* Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Test hipMemset small size buffers with unique memset values.
|
||||
2) Test hipMemset, hipMemsetD8, hipMemsetD16, hipMemsetD32 apis with unique
|
||||
number of elements and memset values.
|
||||
3) Test hipMemsetAsync, hipMemsetD8Async, hipMemsetD16Async, hipMemsetD32Async
|
||||
apis with unique number of elements and memset values.
|
||||
4) Test two memset async operations at the same time.
|
||||
*/
|
||||
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
|
||||
// Table with unique number of elements and memset values.
|
||||
// (N, memsetval, memsetD32val, memsetD16val, memsetD8val)
|
||||
typedef std::tuple<size_t, char, int, int16_t, char> tupletype;
|
||||
static constexpr std::initializer_list<tupletype> tableItems {
|
||||
std::make_tuple((4*1024*1024), 0x42, 0xDEADBEEF, 0xDEAD, 0xDE),
|
||||
std::make_tuple((10) , 0x42, 0x101 , 0x10, 0x1),
|
||||
std::make_tuple((10013) , 0x5a, 0xDEADBEEF, 0xDEAD, 0xDE),
|
||||
std::make_tuple((256*1024*1024), 0xa6, 0xCAFEBABE, 0xCAFE, 0xCA)
|
||||
};
|
||||
|
||||
enum MemsetType {
|
||||
hipMemsetTypeDefault,
|
||||
hipMemsetTypeD8,
|
||||
hipMemsetTypeD16,
|
||||
hipMemsetTypeD32
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
static bool testhipMemset(T *A_h, T *A_d, T memsetval, enum MemsetType type,
|
||||
size_t numElements) {
|
||||
size_t Nbytes = numElements * sizeof(T);
|
||||
bool testResult = true;
|
||||
constexpr auto MAX_OFFSET = 3; // To memset on unaligned ptr.
|
||||
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
A_h = reinterpret_cast<T*> (malloc(Nbytes));
|
||||
REQUIRE(A_h != nullptr);
|
||||
|
||||
for (int offset = MAX_OFFSET; offset >= 0; offset --) {
|
||||
if (type == hipMemsetTypeDefault) {
|
||||
HIP_CHECK(hipMemset(A_d + offset, memsetval, numElements - offset));
|
||||
|
||||
} else if (type == hipMemsetTypeD8) {
|
||||
HIP_CHECK(hipMemsetD8((hipDeviceptr_t)(A_d + offset), memsetval,
|
||||
numElements - offset));
|
||||
|
||||
} else if (type == hipMemsetTypeD16) {
|
||||
HIP_CHECK(hipMemsetD16((hipDeviceptr_t)(A_d + offset), memsetval,
|
||||
numElements - offset));
|
||||
|
||||
} else if (type == hipMemsetTypeD32) {
|
||||
HIP_CHECK(hipMemsetD32((hipDeviceptr_t)(A_d + offset), memsetval,
|
||||
numElements - offset));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
for (size_t i = offset; i < numElements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
testResult = false;
|
||||
CAPTURE(i, A_h[i], memsetval);
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
free(A_h);
|
||||
return testResult;
|
||||
}
|
||||
|
||||
|
||||
template<typename T>
|
||||
static bool testhipMemsetAsync(T *A_h, T *A_d, T memsetval,
|
||||
enum MemsetType type, size_t numElements) {
|
||||
size_t Nbytes = numElements * sizeof(T);
|
||||
bool testResult = true;
|
||||
constexpr auto MAX_OFFSET = 3; // To memset on unaligned ptr.
|
||||
hipStream_t stream;
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
A_h = reinterpret_cast<T*> (malloc(Nbytes));
|
||||
REQUIRE(A_h != nullptr);
|
||||
|
||||
for (int offset = MAX_OFFSET; offset >= 0; offset --) {
|
||||
if (type == hipMemsetTypeDefault) {
|
||||
HIP_CHECK(hipMemsetAsync(A_d + offset, memsetval, numElements - offset,
|
||||
stream));
|
||||
|
||||
} else if (type == hipMemsetTypeD8) {
|
||||
HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)(A_d + offset), memsetval,
|
||||
numElements - offset, stream));
|
||||
|
||||
} else if (type == hipMemsetTypeD16) {
|
||||
HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)(A_d + offset), memsetval,
|
||||
numElements - offset, stream));
|
||||
|
||||
} else if (type == hipMemsetTypeD32) {
|
||||
HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)(A_d + offset), memsetval,
|
||||
numElements - offset, stream));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
for (size_t i = offset; i < numElements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
testResult = false;
|
||||
CAPTURE(i, A_h[i], memsetval);
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
free(A_h);
|
||||
return testResult;
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Test hipMemset, hipMemsetD8, hipMemsetD16, hipMemsetD32 apis with unique
|
||||
* number of elements and memset values.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset_SetMemoryWithOffset") {
|
||||
char memsetval;
|
||||
int memsetD32val;
|
||||
int16_t memsetD16val;
|
||||
char memsetD8val;
|
||||
size_t N;
|
||||
bool ret;
|
||||
|
||||
std::tie(N, memsetval, memsetD32val, memsetD16val, memsetD8val) =
|
||||
GENERATE(table<size_t, char, int, int16_t, char>(tableItems));
|
||||
|
||||
|
||||
SECTION("Memset with hipMemsetTypeDefault") {
|
||||
char *cA_d{nullptr}, *cA_h{nullptr};
|
||||
ret = testhipMemset(cA_h, cA_d, memsetval, hipMemsetTypeDefault, N);
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("Memset with hipMemsetTypeD32") {
|
||||
int32_t *iA_d{nullptr}, *iA_h{nullptr};
|
||||
ret = testhipMemset(iA_h, iA_d, memsetD32val, hipMemsetTypeD32, N);
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("Memset with hipMemsetTypeD16") {
|
||||
int16_t *siA_d{nullptr}, *siA_h{nullptr};
|
||||
ret = testhipMemset(siA_h, siA_d, memsetD16val, hipMemsetTypeD16, N);
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("Memset with hipMemsetTypeD8") {
|
||||
char *cA_d{nullptr}, *cA_h{nullptr};
|
||||
ret = testhipMemset(cA_h, cA_d, memsetD8val, hipMemsetTypeD8, N);
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Test hipMemsetAsync, hipMemsetD8Async, hipMemsetD16Async, hipMemsetD32Async
|
||||
* apis with unique number of elements and memset values.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemsetAsync_SetMemoryWithOffset") {
|
||||
char memsetval;
|
||||
int memsetD32val;
|
||||
int16_t memsetD16val;
|
||||
char memsetD8val;
|
||||
size_t N;
|
||||
bool ret;
|
||||
|
||||
std::tie(N, memsetval, memsetD32val, memsetD16val, memsetD8val) =
|
||||
GENERATE(table<size_t, char, int, int16_t, char>(tableItems));
|
||||
|
||||
|
||||
SECTION("Memset with hipMemsetTypeDefault") {
|
||||
char *cA_d{nullptr}, *cA_h{nullptr};
|
||||
ret = testhipMemsetAsync(cA_h, cA_d, memsetval, hipMemsetTypeDefault, N);
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("Memset with hipMemsetTypeD32") {
|
||||
int32_t *iA_d{nullptr}, *iA_h{nullptr};
|
||||
ret = testhipMemsetAsync(iA_h, iA_d, memsetD32val, hipMemsetTypeD32, N);
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("Memset with hipMemsetTypeD16") {
|
||||
int16_t *siA_d{nullptr}, *siA_h{nullptr};
|
||||
ret = testhipMemsetAsync(siA_h, siA_d, memsetD16val, hipMemsetTypeD16, N);
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("Memset with hipMemsetTypeD8") {
|
||||
char *cA_d{nullptr}, *cA_h{nullptr};
|
||||
ret = testhipMemsetAsync(cA_h, cA_d, memsetD8val, hipMemsetTypeD8, N);
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test hipMemset small size buffers with unique memset values.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset_SmallBufferSizes") {
|
||||
char *A_d, *A_h;
|
||||
constexpr int memsetval = 0x24;
|
||||
|
||||
auto numElements = GENERATE(range(1, 4));
|
||||
int numBytes = numElements * sizeof(char);
|
||||
|
||||
HIP_CHECK(hipMalloc(&A_d, numBytes));
|
||||
A_h = reinterpret_cast<char*> (malloc(numBytes));
|
||||
|
||||
HIP_CHECK(hipMemset(A_d, memsetval, numBytes));
|
||||
HIP_CHECK(hipMemcpy(A_h, A_d, numBytes, hipMemcpyDeviceToHost));
|
||||
|
||||
for (int i = 0; i < numBytes; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
INFO("Mismatch at index:" << i << " computed:" << A_h[i]
|
||||
<< " memsetval:" << memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Test two memset async operations at the same time.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset_2AsyncOperations") {
|
||||
std::vector<float> v;
|
||||
v.resize(2048);
|
||||
float* p2, *p3;
|
||||
hipMalloc(reinterpret_cast<void**>(&p2), 4096 + 4096*2);
|
||||
p3 = p2+2048;
|
||||
hipStream_t s;
|
||||
hipStreamCreate(&s);
|
||||
hipMemsetAsync(p2, 0, 32*32*4, s);
|
||||
hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s);
|
||||
hipStreamSynchronize(s);
|
||||
for (int i = 0; i < 256; ++i) {
|
||||
hipMemsetAsync(p2, 0, 32*32*4, s);
|
||||
hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s);
|
||||
}
|
||||
hipStreamSynchronize(s);
|
||||
hipDeviceSynchronize();
|
||||
hipMemcpy(&v[0], p2, 1024, hipMemcpyDeviceToHost);
|
||||
hipMemcpy(&v[1024], p3, 1024, hipMemcpyDeviceToHost);
|
||||
|
||||
REQUIRE(v[0] == 0);
|
||||
REQUIRE(v[1024] == 1.75f);
|
||||
}
|
||||
@@ -0,0 +1,175 @@
|
||||
/*
|
||||
* Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) hipMemset2D api with basic functionality.
|
||||
2) hipMemset2DAsync api with basic functionality.
|
||||
3) hipMemset2D api with partial memset and unique width/height.
|
||||
*/
|
||||
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
|
||||
// Table with unique width/height and memset values.
|
||||
// (width2D, height2D, memsetWidth, memsetHeight)
|
||||
typedef std::tuple<int, int, int, int> tupletype;
|
||||
|
||||
static constexpr std::initializer_list<tupletype> tableItems {
|
||||
std::make_tuple(20, 20, 20, 20),
|
||||
std::make_tuple(10, 10, 4, 4),
|
||||
std::make_tuple(100, 100, 20, 40),
|
||||
std::make_tuple(256, 256, 39, 19),
|
||||
std::make_tuple(100, 100, 20, 0),
|
||||
std::make_tuple(100, 100, 0, 20),
|
||||
std::make_tuple(100, 100, 0, 0),
|
||||
};
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* Basic Functionality of hipMemset2D
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset2D_BasicFunctional") {
|
||||
constexpr int memsetval = 0x24;
|
||||
constexpr size_t numH = 256;
|
||||
constexpr size_t numW = 256;
|
||||
size_t pitch_A;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH;
|
||||
size_t elements = numW * numH;
|
||||
char *A_d, *A_h;
|
||||
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width,
|
||||
numH));
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH));
|
||||
HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
INFO("Memset2D mismatch at index:" << i << " computed:"
|
||||
<< A_h[i] << " memsetval:" << memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
hipFree(A_d);
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Basic Functionality of hipMemset2DAsync
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset2DAsync_BasicFunctional") {
|
||||
constexpr int memsetval = 0x26;
|
||||
constexpr size_t numH = 256;
|
||||
constexpr size_t numW = 256;
|
||||
size_t pitch_A;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH;
|
||||
size_t elements = numW * numH;
|
||||
char *A_d, *A_h;
|
||||
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
|
||||
width, numH));
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i=0; i < elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
INFO("Memset2DAsync mismatch at index:" << i << " computed:"
|
||||
<< A_h[i] << " memsetval:" << memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
hipFree(A_d);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Memset partial buffer with unique Width and Height
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset2D_UniqueWidthHeight") {
|
||||
int width2D, height2D;
|
||||
int memsetWidth, memsetHeight;
|
||||
char *A_d, *A_h;
|
||||
size_t pitch_A;
|
||||
constexpr int memsetval = 0x26;
|
||||
|
||||
std::tie(width2D, height2D, memsetWidth, memsetHeight) =
|
||||
GENERATE(table<int, int, int, int>(tableItems));
|
||||
|
||||
size_t width = width2D * sizeof(char);
|
||||
size_t sizeElements = width * height2D;
|
||||
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
|
||||
width, height2D));
|
||||
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
|
||||
for (size_t index = 0; index < sizeElements; index++) {
|
||||
A_h[index] = 'c';
|
||||
}
|
||||
|
||||
INFO("2D Dimension: Width:" << width2D << " Height:" << height2D <<
|
||||
" MemsetWidth:" << memsetWidth << " MemsetHeight:" << memsetHeight);
|
||||
|
||||
HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, memsetWidth, memsetHeight));
|
||||
HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, width2D, height2D,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (int row = 0; row < memsetHeight; row++) {
|
||||
for (int column = 0; column < memsetWidth; column++) {
|
||||
if (A_h[(row * width) + column] != memsetval) {
|
||||
INFO("A_h[" << row << "][" << column << "]" <<
|
||||
" didnot match " << memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
hipFree(A_d);
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
@@ -0,0 +1,185 @@
|
||||
/*
|
||||
* Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Order of execution of device kernel and hipMemset2DAsync api
|
||||
2) hipMemSet2DAsync execution in multiple threads
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_helper.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
|
||||
|
||||
/* Defines */
|
||||
#define NUM_THREADS 1000
|
||||
#define ITER 100
|
||||
#define NUM_H 256
|
||||
#define NUM_W 256
|
||||
|
||||
|
||||
|
||||
void queueJobsForhipMemset2DAsync(char* A_d, char* A_h, size_t pitch,
|
||||
size_t width, hipStream_t stream) {
|
||||
constexpr int memsetval = 0x22;
|
||||
HIPCHECK(hipMemset2DAsync(A_d, pitch, memsetval, NUM_W, NUM_H, stream));
|
||||
HIPCHECK(hipMemcpy2DAsync(A_h, width, A_d, pitch, NUM_W, NUM_H,
|
||||
hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Order of execution of device kernel and hipMemset2DAsync api.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset2DAsync_WithKernel") {
|
||||
constexpr auto N = 4 * 1024 * 1024;
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
constexpr int memsetval = 0x22;
|
||||
char *A_d, *A_h, *B_d, *B_h, *C_d;
|
||||
size_t pitch_A, pitch_B, pitch_C;
|
||||
size_t width = NUM_W * sizeof(char);
|
||||
size_t sizeElements = width * NUM_H;
|
||||
size_t elements = NUM_W * NUM_H;
|
||||
unsigned blocks{};
|
||||
int validateCount{};
|
||||
hipStream_t stream;
|
||||
|
||||
blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
|
||||
width, NUM_H));
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&B_d), &pitch_B,
|
||||
width, NUM_H));
|
||||
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
B_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
REQUIRE(B_h != nullptr);
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&C_d), &pitch_C,
|
||||
width, NUM_H));
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
B_h[i] = i;
|
||||
}
|
||||
HIP_CHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H,
|
||||
hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
|
||||
for (size_t k = 0; k < ITER; k++) {
|
||||
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, stream, B_d, C_d, elements);
|
||||
|
||||
HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t p = 0 ; p < elements ; p++) {
|
||||
if (A_h[p] == memsetval) {
|
||||
validateCount+= 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
REQUIRE(static_cast<size_t>(validateCount) == (ITER * elements));
|
||||
|
||||
HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d));
|
||||
free(A_h); free(B_h);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* hipMemSet2DAsync execution in multiple threads.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset2DAsync_MultiThread") {
|
||||
constexpr auto N = 4 * 1024 * 1024;
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
constexpr auto memPerThread = 200;
|
||||
constexpr int memsetval = 0x22;
|
||||
char *A_d, *A_h, *B_d, *B_h, *C_d;
|
||||
size_t pitch_A, pitch_B, pitch_C;
|
||||
size_t width = NUM_W * sizeof(char);
|
||||
size_t sizeElements = width * NUM_H;
|
||||
size_t elements = NUM_W * NUM_H;
|
||||
unsigned blocks{};
|
||||
int validateCount{};
|
||||
hipStream_t stream;
|
||||
|
||||
blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
auto thread_count = HipTest::getHostThreadCount(memPerThread, NUM_THREADS);
|
||||
if (thread_count == 0) {
|
||||
WARN("Resources not available for thread creation");
|
||||
return;
|
||||
}
|
||||
|
||||
std::thread *t = new std::thread[thread_count];
|
||||
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
|
||||
width, NUM_H));
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&B_d), &pitch_B,
|
||||
width, NUM_H));
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
B_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
REQUIRE(B_h != nullptr);
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&C_d), &pitch_C,
|
||||
width, NUM_H));
|
||||
|
||||
for (size_t i = 0 ; i < elements ; i++) {
|
||||
B_h[i] = i;
|
||||
}
|
||||
HIP_CHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H,
|
||||
hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
for (int i = 0 ; i < ITER ; i++) {
|
||||
for (size_t k = 0 ; k < thread_count; k++) {
|
||||
if (k%2) {
|
||||
t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, A_h, pitch_A,
|
||||
width, stream);
|
||||
} else {
|
||||
t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, B_h, pitch_A,
|
||||
width, stream);
|
||||
}
|
||||
}
|
||||
for (size_t j = 0 ; j < thread_count; j++) {
|
||||
t[j].join();
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
for (size_t k = 0 ; k < elements ; k++) {
|
||||
if ((A_h[k] == memsetval) && (B_h[k] == memsetval)) {
|
||||
validateCount+= 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
REQUIRE(static_cast<size_t>(validateCount) == (ITER * elements));
|
||||
|
||||
HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d));
|
||||
free(A_h); free(B_h);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
|
||||
delete[] t;
|
||||
}
|
||||
@@ -0,0 +1,128 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Functional test for Memset3D and Memset3DAsync
|
||||
*/
|
||||
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
|
||||
/**
|
||||
* Basic Functional test of hipMemset3D
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3D_BasicFunctional") {
|
||||
constexpr int memsetval = 0x22;
|
||||
constexpr size_t numH = 256;
|
||||
constexpr size_t numW = 256;
|
||||
constexpr size_t depth = 10;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW * numH * depth;
|
||||
char *A_h;
|
||||
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
|
||||
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
A_h = reinterpret_cast<char *>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent));
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
|
||||
myparms.srcPtr = devPitchedPtr;
|
||||
myparms.extent = extent;
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
INFO("Memset3D mismatch at index:" << i << " computed:"
|
||||
<< A_h[i] << " memsetval:" << memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipFree(devPitchedPtr.ptr));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* Basic Functional test of hipMemset3DAsync
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_BasicFunctional") {
|
||||
constexpr int memsetval = 0x22;
|
||||
constexpr size_t numH = 256;
|
||||
constexpr size_t numW = 256;
|
||||
constexpr size_t depth = 10;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW * numH * depth;
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
char *A_h;
|
||||
|
||||
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
A_h = reinterpret_cast<char *>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
|
||||
myparms.srcPtr = devPitchedPtr;
|
||||
myparms.extent = extent;
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
INFO("Memset3DAsync mismatch at index:" << i << " computed:"
|
||||
<< A_h[i] << " memsetval:" << memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipFree(devPitchedPtr.ptr));
|
||||
free(A_h);
|
||||
}
|
||||
@@ -0,0 +1,515 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
|
||||
1) Passing width as 0 in extent, verify hipMemset3D api returns success and
|
||||
doesn't modify the buffer passed.
|
||||
2) Passing width as 0 in extent, verify hipMemset3DAsync api returns success
|
||||
and doesn't modify the buffer passed.
|
||||
|
||||
3) Passing height as 0 in extent, verify hipMemset3D api returns success and
|
||||
doesn't modify the buffer passed.
|
||||
4) Passing height as 0 in extent, verify hipMemset3DAsync api returns success
|
||||
and doesn't modify the buffer passed.
|
||||
|
||||
5) Passing depth as 0 in extent, verify hipMemset3D api returns success and
|
||||
doesn't modify the buffer passed.
|
||||
6) Passing depth as 0 in extent, verify hipMemset3DAsync api returns success
|
||||
and doesn't modify the buffer passed.
|
||||
|
||||
7) When extent passed with width, height and depth all as zeroes, verify
|
||||
hipMemset3D api returns success and doesn't modify the buffer passed.
|
||||
8) When extent passed with width, height and depth all as zeroes, verify
|
||||
hipMemset3DAsync api returns success and doesn't modify the buffer passed.
|
||||
|
||||
9) Validate data after performing memory set operation with max memset value
|
||||
for hipMemset3D api.
|
||||
10) Validate data after performing memory set operation with max memset value
|
||||
for hipMemset3DAsync api.
|
||||
|
||||
11) Select random slice of 3d array and Memset complete slice with
|
||||
hipMemset3D api.
|
||||
12) Select random slice of 3d array and Memset complete slice with
|
||||
hipMemset3DAsync api.
|
||||
|
||||
13) Seek device pitched ptr to desired portion of 3d array and memset the
|
||||
portion with hipMemset3D api.
|
||||
14) Seek device pitched ptr to desired portion of 3d array and memset the
|
||||
portion with hipMemset3DAsync api.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
/*
|
||||
* Defines
|
||||
*/
|
||||
#define MEMSETVAL 1
|
||||
#define TESTVAL 2
|
||||
#define NUMH_EXT 256
|
||||
#define NUMW_EXT 100
|
||||
#define DEPTH_EXT 10
|
||||
#define NUMH_MAX 256
|
||||
#define NUMW_MAX 256
|
||||
#define DEPTH_MAX 10
|
||||
#define ZSIZE_S 32
|
||||
#define YSIZE_S 32
|
||||
#define XSIZE_S 32
|
||||
#define ZSIZE_P 30
|
||||
#define YSIZE_P 30
|
||||
#define XSIZE_P 30
|
||||
#define ZPOS_START 10
|
||||
#define ZSET_LEN 10
|
||||
#define ZPOS_END 19
|
||||
#define YPOS_START 10
|
||||
#define YSET_LEN 10
|
||||
#define YPOS_END 19
|
||||
#define XPOS_START 10
|
||||
#define XSET_LEN 10
|
||||
#define XPOS_END 19
|
||||
|
||||
|
||||
/**
|
||||
* Memset with extent passed and verify data to be intact
|
||||
*/
|
||||
static void testMemsetWithExtent(bool bAsync, hipExtent tstExtent) {
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
hipError_t ret;
|
||||
char *A_h;
|
||||
size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT;
|
||||
size_t width = numW * sizeof(char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
|
||||
A_h = reinterpret_cast<char *>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
memset(A_h, 0, sizeElements);
|
||||
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
if (bAsync) {
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
ret = hipMemset3DAsync(devPitchedPtr, MEMSETVAL, extent, stream);
|
||||
INFO("testMemsetWithExtent(" << extent.width << "," << extent.height
|
||||
<< "," << extent.depth << ") memset "
|
||||
<< MEMSETVAL << ", ret : " << ret);
|
||||
REQUIRE(ret == hipSuccess);
|
||||
|
||||
ret = hipMemset3DAsync(devPitchedPtr, TESTVAL, tstExtent, stream);
|
||||
INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height
|
||||
<< "," << tstExtent.depth << ") memset "
|
||||
<< TESTVAL << "ret : " << ret);
|
||||
REQUIRE(ret == hipSuccess);
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
} else {
|
||||
ret = hipMemset3D(devPitchedPtr, MEMSETVAL, extent);
|
||||
INFO("testMemsetWithExtent(" << extent.width << "," << extent.height
|
||||
<< "," << extent.depth << ") memset "
|
||||
<< MEMSETVAL << ",ret : " << ret);
|
||||
REQUIRE(ret == hipSuccess);
|
||||
|
||||
ret = hipMemset3D(devPitchedPtr, TESTVAL, tstExtent);
|
||||
INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height
|
||||
<< "," << tstExtent.depth << ") memset "
|
||||
<< TESTVAL << ",ret : " << ret);
|
||||
REQUIRE(ret == hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
|
||||
myparms.srcPtr = devPitchedPtr;
|
||||
myparms.extent = extent;
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
if (A_h[i] != MEMSETVAL) {
|
||||
INFO("testMemsetWithExtent: index:" << i << ",computed:"
|
||||
<< std::hex << static_cast<int>(A_h[i]) << ",memsetval:"
|
||||
<< std::hex << MEMSETVAL);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(devPitchedPtr.ptr));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Validates data after performing memory set operation with max memset value
|
||||
*/
|
||||
static void testMemsetMaxValue(bool bAsync) {
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
unsigned char *A_h;
|
||||
int memsetval = std::numeric_limits<unsigned char>::max();
|
||||
size_t numH = NUMH_MAX, numW = NUMW_MAX, depth = DEPTH_MAX;
|
||||
size_t width = numW * sizeof(unsigned char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
|
||||
A_h = reinterpret_cast<unsigned char *> (malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
memset(A_h, 0, sizeElements);
|
||||
|
||||
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
if (bAsync) {
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
} else {
|
||||
HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent));
|
||||
}
|
||||
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
|
||||
myparms.srcPtr = devPitchedPtr;
|
||||
myparms.extent = extent;
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
INFO("testMemsetMaxValue: index:" << i << ",computed:"
|
||||
<< std::hex << static_cast<int>(A_h[i]) << ",memsetval:"
|
||||
<< std::hex << memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipFree(devPitchedPtr.ptr));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* Function seeks device ptr to random slice and performs Memset operation
|
||||
* on the slice selected.
|
||||
*/
|
||||
static void seekAndSet3DArraySlice(bool bAsync) {
|
||||
char array3D[ZSIZE_S][YSIZE_S][XSIZE_S]{};
|
||||
dim3 arr_dimensions = dim3(ZSIZE_S, YSIZE_S, XSIZE_S);
|
||||
hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x,
|
||||
arr_dimensions.y, arr_dimensions.z);
|
||||
hipPitchedPtr devicePitchedPointer;
|
||||
int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL;
|
||||
|
||||
HIP_CHECK(hipMalloc3D(&devicePitchedPointer, extent));
|
||||
HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent));
|
||||
|
||||
// select random slice for memset
|
||||
unsigned int seed = time(nullptr);
|
||||
int slice_index = rand_r(&seed) % ZSIZE_S;
|
||||
|
||||
INFO("memset3d for sliceindex " << slice_index);
|
||||
|
||||
// Get attributes from device pitched pointer
|
||||
size_t pitch = devicePitchedPointer.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
|
||||
// Point devptr to selected slice
|
||||
char *devPtrSlice = (reinterpret_cast<char *>(devicePitchedPointer.ptr))
|
||||
+ slice_index * slicePitch;
|
||||
hipExtent extentSlice = make_hipExtent(sizeof(char) * arr_dimensions.x,
|
||||
arr_dimensions.y, 1);
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrSlice, pitch,
|
||||
arr_dimensions.x, arr_dimensions.y);
|
||||
|
||||
if (bAsync) {
|
||||
// Memset selected slice (Async)
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked,
|
||||
extentSlice, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
} else {
|
||||
// Memset selected slice
|
||||
HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, extentSlice));
|
||||
}
|
||||
|
||||
// Copy result back to host buffer
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x,
|
||||
arr_dimensions.x, arr_dimensions.y);
|
||||
myparms.srcPtr = devicePitchedPointer;
|
||||
myparms.extent = extent;
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (int z = 0; z < ZSIZE_S; z++) {
|
||||
for (int y = 0; y < YSIZE_S; y++) {
|
||||
for (int x = 0; x < XSIZE_S; x++) {
|
||||
if (z == slice_index) {
|
||||
if (array3D[z][y][x] != memsetval4seeked) {
|
||||
INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" << z
|
||||
<< "," << y << "," << x << ") " << "computed:" << std::hex
|
||||
<< array3D[z][y][x] << ", memsetval:" << std::hex
|
||||
<< memsetval4seeked);
|
||||
REQUIRE(false);
|
||||
}
|
||||
} else {
|
||||
if (array3D[z][y][x] != memsetval) {
|
||||
INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" << z
|
||||
<< "," << y << "," << x << ") " << "computed:" << std::hex
|
||||
<< array3D[z][y][x] << ", memsetval:" << std::hex
|
||||
<< memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(devicePitchedPointer.ptr));
|
||||
}
|
||||
|
||||
/**
|
||||
* Function seeks device ptr to selected portion of 3d array
|
||||
* and performs Memset operation on the portion.
|
||||
*/
|
||||
static void seekAndSet3DArrayPortion(bool bAsync) {
|
||||
char array3D[ZSIZE_P][YSIZE_P][XSIZE_P]{};
|
||||
dim3 arr_dimensions = dim3(ZSIZE_P, YSIZE_P, XSIZE_P);
|
||||
hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x,
|
||||
arr_dimensions.y, arr_dimensions.z);
|
||||
hipPitchedPtr devicePitchedPointer;
|
||||
int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL;
|
||||
|
||||
HIP_CHECK(hipMalloc3D(&devicePitchedPointer, extent));
|
||||
HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent));
|
||||
|
||||
// For memsetting extent/size(10,10,10) in the mid portion of cube(30,30,30),
|
||||
// seek device ptr to (10,10,10) and then memset 10 bytes across x,y,z axis.
|
||||
size_t pitch = devicePitchedPointer.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
int slice_index = ZPOS_START, y = YPOS_START, x = XPOS_START;
|
||||
|
||||
// Select 10th slice
|
||||
char *devPtrSlice = (reinterpret_cast<char *>(devicePitchedPointer.ptr))
|
||||
+ slice_index * slicePitch;
|
||||
|
||||
// Now select row at height as 10
|
||||
char *current_row = reinterpret_cast<char *>(devPtrSlice + y * pitch);
|
||||
|
||||
// Now select index of selected row as 10
|
||||
char *devPtrIndexed = ¤t_row[x];
|
||||
|
||||
// Make dev Pitchedptr, extent
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrIndexed, pitch,
|
||||
arr_dimensions.x, arr_dimensions.y);
|
||||
hipExtent setExtent = make_hipExtent(sizeof(char) * XSET_LEN, YSET_LEN,
|
||||
ZSET_LEN);
|
||||
|
||||
if (bAsync) {
|
||||
// Memset selected portion (Async)
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked,
|
||||
setExtent, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
} else {
|
||||
// Memset selected portion
|
||||
HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, setExtent));
|
||||
}
|
||||
|
||||
// Copy result back to host buffer
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x,
|
||||
arr_dimensions.x, arr_dimensions.y);
|
||||
myparms.srcPtr = devicePitchedPointer;
|
||||
myparms.extent = extent;
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (int z = 0; z < ZSIZE_P; z++) {
|
||||
for (int y = 0; y < YSIZE_P; y++) {
|
||||
for (int x = 0; x < XSIZE_P; x++) {
|
||||
if ((z >= ZPOS_START && z <= ZPOS_END) &&
|
||||
(y >= YPOS_START && y <= YPOS_END) &&
|
||||
(x >= XPOS_START && x <= XPOS_END)) {
|
||||
if (array3D[z][y][x] != memsetval4seeked) {
|
||||
INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" << z
|
||||
<< "," << y << "," << x << ") " << "computed:" << std::hex
|
||||
<< array3D[z][y][x] << ", memsetval:" << std::hex
|
||||
<< memsetval4seeked);
|
||||
REQUIRE(false);
|
||||
}
|
||||
} else {
|
||||
if (array3D[z][y][x] != memsetval) {
|
||||
INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" << z
|
||||
<< "," << y << "," << x << ") " << "computed:" << std::hex
|
||||
<< array3D[z][y][x] << ", memsetval:" << std::hex
|
||||
<< memsetval);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(devicePitchedPointer.ptr));
|
||||
}
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* Test Memset3D with different combinations of extent
|
||||
* taking zero and non-zero fields.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3D_MemsetWithExtent") {
|
||||
hipExtent testExtent;
|
||||
size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT;
|
||||
|
||||
SECTION("Memset with extent width(0)") {
|
||||
// Memset with extent width(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(0, numH, depth);
|
||||
testMemsetWithExtent(0, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent height(0)") {
|
||||
// Memset with extent height(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(numW, 0, depth);
|
||||
testMemsetWithExtent(0, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent depth(0)") {
|
||||
// Memset with extent depth(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(numW, numH, 0);
|
||||
testMemsetWithExtent(0, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent width,height,depth as 0") {
|
||||
// Memset with extent width,height,depth as 0 and verify data to be intact
|
||||
testExtent = make_hipExtent(0, 0, 0);
|
||||
testMemsetWithExtent(0, testExtent);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Test Memset3DAsync with different combinations of extent
|
||||
* taking zero and non-zero fields.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_MemsetWithExtent") {
|
||||
hipExtent testExtent;
|
||||
size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT;
|
||||
|
||||
SECTION("Memset with extent width(0)") {
|
||||
// Memset with extent width(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(0, numH, depth);
|
||||
testMemsetWithExtent(1, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent height(0)") {
|
||||
// Memset with extent height(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(numW, 0, depth);
|
||||
testMemsetWithExtent(1, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent depth(0)") {
|
||||
// Memset with extent depth(0) and verify data to be intact
|
||||
testExtent = make_hipExtent(numW, numH, 0);
|
||||
testMemsetWithExtent(1, testExtent);
|
||||
}
|
||||
|
||||
SECTION("Memset with extent width,height,depth as 0") {
|
||||
// Memset with extent width,height,depth as 0 and verify data to be intact
|
||||
testExtent = make_hipExtent(0, 0, 0);
|
||||
testMemsetWithExtent(1, testExtent);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Memset3D with max unsigned char and verify memset operation is success
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3D_MemsetMaxValue") {
|
||||
testMemsetMaxValue(0);
|
||||
}
|
||||
|
||||
/**
|
||||
* Memset3DAsync with max unsigned char and verify memset operation is success
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_MemsetMaxValue") {
|
||||
testMemsetMaxValue(1);
|
||||
}
|
||||
|
||||
/**
|
||||
* Seek and set random slice of 3d array, verify memset is success
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3D_SeekSetSlice") {
|
||||
seekAndSet3DArraySlice(0);
|
||||
}
|
||||
|
||||
/**
|
||||
* Seek and set random slice of 3d array with async, verify memset is success
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_SeekSetSlice") {
|
||||
seekAndSet3DArraySlice(1);
|
||||
}
|
||||
|
||||
/**
|
||||
* Memset3D selected portion of 3d array
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3D_SeekSetArrayPortion") {
|
||||
seekAndSet3DArrayPortion(0);
|
||||
}
|
||||
|
||||
/**
|
||||
* Memset3DAsync selected portion of 3d array
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_SeekSetArrayPortion") {
|
||||
seekAndSet3DArrayPortion(1);
|
||||
}
|
||||
@@ -0,0 +1,236 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Test hipMemset3D() with uninitialized devPitchedPtr.
|
||||
2) Test hipMemset3DAsync() with uninitialized devPitchedPtr.
|
||||
|
||||
3) Reset devPitchedPtr to zero and check return value for hipMemset3D().
|
||||
4) Reset devPitchedPtr to zero and check return value for hipMemset3DAsync().
|
||||
|
||||
5) Test hipMemset3D() with extent.width as max size_t and keeping height,
|
||||
depth as valid values.
|
||||
6) Test hipMemset3DAsync() with extent.width as max size_t and keeping height,
|
||||
depth as valid values.
|
||||
7) Test hipMemset3D() with extent.height as max size_t and keeping width,
|
||||
depth as valid values.
|
||||
8) Test hipMemset3DAsync() with extent.height as max size_t and keeping width,
|
||||
depth as valid values.
|
||||
9) Test hipMemset3D() with extent.depth as max size_t and keeping height,
|
||||
width as valid values.
|
||||
10) Test hipMemset3DAsync() with extent.depth as max size_t and keeping height,
|
||||
width as valid values.
|
||||
|
||||
11) Device Ptr out bound and extent(0) passed for hipMemset3D().
|
||||
12) Device Ptr out bound and extent(0) passed for hipMemset3DAsync().
|
||||
|
||||
13) Device Ptr out bound and valid extent passed for hipMemset3D().
|
||||
14) Device Ptr out bound and valid extent passed for hipMemset3DAsync().
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
TEST_CASE("Unit_hipMemset3D_Negative") {
|
||||
hipError_t ret;
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
constexpr int memsetval = 1;
|
||||
constexpr size_t numH = 256, numW = 256;
|
||||
constexpr size_t depth = 10;
|
||||
constexpr size_t width = numW * sizeof(char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
|
||||
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
|
||||
SECTION("Using uninitialized devpitched ptr") {
|
||||
hipPitchedPtr devPitchedUnPtr;
|
||||
|
||||
ret = hipMemset3D(devPitchedUnPtr, memsetval, extent);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Reset devPitchedPtr to zero") {
|
||||
hipPitchedPtr rdevPitchedPtr{};
|
||||
|
||||
ret = hipMemset3D(rdevPitchedPtr, memsetval, extent);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Pass extent fields as max size_t") {
|
||||
hipExtent extMW = make_hipExtent(std::numeric_limits<std::size_t>::max(),
|
||||
numH,
|
||||
depth);
|
||||
hipExtent extMH = make_hipExtent(width,
|
||||
std::numeric_limits<std::size_t>::max(),
|
||||
depth);
|
||||
hipExtent extMD = make_hipExtent(width,
|
||||
numH,
|
||||
std::numeric_limits<std::size_t>::max());
|
||||
|
||||
ret = hipMemset3D(devPitchedPtr, memsetval, extMW);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
|
||||
ret = hipMemset3D(devPitchedPtr, memsetval, extMH);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
|
||||
if ((TestContext::get()).isAmd()) {
|
||||
ret = hipMemset3D(devPitchedPtr, memsetval, extMD);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
} else {
|
||||
WARN("Test is skipped for max depth."
|
||||
<< "Cuda doesn't check the maximum depth of extent field");
|
||||
}
|
||||
}
|
||||
|
||||
SECTION("Device Ptr out bound and extent(0) passed for memset") {
|
||||
size_t pitch = devPitchedPtr.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
constexpr auto advanceOffset = 10;
|
||||
|
||||
// Point devptr to end of allocated memory
|
||||
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
|
||||
+ depth * slicePitch;
|
||||
|
||||
// Advance devptr further to go out of boundary
|
||||
devPtrMod = devPtrMod + advanceOffset;
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
|
||||
numW * sizeof(char), numH);
|
||||
hipExtent extent0{};
|
||||
ret = hipMemset3D(modDevPitchedPtr, memsetval, extent0);
|
||||
|
||||
// api expected to check extent0 and return success before going for ptr.
|
||||
REQUIRE(ret == hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Device Ptr out bound and valid extent passed for memset") {
|
||||
size_t pitch = devPitchedPtr.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
constexpr auto advanceOffset = 10;
|
||||
|
||||
// Point devptr to end of allocated memory
|
||||
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
|
||||
+ depth * slicePitch;
|
||||
|
||||
// Advance devptr further to go out of boundary
|
||||
devPtrMod = devPtrMod + advanceOffset;
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
|
||||
numW * sizeof(char), numH);
|
||||
ret = hipMemset3D(modDevPitchedPtr, memsetval, extent);
|
||||
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(devPitchedPtr.ptr));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemset3DAsync_Negative") {
|
||||
hipError_t ret;
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
hipStream_t stream;
|
||||
constexpr int memsetval = 1;
|
||||
constexpr size_t numH = 256;
|
||||
constexpr size_t numW = 256;
|
||||
constexpr size_t depth = 10;
|
||||
constexpr size_t width = numW * sizeof(char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
|
||||
SECTION("Using uninitialized devpitched ptr") {
|
||||
hipPitchedPtr devPitchedUnPtr;
|
||||
|
||||
ret = hipMemset3DAsync(devPitchedUnPtr, memsetval, extent, stream);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Reset devPitchedPtr to zero") {
|
||||
hipPitchedPtr rdevPitchedPtr{};
|
||||
|
||||
ret = hipMemset3DAsync(rdevPitchedPtr, memsetval, extent, stream);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Pass extent fields as max size_t") {
|
||||
hipExtent extMW = make_hipExtent(std::numeric_limits<std::size_t>::max(),
|
||||
numH,
|
||||
depth);
|
||||
hipExtent extMH = make_hipExtent(width,
|
||||
std::numeric_limits<std::size_t>::max(),
|
||||
depth);
|
||||
hipExtent extMD = make_hipExtent(width,
|
||||
numH,
|
||||
std::numeric_limits<std::size_t>::max());
|
||||
|
||||
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMW, stream);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
|
||||
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMH, stream);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
|
||||
if ((TestContext::get()).isAmd()) {
|
||||
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMD, stream);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
} else {
|
||||
WARN("Test is skipped for max depth."
|
||||
<< "Cuda doesn't check the maximum depth of extent field");
|
||||
}
|
||||
}
|
||||
|
||||
SECTION("Device Ptr out bound and extent(0) passed for memset") {
|
||||
size_t pitch = devPitchedPtr.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
constexpr auto advanceOffset = 10;
|
||||
|
||||
// Point devptr to end of allocated memory
|
||||
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
|
||||
+ depth * slicePitch;
|
||||
|
||||
// Advance devptr further to go out of boundary
|
||||
devPtrMod = devPtrMod + advanceOffset;
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
|
||||
numW * sizeof(char), numH);
|
||||
hipExtent extent0{};
|
||||
ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent0, stream);
|
||||
|
||||
// api expected to check extent0 and return success before going for ptr.
|
||||
REQUIRE(ret == hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Device Ptr out bound and valid extent passed for memset") {
|
||||
size_t pitch = devPitchedPtr.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
constexpr auto advanceOffset = 10;
|
||||
|
||||
// Point devptr to end of allocated memory
|
||||
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
|
||||
+ depth * slicePitch;
|
||||
|
||||
// Advance devptr further to go out of boundary
|
||||
devPtrMod = devPtrMod + advanceOffset;
|
||||
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
|
||||
numW * sizeof(char), numH);
|
||||
ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent, stream);
|
||||
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK(hipFree(devPitchedPtr.ptr));
|
||||
}
|
||||
@@ -0,0 +1,266 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Validate Async behavior of hipMemset3DAsync with commands queued
|
||||
concurrently from multiple threads.
|
||||
2) Validate hipMemset3DAsync behavior when api is queued along with kernel
|
||||
function operating on same memory.
|
||||
3) Perform regression of hipMemset3D api in loop with device memory allocated
|
||||
on different gpus.
|
||||
4) Perform regression of hipMemset3DAsync api in loop with device memory
|
||||
allocated on different gpus.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
|
||||
/*
|
||||
* Defines
|
||||
*/
|
||||
#define MAX_REGRESS_ITERS 2
|
||||
#define MAX_THREADS 10
|
||||
|
||||
/**
|
||||
* kernel function sets device memory with value passed
|
||||
*/
|
||||
static __global__ void func_set_value(hipPitchedPtr devicePitchedPointer,
|
||||
hipExtent extent,
|
||||
unsigned char val) {
|
||||
// Index Calculation
|
||||
size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
size_t y = threadIdx.y + blockDim.y * blockIdx.y;
|
||||
size_t z = threadIdx.z + blockDim.z * blockIdx.z;
|
||||
|
||||
// Get attributes from device pitched pointer
|
||||
char *devicePointer = reinterpret_cast<char *>(devicePitchedPointer.ptr);
|
||||
size_t pitch = devicePitchedPointer.pitch;
|
||||
size_t slicePitch = pitch * extent.height;
|
||||
|
||||
// Loop over the device buffer
|
||||
if (z < extent.depth) {
|
||||
char *current_slice_index = devicePointer + z * slicePitch;
|
||||
if (y < extent.height) {
|
||||
// Get data array containing all elements from the current row
|
||||
char *current_row = reinterpret_cast<char *>(current_slice_index
|
||||
+ y * pitch);
|
||||
if (x < extent.width) {
|
||||
current_row[x] = val;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Thread function queues kernel function and memset cmds
|
||||
*/
|
||||
static void threadFunc(hipStream_t stream, hipPitchedPtr devpPtr,
|
||||
int memsetval, int testval, hipExtent extent, hipMemcpy3DParms myparms) {
|
||||
// Kernel Launch Configuration
|
||||
constexpr auto size = 8;
|
||||
dim3 threadsPerBlock = dim3(size, size, size);
|
||||
dim3 blocks;
|
||||
blocks = dim3((extent.width + threadsPerBlock.x - 1) / threadsPerBlock.x,
|
||||
(extent.height + threadsPerBlock.y - 1) / threadsPerBlock.y,
|
||||
(extent.depth + threadsPerBlock.z - 1) / threadsPerBlock.z);
|
||||
|
||||
hipLaunchKernelGGL(func_set_value, dim3(blocks), dim3(threadsPerBlock), 0,
|
||||
stream, devpPtr, extent, memsetval);
|
||||
HIPCHECK(hipMemset3DAsync(devpPtr, testval, extent, stream));
|
||||
HIPCHECK(hipMemcpy3DAsync(&myparms, stream));
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Performs api regression in loop
|
||||
*/
|
||||
bool loopRegression(bool bAsync) {
|
||||
bool testPassed = true;
|
||||
char *A_h;
|
||||
constexpr int memsetval = 1;
|
||||
constexpr size_t numH = 256, numW = 100, depth = 10;
|
||||
int numGpu = 0, hasPeerAccess = 0;
|
||||
size_t width = numW * sizeof(char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
std::vector<hipPitchedPtr> devPitchedPtrlist;
|
||||
hipPitchedPtr pitchedPtr, devpPtr;
|
||||
|
||||
A_h = reinterpret_cast<char *>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
memset(A_h, 0, sizeElements);
|
||||
|
||||
// Populate hipMemcpy3D parameters
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
|
||||
myparms.extent = extent;
|
||||
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipGetDeviceCount(&numGpu));
|
||||
REQUIRE(numGpu > 0);
|
||||
|
||||
// Alloc 3D arrays in all GPUs
|
||||
for (int j = 0; j < numGpu; j++) {
|
||||
HIP_CHECK(hipSetDevice(j));
|
||||
HIP_CHECK(hipMalloc3D(&pitchedPtr, extent));
|
||||
devPitchedPtrlist.push_back(pitchedPtr);
|
||||
}
|
||||
|
||||
for (int itern = 0; itern < MAX_REGRESS_ITERS; itern++) {
|
||||
// Validate hipMemset3D data consistency in multiple iters
|
||||
for (int i = 0; i < numGpu; i++) {
|
||||
for (int j = 0; j < numGpu; j++) {
|
||||
HIP_CHECK(hipDeviceCanAccessPeer(&hasPeerAccess, i, j));
|
||||
if (!hasPeerAccess) {
|
||||
// Skip and continue if no peer access
|
||||
continue;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipSetDevice(i));
|
||||
devpPtr = devPitchedPtrlist[j];
|
||||
HIP_CHECK(hipDeviceEnablePeerAccess(j, 0));
|
||||
HIP_CHECK(hipMemset3D(devpPtr, 0, extent));
|
||||
|
||||
if (bAsync) {
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMemset3DAsync(devpPtr, memsetval, extent, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
} else {
|
||||
HIP_CHECK(hipMemset3D(devpPtr, memsetval, extent));
|
||||
}
|
||||
|
||||
myparms.srcPtr = devpPtr;
|
||||
memset(A_h, 0, sizeElements);
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (size_t indx = 0; indx < elements; indx++) {
|
||||
if (A_h[indx] != memsetval) {
|
||||
testPassed = false;
|
||||
printf("RegressIter : mismatch at index:%d computed:%02x, "
|
||||
"memsetval:%02x\n", static_cast<int>(indx),
|
||||
static_cast<int>(A_h[indx]), static_cast<int>(memsetval));
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
for (int j = 0; j < numGpu; j++) {
|
||||
HIP_CHECK(hipFree(devPitchedPtrlist[j].ptr));
|
||||
}
|
||||
|
||||
free(A_h);
|
||||
return testPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Perform regression of hipMemset3D api with device memory allocated
|
||||
* on different gpus.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3D_RegressInLoop") {
|
||||
bool TestPassed = false;
|
||||
|
||||
TestPassed = loopRegression(0);
|
||||
REQUIRE(TestPassed == true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Perform regression of hipMemset3DAsync api with device memory allocated
|
||||
* on different gpus.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_RegressInLoop") {
|
||||
bool TestPassed = false;
|
||||
|
||||
TestPassed = loopRegression(1);
|
||||
REQUIRE(TestPassed == true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Async commands queued concurrently and executed
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset3DAsync_ConcurrencyMthread") {
|
||||
char *A_h;
|
||||
constexpr int memsetval = 1, testval = 2;
|
||||
constexpr size_t numH = 256, numW = 100, depth = 10;
|
||||
size_t width = numW * sizeof(char);
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
hipPitchedPtr devpPtr;
|
||||
hipStream_t stream;
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMalloc3D(&devpPtr, extent));
|
||||
|
||||
A_h = reinterpret_cast<char *>(malloc(sizeElements));
|
||||
REQUIRE(A_h != nullptr);
|
||||
memset(A_h, 0, sizeElements);
|
||||
|
||||
// Populate hipMemcpy3D parameters
|
||||
hipMemcpy3DParms myparms{};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.srcPtr = devpPtr;
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
|
||||
myparms.extent = extent;
|
||||
|
||||
#if HT_NVIDIA
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
|
||||
std::vector<std::thread> threadlist;
|
||||
|
||||
// Queue cmds concurrently from multiple threads on same stream
|
||||
for (int i = 0; i < MAX_THREADS; i++) {
|
||||
threadlist.push_back(std::thread(threadFunc, stream, devpPtr, memsetval,
|
||||
testval, extent, myparms));
|
||||
}
|
||||
|
||||
for (auto &t : threadlist) {
|
||||
t.join();
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
|
||||
for (size_t k = 0 ; k < elements ; k++) {
|
||||
if (A_h[k] != testval) {
|
||||
CAPTURE(A_h[k], testval, k);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
free(A_h);
|
||||
HIP_CHECK(hipFree(devpPtr.ptr));
|
||||
}
|
||||
@@ -0,0 +1,193 @@
|
||||
/*
|
||||
* Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
* Test for checking order of execution of device kernel and
|
||||
* hipMemsetAsync apis on all gpus
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
|
||||
#define ITER 6
|
||||
#define N 1024 * 1024
|
||||
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
static unsigned blocks = 0;
|
||||
|
||||
|
||||
template <typename T>
|
||||
class MemSetKernelTest {
|
||||
public:
|
||||
T *A_h, *B_d, *B_h, *C_d;
|
||||
T memSetVal;
|
||||
size_t Nbytes;
|
||||
bool testResult = true;
|
||||
int validateCount = 0;
|
||||
hipStream_t stream;
|
||||
|
||||
void memAllocate(T memSetValue) {
|
||||
memSetVal = memSetValue;
|
||||
Nbytes = N * sizeof(T);
|
||||
|
||||
A_h = reinterpret_cast<T*>(malloc(Nbytes));
|
||||
HIP_ASSERT(A_h != nullptr);
|
||||
HIP_CHECK(hipMalloc(&B_d , Nbytes));
|
||||
B_h = reinterpret_cast<T*>(malloc(Nbytes));
|
||||
HIP_ASSERT(B_h != nullptr);
|
||||
HIP_CHECK(hipMalloc(&C_d , Nbytes));
|
||||
|
||||
for (int i = 0 ; i < N ; i++) {
|
||||
B_h[i] = i;
|
||||
}
|
||||
HIP_CHECK(hipMemcpy(B_d , B_h , Nbytes , hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
}
|
||||
|
||||
void memDeallocate() {
|
||||
HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d));
|
||||
free(B_h); free(A_h);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void validateExecutionOrder() {
|
||||
for (int p = 0 ; p < N ; p++) {
|
||||
if (A_h[p] == memSetVal) {
|
||||
validateCount+= 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
bool resultAfterAllIterations() {
|
||||
testResult = (validateCount == (ITER * N)) ? true : false;
|
||||
memDeallocate();
|
||||
return testResult;
|
||||
}
|
||||
};
|
||||
|
||||
static bool testhipMemsetAsyncWithKernel() {
|
||||
MemSetKernelTest<char> obj;
|
||||
constexpr char memsetval = 0x42;
|
||||
|
||||
obj.memAllocate(memsetval);
|
||||
for (int k = 0 ; k < ITER ; k++) {
|
||||
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
|
||||
HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream));
|
||||
HIP_CHECK(hipStreamSynchronize(obj.stream));
|
||||
HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
obj.validateExecutionOrder();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
static bool testhipMemsetD32AsyncWithKernel() {
|
||||
MemSetKernelTest <int32_t> obj;
|
||||
constexpr int memsetD32val = 0xDEADBEEF;
|
||||
|
||||
obj.memAllocate(memsetD32val);
|
||||
for (int k = 0 ; k < ITER ; k++) {
|
||||
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
|
||||
HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)obj.C_d , obj.memSetVal,
|
||||
N, obj.stream));
|
||||
HIP_CHECK(hipStreamSynchronize(obj.stream));
|
||||
HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
obj.validateExecutionOrder();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
static bool testhipMemsetD16AsyncWithKernel() {
|
||||
MemSetKernelTest <int16_t> obj;
|
||||
constexpr int16_t memsetD16val = 0xDEAD;
|
||||
|
||||
obj.memAllocate(memsetD16val);
|
||||
for (int k = 0 ; k < ITER ; k++) {
|
||||
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
|
||||
HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)obj.C_d , obj.memSetVal,
|
||||
N, obj.stream));
|
||||
HIP_CHECK(hipStreamSynchronize(obj.stream));
|
||||
HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
obj.validateExecutionOrder();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
static bool testhipMemsetD8AsyncWithKernel() {
|
||||
MemSetKernelTest <char> obj;
|
||||
constexpr char memsetD8val = 0xDE;
|
||||
|
||||
obj.memAllocate(memsetD8val);
|
||||
for (int k = 0; k < ITER; k++) {
|
||||
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
|
||||
HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)obj.C_d, obj.memSetVal,
|
||||
N, obj.stream));
|
||||
HIP_CHECK(hipStreamSynchronize(obj.stream));
|
||||
HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
obj.validateExecutionOrder();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
|
||||
/*
|
||||
* Test for checking order of execution of device kernel and
|
||||
* hipMemsetAsync apis on all gpus
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemsetAsync_VerifyExecutionWithKernel") {
|
||||
int numDevices = 0;
|
||||
bool ret;
|
||||
|
||||
blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HIP_CHECK(hipGetDeviceCount(&numDevices));
|
||||
REQUIRE(numDevices > 0);
|
||||
|
||||
auto devNum = GENERATE_COPY(range(0, numDevices));
|
||||
HIP_CHECK(hipSetDevice(devNum));
|
||||
|
||||
SECTION("hipMemsetAsync With Kernel") {
|
||||
ret = testhipMemsetAsyncWithKernel();
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD32Async With Kernel") {
|
||||
ret = testhipMemsetD32AsyncWithKernel();
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD16Async With Kernel") {
|
||||
ret = testhipMemsetD16AsyncWithKernel();
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD8Async With Kernel") {
|
||||
ret = testhipMemsetD8AsyncWithKernel();
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,243 @@
|
||||
/*
|
||||
* Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
* Test that validates functionality of hipmemsetAsync apis over multi threads
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
|
||||
|
||||
#define NUM_THREADS 20
|
||||
#define ITER 10
|
||||
#define N (4*1024*1024)
|
||||
|
||||
|
||||
template <typename T>
|
||||
class MemSetAsyncMthreadTest {
|
||||
public:
|
||||
T *A_h, *A_d, *B_h;
|
||||
T memSetVal;
|
||||
size_t Nbytes;
|
||||
bool testResult = true;
|
||||
int validateCount = 0;
|
||||
hipStream_t stream;
|
||||
|
||||
void memAllocate(T memSetValue) {
|
||||
memSetVal = memSetValue;
|
||||
Nbytes = N * sizeof(T);
|
||||
|
||||
A_h = reinterpret_cast<T*>(malloc(Nbytes));
|
||||
HIP_ASSERT(A_h != nullptr);
|
||||
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
B_h = reinterpret_cast<T*>(malloc(Nbytes));
|
||||
HIP_ASSERT(B_h != nullptr);
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
}
|
||||
|
||||
void threadCompleteStatus() {
|
||||
for (int k = 0 ; k < N ; k++) {
|
||||
if ((A_h[k] == memSetVal) && (B_h[k] == memSetVal)) {
|
||||
validateCount+= 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
bool resultAfterAllIterations() {
|
||||
memDeallocate();
|
||||
testResult = (validateCount == (ITER * N)) ? true: false;
|
||||
return testResult;
|
||||
}
|
||||
|
||||
void memDeallocate() {
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
free(A_h);
|
||||
free(B_h);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
void queueJobsForhipMemsetAsync(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
|
||||
hipStream_t stream) {
|
||||
HIPCHECK(hipMemsetAsync(A_d, memSetVal, N, stream));
|
||||
HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void queueJobsForhipMemsetD32Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
|
||||
hipStream_t stream) {
|
||||
HIPCHECK(hipMemsetD32Async((hipDeviceptr_t)A_d, memSetVal, N, stream));
|
||||
HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void queueJobsForhipMemsetD16Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
|
||||
hipStream_t stream) {
|
||||
HIPCHECK(hipMemsetD16Async((hipDeviceptr_t)A_d, memSetVal, N, stream));
|
||||
HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void queueJobsForhipMemsetD8Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
|
||||
hipStream_t stream) {
|
||||
HIPCHECK(hipMemsetD8Async((hipDeviceptr_t)A_d, memSetVal, N, stream));
|
||||
HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
/* Queue hipMemsetAsync jobs on multiple threads and verify they all
|
||||
* finished on all threads successfully
|
||||
*/
|
||||
bool testhipMemsetAsyncWithMultiThread() {
|
||||
MemSetAsyncMthreadTest <char> obj;
|
||||
constexpr char memsetval = 0x42;
|
||||
obj.memAllocate(memsetval);
|
||||
std::thread t[NUM_THREADS];
|
||||
|
||||
for (int i = 0 ; i < ITER ; i++) {
|
||||
for (int k = 0 ; k < NUM_THREADS ; k++) {
|
||||
if (k%2) {
|
||||
t[k] = std::thread(queueJobsForhipMemsetAsync<char>, obj.A_d, obj.A_h,
|
||||
obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
} else {
|
||||
t[k] = std::thread(queueJobsForhipMemsetAsync<char>, obj.A_d, obj.B_h,
|
||||
obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
}
|
||||
}
|
||||
|
||||
for (int j = 0 ; j < NUM_THREADS ; j++) {
|
||||
t[j].join();
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(obj.stream));
|
||||
obj.threadCompleteStatus();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
bool testhipMemsetD32AsyncWithMultiThread() {
|
||||
MemSetAsyncMthreadTest <int32_t> obj;
|
||||
constexpr int memsetD32val = 0xDEADBEEF;
|
||||
obj.memAllocate(memsetD32val);
|
||||
std::thread t[NUM_THREADS];
|
||||
|
||||
for (int i = 0 ; i < ITER ; i++) {
|
||||
for (int k = 0 ; k < NUM_THREADS ; k++) {
|
||||
if (k%2) {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD32Async<int32_t>, obj.A_d,
|
||||
obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
} else {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD32Async<int32_t>, obj.A_d,
|
||||
obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
}
|
||||
}
|
||||
|
||||
for (int j = 0 ; j < NUM_THREADS ; j++) {
|
||||
t[j].join();
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(obj.stream));
|
||||
obj.threadCompleteStatus();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
bool testhipMemsetD16AsyncWithMultiThread() {
|
||||
MemSetAsyncMthreadTest <int16_t> obj;
|
||||
constexpr int16_t memsetD16val = 0xDEAD;
|
||||
obj.memAllocate(memsetD16val);
|
||||
std::thread t[NUM_THREADS];
|
||||
|
||||
for (int i = 0 ; i < ITER ; i++) {
|
||||
for (int k = 0 ; k < NUM_THREADS ; k++) {
|
||||
if (k%2) {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD16Async<int16_t>, obj.A_d,
|
||||
obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
} else {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD16Async<int16_t>, obj.A_d,
|
||||
obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
}
|
||||
}
|
||||
|
||||
for (int j = 0 ; j < NUM_THREADS ; j++) {
|
||||
t[j].join();
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(obj.stream));
|
||||
obj.threadCompleteStatus();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
bool testhipMemsetD8AsyncWithMultiThread() {
|
||||
MemSetAsyncMthreadTest <char> obj;
|
||||
constexpr char memsetD8val = 0xDE;
|
||||
obj.memAllocate(memsetD8val);
|
||||
std::thread t[NUM_THREADS];
|
||||
|
||||
for (int i = 0 ; i < ITER ; i++) {
|
||||
for (int k = 0 ; k < NUM_THREADS ; k++) {
|
||||
if (k%2) {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD8Async<char>, obj.A_d,
|
||||
obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
} else {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD8Async<char>, obj.A_d,
|
||||
obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
}
|
||||
}
|
||||
for (int j = 0 ; j < NUM_THREADS ; j++) {
|
||||
t[j].join();
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(obj.stream));
|
||||
obj.threadCompleteStatus();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
|
||||
/*
|
||||
* Test that validates functionality of hipmemsetAsync apis over multi threads
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemsetAsync_QueueJobsMultithreaded") {
|
||||
bool ret;
|
||||
|
||||
SECTION("hipMemsetAsync With MultiThread") {
|
||||
ret = testhipMemsetAsyncWithMultiThread();
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD32Async With MultiThread") {
|
||||
ret = testhipMemsetD32AsyncWithMultiThread();
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD16Async With MultiThread") {
|
||||
ret = testhipMemsetD16AsyncWithMultiThread();
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD8Async With MultiThread") {
|
||||
ret = testhipMemsetD8AsyncWithMultiThread();
|
||||
REQUIRE(ret == true);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,127 @@
|
||||
/*
|
||||
* Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Test hipMemset apis with invalid pointer and invalid 2D pitch.
|
||||
2) Test hipMemsetAsync apis with invalid pointer and invalid 2D pitch.
|
||||
*/
|
||||
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#define N 50
|
||||
#define MEMSETVAL 0x42
|
||||
|
||||
/**
|
||||
* Testcase validates hipMemset apis behavior with
|
||||
* invalid pointer and invalid 2D pitch value.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemset_InvalidPtrTests") {
|
||||
hipError_t ret;
|
||||
constexpr int Nbytes = N*sizeof(char);
|
||||
char *A_d;
|
||||
|
||||
SECTION("hipMemset with null") {
|
||||
ret = hipMemset(NULL, MEMSETVAL , Nbytes);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("hipMemset with hostptr") {
|
||||
char *A_h;
|
||||
A_h = reinterpret_cast<char *>(malloc(Nbytes));
|
||||
|
||||
ret = hipMemset(A_h, MEMSETVAL, Nbytes);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD32 with null") {
|
||||
ret = hipMemsetD32(NULL, MEMSETVAL , Nbytes);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD16 with null") {
|
||||
ret = hipMemsetD16(NULL, MEMSETVAL , Nbytes);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD8 with null") {
|
||||
ret = hipMemsetD8(NULL, MEMSETVAL , Nbytes);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("hipMemset2D with null") {
|
||||
constexpr size_t NUM_H = 256, NUM_W = 256;
|
||||
size_t pitch_A;
|
||||
size_t width = NUM_W * sizeof(char);
|
||||
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
|
||||
width , NUM_H));
|
||||
ret = hipMemset2D(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
|
||||
hipFree(A_d);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Testcase validates hipMemsetAsync apis behavior with
|
||||
* invalid pointer and invalid 2D pitch value.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMemsetAsync_InvalidPtrTests") {
|
||||
hipError_t ret;
|
||||
constexpr int Nbytes = N*sizeof(char);
|
||||
char *A_d;
|
||||
|
||||
SECTION("hipMemsetAsync with null") {
|
||||
ret = hipMemsetAsync(NULL, MEMSETVAL, Nbytes , 0);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD32Async with null") {
|
||||
ret = hipMemsetD32Async(NULL, MEMSETVAL , Nbytes, 0);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD16Async with null") {
|
||||
ret = hipMemsetD16Async(NULL, MEMSETVAL , Nbytes, 0);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("hipMemsetD8Async with null") {
|
||||
ret = hipMemsetD8Async(NULL, MEMSETVAL , Nbytes, 0);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("hipMemset2DAsync with null") {
|
||||
constexpr size_t NUM_H = 256, NUM_W = 256;
|
||||
size_t pitch_A;
|
||||
size_t width = NUM_W * sizeof(char);
|
||||
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
|
||||
width , NUM_H));
|
||||
ret = hipMemset2DAsync(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H, 0);
|
||||
REQUIRE(ret != hipSuccess);
|
||||
|
||||
hipFree(A_d);
|
||||
}
|
||||
}
|
||||
@@ -9,6 +9,8 @@ set(TEST_SRC
|
||||
hipStreamCreateWithFlags.cc
|
||||
hipStreamCreateWithPriority.cc
|
||||
hipStreamWithCUMask.cc
|
||||
hipStreamGetCUMask.cc
|
||||
hipAPIStreamDisable.cc
|
||||
)
|
||||
else()
|
||||
set(TEST_SRC
|
||||
@@ -20,6 +22,7 @@ set(TEST_SRC
|
||||
hipStreamAddCallback.cc
|
||||
hipStreamCreateWithFlags.cc
|
||||
hipStreamCreateWithPriority.cc
|
||||
hipAPIStreamDisable.cc
|
||||
)
|
||||
endif()
|
||||
|
||||
|
||||
@@ -0,0 +1,68 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
#include <hip_test_common.hh>
|
||||
#include "hip/math_functions.h"
|
||||
|
||||
#define NUM_STREAMS 8
|
||||
|
||||
namespace hipAPIStreamDisableTest {
|
||||
const int NN = 1 << 21;
|
||||
|
||||
__global__ void kernel(float* x, float* y, int n) {
|
||||
int tid = threadIdx.x;
|
||||
if (tid < 1) {
|
||||
for (int i = 0; i < n; i++) {
|
||||
x[i] = sqrt(powf(3.14159, i));
|
||||
}
|
||||
y[tid] = y[tid] + 1.0f;
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void nKernel(float* y) {
|
||||
int tid = threadIdx.x;
|
||||
y[tid] = y[tid] + 1.0f;
|
||||
}
|
||||
} // namespace hipAPIStreamDisableTest
|
||||
|
||||
/**
|
||||
* Validate basic multistream functionalities
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamCreate_MultistreamBasicFunctionalities") {
|
||||
hipStream_t streams[NUM_STREAMS];
|
||||
float *data[NUM_STREAMS], *yd, *xd;
|
||||
float y = 1.0f, x = 1.0f;
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&yd), sizeof(float)));
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&xd), sizeof(float)));
|
||||
HIP_CHECK(hipMemcpy(yd, &y, sizeof(float), hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(xd, &x, sizeof(float), hipMemcpyHostToDevice));
|
||||
|
||||
for (int i = 0; i < NUM_STREAMS; i++) {
|
||||
HIP_CHECK(hipStreamCreate(&streams[i]));
|
||||
HIP_CHECK(hipMalloc(&data[i],
|
||||
(hipAPIStreamDisableTest::NN)*sizeof(float)));
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipAPIStreamDisableTest::kernel),
|
||||
dim3(1), dim3(1), 0, streams[i], data[i], xd,
|
||||
hipAPIStreamDisableTest::NN);
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipAPIStreamDisableTest::nKernel),
|
||||
dim3(1), dim3(1), 0, 0, yd);
|
||||
}
|
||||
HIP_CHECK(hipMemcpy(&x, xd, sizeof(float), hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipMemcpy(&y, yd, sizeof(float), hipMemcpyDeviceToHost));
|
||||
REQUIRE(x == y);
|
||||
}
|
||||
@@ -0,0 +1,177 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Test to verify hipExtStreamGetCUMask api returning default CU Mask or global CU Mask.
|
||||
2) Test to verify hipExtStreamGetCUMask api returns custom mask set.
|
||||
3) Negative tests for hipExtStreamGetCUMask api.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <vector>
|
||||
|
||||
|
||||
/**
|
||||
* Scenario to verify hipExtStreamGetCUMask api returning default CU Mask or global CU Mask.
|
||||
* Scenario to verify hipExtStreamGetCUMask api returns custom mask set.
|
||||
*/
|
||||
TEST_CASE("Unit_hipExtStreamGetCUMask_verifyDefaultAndCustomMask") {
|
||||
constexpr int maxNum = 6;
|
||||
std::vector<uint32_t> cuMask(maxNum);
|
||||
hipDeviceProp_t props;
|
||||
std::stringstream ss;
|
||||
char* gCUMask{nullptr};
|
||||
std::string globalCUMask("");
|
||||
std::vector<uint32_t> defaultCUMask;
|
||||
|
||||
int nGpu = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&nGpu));
|
||||
if (nGpu < 1) {
|
||||
INFO("info: didn't find any GPU! skipping the test!");
|
||||
return;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipSetDevice(0));
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, 0));
|
||||
INFO("info: running on bus " << "0x" << props.pciBusID << " " <<
|
||||
props.name << " with " << props.multiProcessorCount << " CUs");
|
||||
|
||||
// Get global CU Mask if exists
|
||||
gCUMask = getenv("ROC_GLOBAL_CU_MASK");
|
||||
if (gCUMask != nullptr && gCUMask[0] != '\0') {
|
||||
globalCUMask.assign(gCUMask);
|
||||
|
||||
for_each(globalCUMask.begin(), globalCUMask.end(), [](char & c) {
|
||||
c = ::tolower(c);
|
||||
});
|
||||
}
|
||||
|
||||
// Create default CU Mask
|
||||
uint32_t temp = 0;
|
||||
uint32_t bit_index = 0;
|
||||
for (uint32_t i = 0; i < (uint32_t)props.multiProcessorCount; i++) {
|
||||
temp |= 1UL << bit_index;
|
||||
if (bit_index >= 32) {
|
||||
defaultCUMask.push_back(temp);
|
||||
temp = 0;
|
||||
bit_index = 0;
|
||||
temp |= 1UL << bit_index;
|
||||
}
|
||||
bit_index += 1;
|
||||
}
|
||||
if (bit_index != 0) {
|
||||
defaultCUMask.push_back(temp);
|
||||
}
|
||||
|
||||
SECTION("Verify with default CU Mask or global CU Mask") {
|
||||
// make a default CU mask bit-array where all CUs are active
|
||||
// this default mask is expected to be returned when there is no
|
||||
// custom or global CU mask defined
|
||||
|
||||
HIP_CHECK(hipExtStreamGetCUMask(0, cuMask.size(), &cuMask[0]));
|
||||
|
||||
ss << std::hex;
|
||||
for (int i = cuMask.size() - 1; i >= 0; i--) {
|
||||
ss << cuMask[i];
|
||||
}
|
||||
|
||||
// remove extra 0 from ss if any present
|
||||
size_t found = ss.str().find_first_not_of("0");
|
||||
if (found != std::string::npos) {
|
||||
ss.str(ss.str().substr(found, ss.str().length()));
|
||||
}
|
||||
|
||||
INFO("info: CU mask for the default stream is: 0x" << ss.str().c_str());
|
||||
if (globalCUMask.size() > 0) {
|
||||
if (ss.str().compare(globalCUMask) != 0) {
|
||||
INFO("Expected CU mask:" << globalCUMask.c_str() <<
|
||||
", api returned:" << ss.str().c_str());
|
||||
REQUIRE(false);
|
||||
}
|
||||
} else {
|
||||
for (int i = 0 ; i < min(cuMask.size(), defaultCUMask.size()); i++) {
|
||||
if (cuMask[i] != defaultCUMask[i]) {
|
||||
INFO("Expected CU mask " << defaultCUMask[i] <<
|
||||
", api returned:" << cuMask[i]);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
SECTION("Verify with custom mask set") {
|
||||
std::vector<uint32_t> customMask(defaultCUMask);
|
||||
hipStream_t stream;
|
||||
customMask[0] = 0xe;
|
||||
|
||||
HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, customMask.size(),
|
||||
customMask.data()));
|
||||
ss.str("");
|
||||
for (int i = customMask.size() - 1; i >= 0; i--) {
|
||||
ss << customMask[i];
|
||||
}
|
||||
INFO("info: setting a custom CU mask 0x" << ss.str());
|
||||
|
||||
HIP_CHECK(hipExtStreamGetCUMask(stream, cuMask.size(), &cuMask[0]));
|
||||
ss.str("");
|
||||
for (int i = cuMask.size() - 1; i >= 0; i--) {
|
||||
ss << cuMask[i];
|
||||
}
|
||||
|
||||
size_t found = ss.str().find_first_not_of("0");
|
||||
if (found != std::string::npos) {
|
||||
ss.str(ss.str().substr(found, ss.str().length()));
|
||||
}
|
||||
|
||||
INFO("info: reading back CU mask 0x" << ss.str() <<
|
||||
" for stream " << stream);
|
||||
|
||||
if (!gCUMask) {
|
||||
for (size_t i = 0; i < customMask.size(); i++) {
|
||||
if (customMask[i] != cuMask[i]) {
|
||||
INFO("Error! expected CU mask:" << customMask[i]
|
||||
<< ", Received CU mask:" << cuMask[i]);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Negative tests for hipExtStreamGetCUMask.
|
||||
*/
|
||||
TEST_CASE("Unit_hipExtStreamGetCUMask_Negative") {
|
||||
hipError_t ret;
|
||||
constexpr int maxNum = 6;
|
||||
std::vector<uint32_t> cuMask(maxNum);
|
||||
|
||||
SECTION("cuMask is nullptr") {
|
||||
ret = hipExtStreamGetCUMask(0, cuMask.size(), nullptr);
|
||||
REQUIRE(ret == hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("cuMaskSize is 0") {
|
||||
ret = hipExtStreamGetCUMask(0, 0, &cuMask[0]);
|
||||
REQUIRE(ret == hipErrorInvalidValue);
|
||||
}
|
||||
}
|
||||
@@ -16,16 +16,24 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Test flag value of stream created with hipStreamCreateWithFlags/
|
||||
/hipStreamCreate/hipStreamCreateWithPriority.
|
||||
2) Negative tests for hipStreamGetFlags api.
|
||||
3) Test flag value when streams created with CUMask.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
TEST_CASE("Unit_hipStreamGetFlags_Negative") {
|
||||
// Get flags for uninitialized stream
|
||||
|
||||
/**
|
||||
* Test flag value of stream created with various types.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetFlags_BasicFunctionalities") {
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
|
||||
REQUIRE(hipStreamGetFlags(stream, nullptr) == hipErrorInvalidValue);
|
||||
}
|
||||
TEST_CASE("Unit_hipStreamGetFlags") {
|
||||
hipStream_t stream;
|
||||
unsigned int flags;
|
||||
unsigned int flags;
|
||||
// Check flag value of stream created with hipStreamCreateWithFlags
|
||||
SECTION("Check flag value of streams hipStreamCreateWithFlags") {
|
||||
HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
|
||||
HIP_CHECK(hipStreamGetFlags(stream, &flags));
|
||||
REQUIRE(flags == hipStreamDefault);
|
||||
@@ -34,4 +42,50 @@ TEST_CASE("Unit_hipStreamGetFlags") {
|
||||
HIP_CHECK(hipStreamGetFlags(stream, &flags));
|
||||
REQUIRE(flags == hipStreamNonBlocking);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
// Check flag value of stream created with hipStreamCreate
|
||||
SECTION("Check flag value of streams hipStreamCreate") {
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipStreamGetFlags(stream, &flags));
|
||||
REQUIRE(flags == hipStreamDefault);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
// Check flag value of stream created with hipStreamCreateWithPriority
|
||||
SECTION("Check flag value of streams hipStreamCreateWithPriority") {
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, 0));
|
||||
HIP_CHECK(hipStreamGetFlags(stream, &flags));
|
||||
REQUIRE(flags == hipStreamDefault);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, 0));
|
||||
HIP_CHECK(hipStreamGetFlags(stream, &flags));
|
||||
REQUIRE(flags == hipStreamNonBlocking);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Negative Scenarios
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetFlags_Negative") {
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
// nullptr check
|
||||
REQUIRE(hipStreamGetFlags(stream, nullptr) != hipSuccess);
|
||||
REQUIRE(hipStreamGetFlags(nullptr, hipStreamDefault) != hipSuccess);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
#if HT_AMD
|
||||
/**
|
||||
* Test flag value when streams created with CUMask.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetFlags_StreamsCreatedWithCUMask") {
|
||||
hipStream_t stream;
|
||||
unsigned int flags;
|
||||
const uint32_t cuMask = 0xffffffff;
|
||||
HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, 1, &cuMask));
|
||||
HIP_CHECK(hipStreamGetFlags(stream, &flags));
|
||||
REQUIRE(flags == hipStreamDefault);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
#endif
|
||||
|
||||
@@ -16,11 +16,28 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Negative tests for hipStreamGetPriority api.
|
||||
2) Create stream and check default priority of stream is within range.
|
||||
3) Create stream with high or low priority and check priority is set as expected.
|
||||
4) Create stream with higher priority or lower priority for the priority range returned.
|
||||
5) Create stream with CUMask and check priority is returned as expected.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
/**
|
||||
* Negative tests for hipStreamGetPriority api.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetPriority_Negative") {
|
||||
hipStream_t stream = 0;
|
||||
REQUIRE(hipStreamGetPriority(stream, nullptr) == hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
/**
|
||||
* Create stream and check default priority of stream is within range.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetPriority_default") {
|
||||
int priority_low = 0;
|
||||
int priority_high = 0;
|
||||
@@ -37,6 +54,10 @@ TEST_CASE("Unit_hipStreamGetPriority_default") {
|
||||
REQUIRE(priority >= priority_high);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
/**
|
||||
* Create stream with high priority and check priority is set as expected.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetPriority_high") {
|
||||
int priority_low = 0;
|
||||
int priority_high = 0;
|
||||
@@ -44,12 +65,17 @@ TEST_CASE("Unit_hipStreamGetPriority_high") {
|
||||
HIP_CHECK(hipSetDevice(devID));
|
||||
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high));
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
|
||||
priority_high));
|
||||
int priority = 0;
|
||||
HIP_CHECK(hipStreamGetPriority(stream, &priority));
|
||||
REQUIRE(priority == priority_high);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
/**
|
||||
* Create stream with higher priority for the priority range returned.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetPriority_higher") {
|
||||
int priority_low = 0;
|
||||
int priority_high = 0;
|
||||
@@ -57,12 +83,17 @@ TEST_CASE("Unit_hipStreamGetPriority_higher") {
|
||||
HIP_CHECK(hipSetDevice(devID));
|
||||
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high-1));
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
|
||||
priority_high-1));
|
||||
int priority = 0;
|
||||
HIP_CHECK(hipStreamGetPriority(stream, &priority));
|
||||
REQUIRE(priority == priority_high);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
/**
|
||||
* Create stream with low priority and check priority is set as expected.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetPriority_low") {
|
||||
int priority_low = 0;
|
||||
int priority_high = 0;
|
||||
@@ -70,12 +101,17 @@ TEST_CASE("Unit_hipStreamGetPriority_low") {
|
||||
HIP_CHECK(hipSetDevice(devID));
|
||||
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low));
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
|
||||
priority_low));
|
||||
int priority = 0;
|
||||
HIP_CHECK(hipStreamGetPriority(stream, &priority));
|
||||
REQUIRE(priority_low == priority);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
/**
|
||||
* Create stream with lower priority for the priority range returned.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetPriority_lower") {
|
||||
int priority_low = 0;
|
||||
int priority_high = 0;
|
||||
@@ -83,9 +119,35 @@ TEST_CASE("Unit_hipStreamGetPriority_lower") {
|
||||
HIP_CHECK(hipSetDevice(devID));
|
||||
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_low+1));
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
|
||||
priority_low+1));
|
||||
int priority = 0;
|
||||
HIP_CHECK(hipStreamGetPriority(stream, &priority));
|
||||
REQUIRE(priority_low == priority);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
#if HT_AMD
|
||||
/**
|
||||
* Create stream with CUMask and check priority is returned as expected.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetPriority_StreamsWithCUMask") {
|
||||
hipStream_t stream;
|
||||
int priority = 0;
|
||||
int priority_normal = 0;
|
||||
int priority_low = 0;
|
||||
int priority_high = 0;
|
||||
// Test is to get the Stream Priority Range
|
||||
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
|
||||
priority_normal = priority_low + priority_high;
|
||||
// Check if priorities are indeed supported
|
||||
REQUIRE_FALSE(priority_low == priority_high);
|
||||
// Creating a stream with hipExtStreamCreateWithCUMask and checking
|
||||
// priority.
|
||||
const uint32_t cuMask = 0xffffffff;
|
||||
HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, 1, &cuMask));
|
||||
HIP_CHECK(hipStreamGetPriority(stream, &priority));
|
||||
REQUIRE_FALSE(priority_normal != priority);
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
#endif
|
||||
|
||||
Atsaukties uz šo jaunā problēmā
Block a user