Merge 'develop' into 'amd-staging'

Change-Id: Ic259c16363ab4e869531643c7cce211b2c6c550d


[ROCm/hip-tests commit: 070da7dc56]
This commit is contained in:
Jenkins
2023-06-26 23:12:03 +00:00
57 changed files with 4908 additions and 1129 deletions
@@ -14,9 +14,7 @@
"Unit_hipDeviceReset_Positive_Basic",
"Unit_hipDeviceReset_Positive_Threaded",
"Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic",
"Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic",
"Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Positive_Basic",
"Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Positive_Basic",
"Unit_hipKernelNameRef_Negative_Parameters",
"Unit_hipMemAdvise_AccessedBy_All_Devices",
"Unit_hipMemAdvise_No_Flag_Interference",
@@ -31,7 +29,6 @@
"Unit_hipStreamAttachMemAsync_Positive_AttachGlobal",
"Unit_hipStreamAttachMemAsync_Negative_Parameters",
"Unit_hipMemGetAddressRange_Positive",
"Unit_hipGraphAddMemcpyNode1D_Negative_Basic",
"Unit_hipStreamGetCaptureInfo_Nullstream_CaptureInfo",
"intermittent issue: corrupted double-linked list",
"Unit_hipGraphRetainUserObject_Functional_2",
@@ -44,10 +41,6 @@
"Unit_hipStreamCreateWithFlags_DefaultStreamInteraction",
"Unit_hipStreamWaitEvent_UninitializedStream_Negative",
"Unit_hipDeviceSetSharedMemConfig_Negative_Parameters",
"Disabling tests tracked with SWDEV-394083",
"Unit_hipDeviceSynchronize_Positive_Nullstream",
"Disabling tests tracked with SWDEV-393637",
"Unit_hipDeviceSynchronize_Functional",
"Unit_hipMemset3DSync",
"Unit_hipStreamAddCallback_StrmSyncTiming",
"Disabling test tracked SWDEV-394199",
@@ -71,10 +64,22 @@
"Unit_hipMemcpyPeer_Positive_Synchronization_Behavior",
"SWDEV-398981 fails in stress test",
"Unit_hipStreamCreateWithPriority_MulthreadDefaultflag",
"Disabling below 4 tests temporarily due to change in API behavior",
"Unit_hipDeviceSetCacheConfig_Positive_Basic",
"Unit_hipDeviceSetCacheConfig_Negative_Parameters",
"Unit_hipDeviceSetSharedMemConfig_Positive_Basic",
"Unit_hipDeviceSetSharedMemConfig_Negative_Parameters"
"=== Below tests fail in stress test on 23/06/23 ===",
"Unit_hipIpcMemAccess_ParameterValidation",
"Unit_hipMemcpy2DFromArrayAsync_Positive_Synchronization_Behavior",
"Unit_hipMalloc3D_SmallandBigChunks",
"Unit_hipMalloc3D_MultiThread",
"Unit_hipArrayCreate_DiffSizes",
"Unit_hipArrayCreate_MultiThread",
"hipMemGetInfo_DifferentMallocSmall",
"Unit_hipMemGetInfo_ParaSmall",
"Unit_hipMemGetInfo_ParaNonDiv",
"Unit_hipMemGetInfo_ParaMultiSmall",
"Unit_hipMemGetInfo_Negative",
"Unit_hipMemsetDSync - int8_t",
"Unit_hipGraphClone_Test_hipGraphExecMemcpyNodeSetParams",
"Unit_hipGraphClone_Test_hipGraphMemcpyNodeSetParams1D_and_exec",
"Unit_hipStreamValue_Wait64_Blocking_NoMask_And",
"Unit_hipStreamValue_Wait64_Blocking_NoMask_Nor"
]
}
@@ -166,11 +166,6 @@
"SWDEV-402082 - PAL Backend fails to reserve address on GPU except first one",
"Unit_hipGraphInstantiateWithFlags_FlagAutoFreeOnLaunch_check",
"SWDEV-398981 fails in stress test",
"Unit_hipStreamCreateWithPriority_MulthreadDefaultflag",
"Disabling below 4 tests temporarily due to change in API behavior",
"Unit_hipDeviceSetCacheConfig_Positive_Basic",
"Unit_hipDeviceSetCacheConfig_Negative_Parameters",
"Unit_hipDeviceSetSharedMemConfig_Positive_Basic",
"Unit_hipDeviceSetSharedMemConfig_Negative_Parameters"
"Unit_hipStreamCreateWithPriority_MulthreadDefaultflag"
]
}
@@ -81,4 +81,136 @@ struct vector_info<ushort4>
: type_and_size_and_format<unsigned short, 4, HIP_AD_FORMAT_UNSIGNED_INT16> {};
template <>
struct vector_info<uchar4>
: type_and_size_and_format<unsigned char, 4, HIP_AD_FORMAT_UNSIGNED_INT8> {};
: type_and_size_and_format<unsigned char, 4, HIP_AD_FORMAT_UNSIGNED_INT8> {};
template <
typename T,
typename std::enable_if<std::is_scalar<T>::value == false>::type* = nullptr>
static inline __host__ __device__ constexpr int rank() {
return sizeof(T) / sizeof(decltype(T::x));
}
template<
typename T,
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
static inline bool isEqual(const T &val0, const T &val1) {
return val0.x == val1.x;
}
template<
typename T,
typename std::enable_if<rank<T>() == 2>::type* = nullptr>
static inline bool isEqual(const T &val0, const T &val1) {
return val0.x == val1.x &&
val0.y == val1.y;
}
template<
typename T,
typename std::enable_if<rank<T>() == 4>::type* = nullptr>
static inline bool isEqual(const T &val0, const T &val1) {
return val0.x == val1.x &&
val0.y == val1.y &&
val0.z == val1.z &&
val0.w == val1.w;
}
template<
typename T,
typename std::enable_if<std::is_scalar<T>::value>::type* = nullptr>
static inline bool isEqual(const T &val0, const T &val1) {
return val0 == val1;
}
template<
typename T,
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
const std::string getString(const T& t)
{
std::ostringstream os;
os<< "(" << t.x << ")";
return os.str();
}
template<
typename T,
typename std::enable_if<rank<T>() == 2>::type* = nullptr>
const std::string getString(const T& t)
{
std::ostringstream os;
os<< "(" << t.x << ", " << t.y << ")";
return os.str();
}
template<
typename T,
typename std::enable_if<rank<T>() == 3>::type* = nullptr>
const std::string getString(const T& t)
{
std::ostringstream os;
os<< "(" << t.x << ", " << t.y << ", " << t.z << ")";
return os.str();
}
template<
typename T,
typename std::enable_if<rank<T>() == 4>::type* = nullptr>
const std::string getString(const T& t)
{
std::ostringstream os;
os<< "(" << t.x << ", " << t.y << ", " << t.z << ", " << t.w << ")";
return os.str();
}
template<
typename T,
typename std::enable_if<std::is_scalar<T>::value>::type* = nullptr>
std::string getString(const T& t)
{
std::ostringstream os;
os << t;
return os.str();
}
template<typename T>
static inline T getRandom() {
double r = 0;
if (std::is_signed<T>::value) {
r = (std::rand() - RAND_MAX / 2.0) / (RAND_MAX / 2.0 + 1.);
} else {
r = std::rand() / (RAND_MAX + 1.);
}
return static_cast<T>(std::numeric_limits < T > ::max() * r);
}
template<
typename T,
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
static inline void initVal(T &val) {
val.x = getRandom<decltype(T::x)>();
}
template<
typename T,
typename std::enable_if<rank<T>() == 2>::type* = nullptr>
static inline void initVal(T &val) {
val.x = getRandom<decltype(T::x)>();
val.y = getRandom<decltype(T::x)>();
}
template<
typename T,
typename std::enable_if<rank<T>() == 4>::type* = nullptr>
static inline void initVal(T &val) {
val.x = getRandom<decltype(T::x)>();
val.y = getRandom<decltype(T::x)>();
val.z = getRandom<decltype(T::x)>();
val.w = getRandom<decltype(T::x)>();
}
template<
typename T,
typename std::enable_if<std::is_scalar<T>::value>::type* = nullptr>
static inline void initVal(T &val) {
val = getRandom<T>();
}
@@ -22,6 +22,7 @@ THE SOFTWARE.
#pragma once
#include "hip_test_context.hh"
#include <catch.hpp>
#include <atomic>
#include <chrono>
@@ -30,6 +31,7 @@ THE SOFTWARE.
#include <iomanip>
#include <mutex>
#include <cstdlib>
#include <thread>
#define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__);
@@ -429,6 +431,52 @@ static inline void runKernelForDuration(std::chrono::milliseconds duration,
hipLaunchKernelGGL(waitKernel_used, dim3(1), dim3(1), 0, stream, ticksPerSecond * millis / 1000);
}
class BlockingContext {
std::atomic_bool blocked{true};
hipStream_t stream;
public:
BlockingContext(hipStream_t s) : stream(s), blocked(true) {}
BlockingContext(const BlockingContext& in) {
blocked = in.blocked_val();
stream = in.stream_val();
}
BlockingContext(const BlockingContext&& in) {
blocked = in.blocked_val();
stream = in.stream_val();
}
void reset() { blocked = true; }
BlockingContext& operator=(const BlockingContext& in) {
blocked = in.blocked_val();
stream = in.stream_val();
return *this;
}
void block_stream() {
blocked = true;
auto blocking_callback = [](hipStream_t, hipError_t, void* data) {
auto blocked = reinterpret_cast<std::atomic_bool*>(data);
while (blocked->load()) {
// Yield this thread till we are waiting
std::this_thread::yield();
}
};
HIP_CHECK(hipStreamAddCallback(stream, blocking_callback, (void*)&blocked, 0));
}
void unblock_stream() {
blocked = false;
}
bool is_blocked() const { return hipStreamQuery(stream) == hipErrorNotReady; }
bool blocked_val() const { return blocked.load(); }
hipStream_t stream_val() const { return stream; }
};
} // namespace HipTest
// This must be called in the beginning of image test app's main() to indicate whether image
@@ -42,3 +42,10 @@ THE SOFTWARE.
* This section describes the warp shuffle types & functions of HIP runtime API.
* @}
*/
/**
* @defgroup StreamTest Stream Management
* @{
* This section describes the stream management types & functions of HIP runtime API.
* @}
*/
@@ -8,6 +8,8 @@
#define HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD 0.1
#if HT_NVIDIA
typedef unsigned char uchar;
template<typename T>
typename std::enable_if<sizeof(T) / sizeof(decltype(T::x)) == 4, T>::type
inline __host__ __device__ operator+(const T &a, const T &b)
@@ -47,7 +49,6 @@ inline __host__ __device__ operator*=(T &a, const decltype(T::x) &b)
}
#endif // HT_NVIDIA
// See https://en.wikipedia.org/wiki/SRGB#Transformation
// From CIE 1931 color space to sRGB
inline float hipSRGBMap(float fc) {
double c = static_cast<double>(fc);
@@ -134,7 +134,6 @@ inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hi
HIPCHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0));
}
Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms);
HIP_CHECK(hipGetLastError());
}
template <typename... Attributes>
@@ -28,6 +28,7 @@ add_subdirectory(occupancy)
add_subdirectory(device)
add_subdirectory(printf)
add_subdirectory(texture)
add_subdirectory(surface)
add_subdirectory(streamperthread)
add_subdirectory(kernel)
add_subdirectory(multiThread)
@@ -6,6 +6,9 @@ set(TEST_SRC
hipCGMultiGridGroupType.cc
hipCGMultiGridGroupTypeViaBaseType.cc
hipCGMultiGridGroupTypeViaPublicApi.cc
coalesced_groups_shfl_down.cc
coalesced_groups_shfl_up.cc
simple_coalesced_groups.cc
)
if(HIP_PLATFORM STREQUAL "nvidia")
set_source_files_properties(hipCGMultiGridGroupType.cc PROPERTIES COMPILE_FLAGS "-rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80")
@@ -66,7 +66,7 @@ __global__ void kernel_shfl_down (int * dPtr, int *dResults, int lane_delta, int
}
}
__global__ void kernel_cg_group_partition(int* result, unsigned int tileSz, int cg_sizes) {
__global__ void kernel_cg_group_partition_shfl_down(int* result, unsigned int tileSz, int cg_sizes) {
int id = threadIdx.x + blockIdx.x * blockDim.x;
if (id % cg_sizes == 0) {
@@ -82,7 +82,6 @@ __global__ void kernel_cg_group_partition(int* result, unsigned int tileSz, int
threadBlockCGTy.sync();
coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz);
int threadRank = tiledPartition.thread_rank();
input = tiledPartition.thread_rank();
@@ -110,7 +109,7 @@ void verifyResults(int* ptr, int expectedResult, int numTiles) {
}
}
void compareResults(int* cpu, int* gpu, int size) {
void compareResultsCoalescedGroupsShflDown(int* cpu, int* gpu, int size) {
for (unsigned int i = 0; i < size / sizeof(int); i++) {
if (cpu[i] != gpu[i]) {
INFO(" results do not match.");
@@ -118,7 +117,7 @@ void compareResults(int* cpu, int* gpu, int size) {
}
}
void printResults(int* ptr, int size) {
void printResultsCoalescedGroupsShflDown(int* ptr, int size) {
for (int i = 0; i < size; i++) {
std::cout << ptr[i] << " ";
}
@@ -148,14 +147,14 @@ static void test_group_partition(unsigned int tileSz) {
int* dResult = NULL;
int* hResult = NULL;
hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault);
HIPCHECK(hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault));
memset(hResult, 0, numTiles * sizeof(int));
hipMalloc(&dResult, numTiles * sizeof(int));
HIPCHECK(hipMalloc(&dResult, numTiles * sizeof(int)));
// Launch Kernel
hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock,
hipLaunchKernelGGL(kernel_cg_group_partition_shfl_down, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0, dResult, tileSz, i);
HIP_CHECK(hipGetLastError());
err = hipDeviceSynchronize();
@@ -164,13 +163,13 @@ static void test_group_partition(unsigned int tileSz) {
}
hipMemcpy(hResult, dResult, sizeof(int) * numTiles, hipMemcpyDeviceToHost);
HIPCHECK(hipMemcpy(hResult, dResult, sizeof(int) * numTiles, hipMemcpyDeviceToHost));
verifyResults(hResult, expectedSum, numTiles);
// Free all allocated memory on host and device
hipFree(dResult);
hipFree(hResult);
HIPCHECK(hipFree(dResult));
HIPCHECK(hipHostFree(hResult));
delete[] expectedResult;
printf("\n...PASSED.\n\n");
@@ -199,7 +198,7 @@ static void test_shfl_down() {
int arrSize = blockSize * threadsPerBlock * sizeof(int);
hipHostMalloc(&hPtr, arrSize);
HIPCHECK(hipHostMalloc(&hPtr, arrSize));
// Fill up the array
for (int i = 0; i < WAVE_SIZE; i++) {
hPtr[i] = rand() % 1000;
@@ -210,30 +209,30 @@ static void test_shfl_down() {
cpuResultsArr[i] = (i + lane_delta >= group_size) ? hPtr[i] : hPtr[i + lane_delta];
}
//printf("Array passed to GPU for computation\n");
//printResults(hPtr, WAVE_SIZE);
hipMalloc(&dPtr, group_size_in_bytes);
hipMalloc(&dResults, group_size_in_bytes);
//printResultsCoalescedGroupsShflDown(hPtr, WAVE_SIZE);
HIPCHECK(hipMalloc(&dPtr, group_size_in_bytes));
HIPCHECK(hipMalloc(&dResults, group_size_in_bytes));
hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
HIPCHECK(hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice));
// Launch Kernel
hipLaunchKernelGGL(kernel_shfl_down, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i);
HIP_CHECK(hipGetLastError());
hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
HIPCHECK(hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost));
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
//printf("GPU results: \n");
//printResults(hPtr, WAVE_SIZE);
//printResultsCoalescedGroupsShflDown(hPtr, WAVE_SIZE);
//printf("Printing cpu to be verified array\n");
//printResults(cpuResultsArr, WAVE_SIZE);
//printResultsCoalescedGroupsShflDown(cpuResultsArr, WAVE_SIZE);
compareResults(hPtr, cpuResultsArr, group_size_in_bytes);
compareResultsCoalescedGroupsShflDown(hPtr, cpuResultsArr, group_size_in_bytes);
std::cout << "Results verified!\n";
hipFree(hPtr);
hipFree(dPtr);
HIPCHECK(hipHostFree(hPtr));
HIPCHECK(hipFree(dPtr));
free(cpuResultsArr);
}
}
@@ -246,7 +245,6 @@ TEST_CASE("Unit_coalesced_groups_shfl_down") {
ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess);
hipDeviceProp_t deviceProperties;
ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
// Test shfl_down with random group sizes
for (int i = 0; i < 100; i++) {
@@ -40,7 +40,7 @@ __device__ int prefix_sum_kernel(coalesced_group const& g, int val) {
for (int i = 1; i < sz; i <<= 1) {
int temp = g.shfl_up(val, i);
if (g.thread_rank() >= i) {
if ((int)g.thread_rank() >= i) {
val += temp;
}
}
@@ -60,7 +60,7 @@ __global__ void kernel_shfl_up (int * dPtr, int *dResults, int lane_delta, int c
}
__global__ void kernel_cg_group_partition(int* dPtr, unsigned int tileSz, int cg_sizes) {
__global__ void kernel_cg_group_partition_shfl_up(int* dPtr, unsigned int tileSz, int cg_sizes) {
int id = threadIdx.x + blockIdx.x * blockDim.x;
if (id % cg_sizes == 0) {
@@ -110,7 +110,7 @@ void printResults(int* ptr, int size) {
std::cout << '\n';
}
void verifyResults(int* cpu, int* gpu, int size) {
void verifyResultsCoalescedGroupsShflUp(int* cpu, int* gpu, int size) {
for (unsigned int i = 0; i < size / sizeof(int); i++) {
if (cpu[i] != gpu[i]) {
INFO(" Results do not match.");
@@ -132,14 +132,14 @@ static void test_group_partition(unsigned tileSz) {
int arrSize = blockSize * threadsPerBlock * sizeof(int);
hipHostMalloc(&hPtr, arrSize);
hipMalloc(&dPtr, arrSize);
HIPCHECK(hipHostMalloc(&hPtr, arrSize));
HIPCHECK(hipMalloc(&dPtr, arrSize));
// Launch Kernel
hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock,
hipLaunchKernelGGL(kernel_cg_group_partition_shfl_up, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0, dPtr, tileSz, i);
HIP_CHECK(hipGetLastError());
hipMemcpy(hPtr, dPtr, arrSize, hipMemcpyDeviceToHost);
HIPCHECK(hipMemcpy(hPtr, dPtr, arrSize, hipMemcpyDeviceToHost));
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
@@ -153,12 +153,12 @@ static void test_group_partition(unsigned tileSz) {
//std::cout << "\nPrefix sum results on GPU\n";
//printResults(hPtr, tileSz);
std::cout << "\n";
verifyResults(hPtr, cpuPrefixSum, tileSz);
verifyResultsCoalescedGroupsShflUp(hPtr, cpuPrefixSum, tileSz);
std::cout << "Results verified!\n";
delete[] cpuPrefixSum;
hipFree(hPtr);
hipFree(dPtr);
HIPCHECK(hipHostFree(hPtr));
HIPCHECK(hipFree(dPtr));
}
}
@@ -185,7 +185,7 @@ static void test_shfl_up() {
int arrSize = blockSize * threadsPerBlock * sizeof(int);
hipHostMalloc(&hPtr, arrSize);
HIPCHECK(hipHostMalloc(&hPtr, arrSize));
// Fill up the array
for (int i = 0; i < WAVE_SIZE; i++) {
hPtr[i] = rand() % 1000;
@@ -200,14 +200,14 @@ static void test_shfl_up() {
//printf("Printing cpu results arr\n");
//printResults(cpuResultsArr, WAVE_SIZE);
hipMalloc(&dPtr, group_size_in_bytes);
hipMalloc(&dResults, group_size_in_bytes);
HIPCHECK(hipMalloc(&dPtr, group_size_in_bytes));
HIPCHECK(hipMalloc(&dResults, group_size_in_bytes));
hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
HIPCHECK(hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice));
// Launch Kernel
hipLaunchKernelGGL(kernel_shfl_up, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i);
hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
HIPCHECK(hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipGetLastError());
err = hipDeviceSynchronize();
if (err != hipSuccess) {
@@ -216,22 +216,21 @@ static void test_shfl_up() {
//printf("GPU computation array :\n");
//printResults(hPtr, WAVE_SIZE);
verifyResults(hPtr, cpuResultsArr, group_size_in_bytes);
verifyResultsCoalescedGroupsShflUp(hPtr, cpuResultsArr, group_size_in_bytes);
std::cout << "Results verified!\n";
hipFree(hPtr);
hipFree(dPtr);
HIPCHECK(hipHostFree(hPtr));
HIPCHECK(hipFree(dPtr));
free(cpuResultsArr);
}
}
TEST_CASE("Unit_coalesced_groups_shfl_down") {
TEST_CASE("Unit_coalesced_groups_shfl_up") {
// Use default device for validating the test
int deviceId;
ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess);
hipDeviceProp_t deviceProperties;
ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
for (int i = 0; i < 100; i++) {
test_shfl_up();
@@ -51,7 +51,7 @@ using namespace cooperative_groups;
__device__ int atomicAggInc(int *ptr) {
coalesced_group g = coalesced_threads();
int prev;
int prev = 0;
// elect the first active thread to perform atomic add
if (g.thread_rank() == 0) {
prev = atomicAdd(ptr, g.size());
@@ -104,7 +104,6 @@ __global__ void filter_arr(int *dst, int *nres, const int *src, int n) {
*/
__device__ int reduction_kernel(coalesced_group g, int* x, int val) {
int lane = g.thread_rank();
int sz = g.size();
for (int i = g.size() / 2; i > 0; i /= 2) {
// use lds to store the temporary result
@@ -138,7 +137,6 @@ __global__ void kernel_cg_coalesced_group_partition(unsigned int tileSz, int* re
int id = threadIdx.x + blockIdx.x * blockDim.x;
if (id % cg_sizes == 0) {
coalesced_group threadBlockCGTy = coalesced_threads();
int threadBlockGroupSize = threadBlockCGTy.size();
int* workspace = NULL;
@@ -150,13 +148,11 @@ __global__ void kernel_cg_coalesced_group_partition(unsigned int tileSz, int* re
workspace = sharedMem;
}
int input, outputSum, expectedOutput;
int input, outputSum;
// input to reduction, for each thread, is its' rank in the group
input = threadBlockCGTy.thread_rank();
expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2;
outputSum = reduction_kernel(threadBlockCGTy, workspace, input);
if (threadBlockCGTy.thread_rank() == 0) {
@@ -189,10 +185,8 @@ __global__ void kernel_cg_coalesced_group_partition(unsigned int tileSz, int* re
__global__ void kernel_coalesced_active_groups() {
thread_block threadBlockCGTy = this_thread_block();
int threadBlockGroupSize = threadBlockCGTy.size();
// input to reduction, for each thread, is its' rank in the group
int input = threadBlockCGTy.thread_rank();
if (threadBlockCGTy.thread_rank() == 0) {
printf(" Creating odd and even set of active thread groups based on branch divergence\n\n");
@@ -222,14 +216,14 @@ __global__ void kernel_coalesced_active_groups() {
return;
}
void printResults(int* ptr, int size) {
void printResultsSimpleCoalescedGroups(int* ptr, int size) {
for (int i = 0; i < size; i++) {
std::cout << ptr[i] << " ";
}
std::cout << '\n';
}
void compareResults(int* cpu, int* gpu, int size) {
void compareResultsSimpleCoalescedGroups(int* cpu, int* gpu, int size) {
for (unsigned int i = 0; i < size / sizeof(int); i++) {
if (cpu[i] != gpu[i]) {
INFO(" results do not match.");
@@ -254,7 +248,7 @@ static void test_active_threads_grouping() {
}
// Search if the sum exists in the expected results array
void verifyResults(int* hPtr, int* dPtr, int size) {
void verifyResultsSimpleCoalescedGroups(int* hPtr, int* dPtr, int size) {
int i = 0, j = 0;
for (i = 0; i < size; i++) {
for (j = 0; j < size; j++) {
@@ -294,15 +288,15 @@ static void test_group_partition(unsigned int tileSz, bool useGlobalMem) {
}
int* dResult = NULL;
hipMalloc(&dResult, sizeof(int) * numTiles);
HIPCHECK(hipMalloc(&dResult, sizeof(int) * numTiles));
int* globalMem = NULL;
if (useGlobalMem) {
hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int));
HIPCHECK(hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int)));
}
int* hResult = NULL;
hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault);
HIPCHECK(hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault));
memset(hResult, 0, numTiles * sizeof(int));
// Launch Kernel
@@ -326,13 +320,13 @@ static void test_group_partition(unsigned int tileSz, bool useGlobalMem) {
}
}
hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost);
verifyResults(expectedSum, hResult, numTiles);
HIPCHECK(hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost));
verifyResultsSimpleCoalescedGroups(expectedSum, hResult, numTiles);
// Free all allocated memory on host and device
hipFree(dResult);
hipFree(hResult);
HIPCHECK(hipFree(dResult));
HIPCHECK(hipHostFree(hResult));
if (useGlobalMem) {
hipFree(globalMem);
HIPCHECK(hipFree(globalMem));
}
delete[] expectedSum;
@@ -363,7 +357,7 @@ static void test_shfl_any_to_any() {
int arrSize = blockSize * threadsPerBlock * sizeof(int);
hipHostMalloc(&hPtr, arrSize);
HIPCHECK(hipHostMalloc(&hPtr, arrSize));
// Fill up the array
for (int i = 0; i < WAVE_SIZE; i++) {
hPtr[i] = rand() % 1000;
@@ -382,36 +376,36 @@ static void test_shfl_any_to_any() {
}
//printf("Array passed to GPU for computation\n");
//printResults(hPtr, WAVE_SIZE);
hipMalloc(&dPtr, group_size_in_bytes);
hipMalloc(&dResults, group_size_in_bytes);
//printResultsSimpleCoalescedGroups(hPtr, WAVE_SIZE);
HIPCHECK(hipMalloc(&dPtr, group_size_in_bytes));
HIPCHECK(hipMalloc(&dResults, group_size_in_bytes));
hipMalloc(&dsrcArr, group_size_in_bytes);
hipMemcpy(dsrcArr, srcArr, group_size_in_bytes, hipMemcpyHostToDevice);
HIPCHECK(hipMalloc(&dsrcArr, group_size_in_bytes));
HIPCHECK(hipMemcpy(dsrcArr, srcArr, group_size_in_bytes, hipMemcpyHostToDevice));
hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
HIPCHECK(hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice));
// Launch Kernel
hipLaunchKernelGGL(kernel_shfl_any_to_any, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0 , dPtr, dsrcArr, dResults, i);
HIP_CHECK(hipGetLastError());
hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
HIPCHECK(hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost));
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
//printf("GPU results: \n");
//printResults(hPtr, group_size);
//printResultsSimpleCoalescedGroups(hPtr, group_size);
//printf("Printing cpu to be verified array\n");
//printResults(cpuResultsArr, group_size);
//printResultsSimpleCoalescedGroups(cpuResultsArr, group_size);
//printf("Printing srcLane array that was passed\n");
//printResults(srcArr, group_size);
//printResultsSimpleCoalescedGroups(srcArr, group_size);
//printf("Printing srcLane array on the CPU\n");
//printResults(srcArrCpu, group_size);
compareResults(hPtr, cpuResultsArr, group_size_in_bytes);
//printResultsSimpleCoalescedGroups(srcArrCpu, group_size);
compareResultsSimpleCoalescedGroups(hPtr, cpuResultsArr, group_size_in_bytes);
std::cout << "Results verified!\n";
hipFree(hPtr);
hipFree(dPtr);
HIPCHECK(hipHostFree(hPtr));
HIPCHECK(hipFree(dPtr));
free(srcArr);
free(srcArrCpu);
free(cpuResultsArr);
@@ -440,7 +434,7 @@ static void test_shfl_broadcast() {
int arrSize = blockSize * threadsPerBlock * sizeof(int);
hipHostMalloc(&hPtr, arrSize);
HIPCHECK(hipHostMalloc(&hPtr, arrSize));
// Fill up the array
for (int i = 0; i < WAVE_SIZE; i++) {
hPtr[i] = rand() % 1000;
@@ -455,30 +449,30 @@ static void test_shfl_broadcast() {
cpuResultsArr[i] = srcLaneCpu;
}
printf("Array passed to GPU for computation\n");
printResults(hPtr, WAVE_SIZE);
hipMalloc(&dPtr, group_size_in_bytes);
hipMalloc(&dResults, group_size_in_bytes);
printResultsSimpleCoalescedGroups(hPtr, WAVE_SIZE);
HIPCHECK(hipMalloc(&dPtr, group_size_in_bytes));
HIPCHECK(hipMalloc(&dResults, group_size_in_bytes));
hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
HIPCHECK(hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice));
// Launch Kernel
hipLaunchKernelGGL(kernel_shfl, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0, dPtr, dResults, srcLane, i);
HIP_CHECK(hipGetLastError());
hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
HIPCHECK(hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost));
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
printf("GPU results: \n");
printResults(hPtr, group_size);
printResultsSimpleCoalescedGroups(hPtr, group_size);
printf("Printing cpu to be verified array\n");
printResults(cpuResultsArr, group_size);
printResultsSimpleCoalescedGroups(cpuResultsArr, group_size);
compareResults(hPtr, cpuResultsArr, group_size_in_bytes);
compareResultsSimpleCoalescedGroups(hPtr, cpuResultsArr, group_size_in_bytes);
std::cout << "Results verified!\n";
hipFree(hPtr);
hipFree(dPtr);
HIPCHECK(hipHostFree(hPtr));
HIPCHECK(hipFree(dPtr));
free(cpuResultsArr);
}
}
@@ -489,7 +483,6 @@ TEST_CASE("Unit_coalesced_groups") {
HIP_CHECK(hipGetDevice(&deviceId));
hipDeviceProp_t deviceProperties;
HIP_CHECK(hipGetDeviceProperties(&deviceProperties, deviceId));
int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
std::cout << "Now testing coalesced_groups" << '\n' << std::endl;
@@ -512,7 +505,7 @@ TEST_CASE("Unit_coalesced_groups") {
HIP_CHECK(hipMemcpy(d_data_to_filter, data_to_filter,
sizeof(int) * NUM_ELEMS, hipMemcpyHostToDevice));
hipMemset(d_nres, 0, sizeof(int));
HIPCHECK(hipMemset(d_nres, 0, sizeof(int)));
dim3 dimBlock(NUM_THREADS_PER_BLOCK, 1, 1);
dim3 dimGrid((NUM_ELEMS / NUM_THREADS_PER_BLOCK) + 1, 1, 1);
@@ -38,16 +38,12 @@ TEST_CASE("Unit_hipDeviceSetCacheConfig_Positive_Basic") {
const auto cache_config =
GENERATE(from_range(std::begin(kCacheConfigs), std::end(kCacheConfigs)));
#if HT_AMD
HIP_CHECK_ERROR(hipDeviceSetCacheConfig(cache_config), hipErrorNotSupported);
#elif HT_NVIDIA
HIP_CHECK(hipDeviceSetCacheConfig(cache_config));
#endif
}
TEST_CASE("Unit_hipDeviceSetCacheConfig_Negative_Parameters") {
#if HT_AMD
HIP_CHECK_ERROR(hipDeviceSetCacheConfig(static_cast<hipFuncCache_t>(-1)), hipErrorNotSupported);
HIP_CHECK_ERROR(hipDeviceSetCacheConfig(static_cast<hipFuncCache_t>(-1)), hipSuccess);
#elif HT_NVIDIA
HIP_CHECK_ERROR(hipDeviceSetCacheConfig(static_cast<hipFuncCache_t>(-1)), hipErrorInvalidValue);
#endif
@@ -105,4 +101,4 @@ TEST_CASE("Unit_hipDeviceGetCacheConfig_Positive_Threaded") {
TEST_CASE("Unit_HipDeviceGetCacheConfig_Negative_Parameters") {
HIP_CHECK_ERROR(hipDeviceGetCacheConfig(nullptr), hipErrorInvalidValue);
}
}
@@ -35,11 +35,7 @@ TEST_CASE("Unit_hipDeviceSetSharedMemConfig_Positive_Basic") {
HIP_CHECK(hipSetDevice(device));
INFO("Current device is " << device);
#if HT_AMD
HIP_CHECK_ERROR(hipDeviceSetSharedMemConfig(mem_config), hipErrorNotSupported);
#elif HT_NVIDIA
HIP_CHECK(hipDeviceSetSharedMemConfig(mem_config));
#endif
}
TEST_CASE("Unit_hipDeviceSetSharedMemConfig_Negative_Parameters") {
@@ -113,4 +109,4 @@ TEST_CASE("Unit_hipDeviceGetSharedMemConfig_Positive_Threaded") {
TEST_CASE("Unit_hipDeviceGetSharedMemConfig_Negative_Parameters") {
HIP_CHECK_ERROR(hipDeviceGetSharedMemConfig(nullptr), hipErrorInvalidValue);
}
}
@@ -33,15 +33,15 @@ THE SOFTWARE.
#define NUM_ITERS 1 << 30
static __global__ void Iter(int* Ad, int num) {
int tx = threadIdx.x + blockIdx.x * blockDim.x;
// Kernel loop designed to execute very slowly.
// so we can test timing-related
// behavior below
if (tx == 0) {
for (int i = 0; i < num; i++) {
Ad[tx] += 1;
}
int tx = threadIdx.x + blockIdx.x * blockDim.x;
// Kernel loop designed to execute very slowly.
// so we can test timing-related
// behavior below
if (tx == 0) {
for (int i = 0; i < num; i++) {
Ad[tx] += 1;
}
}
}
TEST_CASE("Unit_hipDeviceSynchronize_Positive_Empty_Streams") {
@@ -61,42 +61,47 @@ TEST_CASE("Unit_hipDeviceSynchronize_Positive_Nullstream") {
INFO("Current device: " << device);
int *A_h = nullptr, *A_d = nullptr;
HipTest::BlockingContext b_context{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), _SIZE, hipHostMallocDefault));
A_h[0] = 1;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), _SIZE));
HIP_CHECK(hipMemcpyAsync(A_d, A_h, _SIZE, hipMemcpyHostToDevice, NULL));
b_context.block_stream();
REQUIRE(b_context.is_blocked());
hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, NULL, A_d, 1 << 30);
HIP_CHECK(hipMemcpyAsync(A_h, A_d, _SIZE, hipMemcpyDeviceToHost, NULL));
CHECK(1 << 30 != A_h[0] - 1);
REQUIRE(1 << 30 != A_h[0] - 1);
b_context.unblock_stream();
HIP_CHECK(hipDeviceSynchronize());
CHECK(1 << 30 == A_h[0] - 1);
REQUIRE(1 << 30 == A_h[0] - 1);
}
TEST_CASE("Unit_hipDeviceSynchronize_Functional") {
int* A[NUM_STREAMS];
int* Ad[NUM_STREAMS];
hipStream_t stream[NUM_STREAMS];
std::vector<HipTest::BlockingContext> b_context;
b_context.reserve(NUM_STREAMS);
for (int i = 0; i < NUM_STREAMS; i++) {
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A[i]), _SIZE,
hipHostMallocDefault));
A[i][0] = 1;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad[i]), _SIZE));
HIP_CHECK(hipStreamCreate(&stream[i]));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A[i]), _SIZE, hipHostMallocDefault));
A[i][0] = 1;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad[i]), _SIZE));
HIP_CHECK(hipStreamCreate(&stream[i]));
b_context.emplace_back(HipTest::BlockingContext(stream[i]));
}
for (int i = 0; i < NUM_STREAMS; i++) {
HIP_CHECK(hipMemcpyAsync(Ad[i], A[i], _SIZE, hipMemcpyHostToDevice,
stream[i]));
HIP_CHECK(hipMemcpyAsync(Ad[i], A[i], _SIZE, hipMemcpyHostToDevice, stream[i]));
}
for (int i = 0; i < NUM_STREAMS; i++) {
hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0,
stream[i], Ad[i], NUM_ITERS);
HIP_CHECK(hipGetLastError());
b_context[i].block_stream();
REQUIRE(b_context[i].is_blocked());
hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, stream[i], Ad[i], NUM_ITERS);
}
for (int i = 0; i < NUM_STREAMS; i++) {
HIP_CHECK(hipMemcpyAsync(A[i], Ad[i], _SIZE, hipMemcpyDeviceToHost,
stream[i]));
HIP_CHECK(hipMemcpyAsync(A[i], Ad[i], _SIZE, hipMemcpyDeviceToHost, stream[i]));
}
@@ -106,7 +111,10 @@ TEST_CASE("Unit_hipDeviceSynchronize_Functional") {
// Conservative implementations which synchronize the hipMemcpyAsync will
// fail, ie if HIP_LAUNCH_BLOCKING=true.
CHECK(NUM_ITERS != A[NUM_STREAMS - 1][0] - 1);
REQUIRE(NUM_ITERS != A[NUM_STREAMS - 1][0] - 1);
for (int i = 0; i < NUM_STREAMS; i++) {
b_context[i].unblock_stream();
}
HIP_CHECK(hipDeviceSynchronize());
CHECK(NUM_ITERS == A[NUM_STREAMS - 1][0] - 1);
REQUIRE(NUM_ITERS == A[NUM_STREAMS - 1][0] - 1);
}
@@ -267,3 +267,7 @@ TEST_CASE("Unit_vectorTypes_CompileTest") {
HIP_CHECK(hipFree(ptr));
REQUIRE(passed == true);
}
//Test to check __half2_raw can be initialized without designator
TEST_CASE("Unit_half2_raw_def") {
constexpr __half2_raw inf = {0x7C00,0x7C00};
}
@@ -4,6 +4,8 @@ set(TEST_SRC
hipGetErrorString.cc
hipGetLastError.cc
hipPeekAtLastError.cc
hipDrvGetErrorName.cc
hipDrvGetErrorString.cc
)
hip_add_exe_to_target(NAME ErrorHandlingTest
@@ -33,6 +33,7 @@ constexpr hipError_t kErrorEnumerators[] = {hipSuccess,
hipErrorProfilerNotInitialized,
hipErrorProfilerAlreadyStarted,
hipErrorProfilerAlreadyStopped,
#if HT_AMD
hipErrorInvalidConfiguration,
hipErrorInvalidPitchValue,
hipErrorInvalidSymbol,
@@ -42,6 +43,7 @@ constexpr hipError_t kErrorEnumerators[] = {hipSuccess,
hipErrorMissingConfiguration,
hipErrorPriorLaunchFailure,
hipErrorInvalidDeviceFunction,
#endif
hipErrorNoDevice,
hipErrorInvalidDevice,
hipErrorInvalidImage,
@@ -95,5 +97,8 @@ constexpr hipError_t kErrorEnumerators[] = {hipSuccess,
hipErrorStreamCaptureWrongThread,
hipErrorGraphExecUpdateFailure,
hipErrorUnknown,
#if HT_AMD
hipErrorRuntimeMemory,
hipErrorRuntimeOther};
hipErrorRuntimeOther
#endif
};
@@ -0,0 +1,363 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
#include <hip_test_common.hh>
#include "errorEnumerators.h"
// Local Function to return the error code in string
static const char *ErrorName(hipError_t enumerator) {
switch (enumerator) {
#if HT_AMD
case hipSuccess:
return "hipSuccess";
case hipErrorInvalidValue:
return "hipErrorInvalidValue";
case hipErrorOutOfMemory:
return "hipErrorOutOfMemory";
case hipErrorNotInitialized:
return "hipErrorNotInitialized";
case hipErrorDeinitialized:
return "hipErrorDeinitialized";
case hipErrorProfilerDisabled:
return "hipErrorProfilerDisabled";
case hipErrorProfilerNotInitialized:
return "hipErrorProfilerNotInitialized";
case hipErrorProfilerAlreadyStarted:
return "hipErrorProfilerAlreadyStarted";
case hipErrorProfilerAlreadyStopped:
return "hipErrorProfilerAlreadyStopped";
case hipErrorInvalidConfiguration:
return "hipErrorInvalidConfiguration";
case hipErrorInvalidSymbol:
return "hipErrorInvalidSymbol";
case hipErrorInvalidDevicePointer:
return "hipErrorInvalidDevicePointer";
case hipErrorInvalidMemcpyDirection:
return "hipErrorInvalidMemcpyDirection";
case hipErrorInsufficientDriver:
return "hipErrorInsufficientDriver";
case hipErrorMissingConfiguration:
return "hipErrorMissingConfiguration";
case hipErrorPriorLaunchFailure:
return "hipErrorPriorLaunchFailure";
case hipErrorInvalidDeviceFunction:
return "hipErrorInvalidDeviceFunction";
case hipErrorNoDevice:
return "hipErrorNoDevice";
case hipErrorInvalidDevice:
return "hipErrorInvalidDevice";
case hipErrorInvalidPitchValue:
return "hipErrorInvalidPitchValue";
case hipErrorInvalidImage:
return "hipErrorInvalidImage";
case hipErrorInvalidContext:
return "hipErrorInvalidContext";
case hipErrorContextAlreadyCurrent:
return "hipErrorContextAlreadyCurrent";
case hipErrorMapFailed:
return "hipErrorMapFailed";
case hipErrorUnmapFailed:
return "hipErrorUnmapFailed";
case hipErrorArrayIsMapped:
return "hipErrorArrayIsMapped";
case hipErrorAlreadyMapped:
return "hipErrorAlreadyMapped";
case hipErrorNoBinaryForGpu:
return "hipErrorNoBinaryForGpu";
case hipErrorAlreadyAcquired:
return "hipErrorAlreadyAcquired";
case hipErrorNotMapped:
return "hipErrorNotMapped";
case hipErrorNotMappedAsArray:
return "hipErrorNotMappedAsArray";
case hipErrorNotMappedAsPointer:
return "hipErrorNotMappedAsPointer";
case hipErrorECCNotCorrectable:
return "hipErrorECCNotCorrectable";
case hipErrorUnsupportedLimit:
return "hipErrorUnsupportedLimit";
case hipErrorContextAlreadyInUse:
return "hipErrorContextAlreadyInUse";
case hipErrorPeerAccessUnsupported:
return "hipErrorPeerAccessUnsupported";
case hipErrorInvalidKernelFile:
return "hipErrorInvalidKernelFile";
case hipErrorInvalidGraphicsContext:
return "hipErrorInvalidGraphicsContext";
case hipErrorInvalidSource:
return "hipErrorInvalidSource";
case hipErrorFileNotFound:
return "hipErrorFileNotFound";
case hipErrorSharedObjectSymbolNotFound:
return "hipErrorSharedObjectSymbolNotFound";
case hipErrorSharedObjectInitFailed:
return "hipErrorSharedObjectInitFailed";
case hipErrorOperatingSystem:
return "hipErrorOperatingSystem";
case hipErrorInvalidHandle:
return "hipErrorInvalidHandle";
case hipErrorIllegalState:
return "hipErrorIllegalState";
case hipErrorNotFound:
return "hipErrorNotFound";
case hipErrorNotReady:
return "hipErrorNotReady";
case hipErrorIllegalAddress:
return "hipErrorIllegalAddress";
case hipErrorLaunchOutOfResources:
return "hipErrorLaunchOutOfResources";
case hipErrorLaunchTimeOut:
return "hipErrorLaunchTimeOut";
case hipErrorPeerAccessAlreadyEnabled:
return "hipErrorPeerAccessAlreadyEnabled";
case hipErrorPeerAccessNotEnabled:
return "hipErrorPeerAccessNotEnabled";
case hipErrorSetOnActiveProcess:
return "hipErrorSetOnActiveProcess";
case hipErrorContextIsDestroyed:
return "hipErrorContextIsDestroyed";
case hipErrorAssert:
return "hipErrorAssert";
case hipErrorHostMemoryAlreadyRegistered:
return "hipErrorHostMemoryAlreadyRegistered";
case hipErrorHostMemoryNotRegistered:
return "hipErrorHostMemoryNotRegistered";
case hipErrorLaunchFailure:
return "hipErrorLaunchFailure";
case hipErrorNotSupported:
return "hipErrorNotSupported";
case hipErrorUnknown:
return "hipErrorUnknown";
case hipErrorRuntimeMemory:
return "hipErrorRuntimeMemory";
case hipErrorRuntimeOther:
return "hipErrorRuntimeOther";
case hipErrorCooperativeLaunchTooLarge:
return "hipErrorCooperativeLaunchTooLarge";
case hipErrorStreamCaptureUnsupported:
return "hipErrorStreamCaptureUnsupported";
case hipErrorStreamCaptureInvalidated:
return "hipErrorStreamCaptureInvalidated";
case hipErrorStreamCaptureMerge:
return "hipErrorStreamCaptureMerge";
case hipErrorStreamCaptureUnmatched:
return "hipErrorStreamCaptureUnmatched";
case hipErrorStreamCaptureUnjoined:
return "hipErrorStreamCaptureUnjoined";
case hipErrorStreamCaptureIsolation:
return "hipErrorStreamCaptureIsolation";
case hipErrorStreamCaptureImplicit:
return "hipErrorStreamCaptureImplicit";
case hipErrorCapturedEvent:
return "hipErrorCapturedEvent";
case hipErrorStreamCaptureWrongThread:
return "hipErrorStreamCaptureWrongThread";
case hipErrorGraphExecUpdateFailure:
return "hipErrorGraphExecUpdateFailure";
case hipErrorTbd:
return "hipErrorTbd";
default:
return "hipErrorUnknown";
#endif
#if HT_NVIDIA
case hipSuccess:
return "CUDA_SUCCESS";
case hipErrorInvalidValue:
return "CUDA_ERROR_INVALID_VALUE";
case hipErrorOutOfMemory:
return "CUDA_ERROR_OUT_OF_MEMORY";
case hipErrorNotInitialized:
return "CUDA_ERROR_NOT_INITIALIZED";
case hipErrorDeinitialized:
return "CUDA_ERROR_DEINITIALIZED";
case hipErrorProfilerDisabled:
return "CUDA_ERROR_PROFILER_DISABLED";
case hipErrorProfilerNotInitialized:
return "CUDA_ERROR_PROFILER_NOT_INITIALIZED";
case hipErrorProfilerAlreadyStarted:
return "CUDA_ERROR_PROFILER_ALREADY_STARTED";
case hipErrorProfilerAlreadyStopped:
return "CUDA_ERROR_PROFILER_ALREADY_STOPPED";
case hipErrorInvalidConfiguration:
return "CUDA_ERROR_UNKNOWN";
case hipErrorInvalidSymbol:
return "CUDA_ERROR_UNKNOWN";
case hipErrorInvalidDevicePointer:
return "CUDA_ERROR_UNKNOWN";
case hipErrorInvalidMemcpyDirection:
return "CUDA_ERROR_UNKNOWN";
case hipErrorInsufficientDriver:
return "CUDA_ERROR_UNKNOWN";
case hipErrorMissingConfiguration:
return "CUDA_ERROR_UNKNOWN";
case hipErrorPriorLaunchFailure:
return "CUDA_ERROR_UNKNOWN";
case hipErrorInvalidDeviceFunction:
return "CUDA_ERROR_UNKNOWN";
case hipErrorNoDevice:
return "CUDA_ERROR_NO_DEVICE";
case hipErrorInvalidDevice:
return "CUDA_ERROR_INVALID_DEVICE";
case hipErrorInvalidPitchValue:
return "CUDA_ERROR_UNKNOWN";
case hipErrorInvalidImage:
return "CUDA_ERROR_INVALID_IMAGE";
case hipErrorInvalidContext:
return "CUDA_ERROR_INVALID_CONTEXT";
case hipErrorContextAlreadyCurrent:
return "CUDA_ERROR_CONTEXT_ALREADY_CURRENT";
case hipErrorMapFailed:
return "CUDA_ERROR_MAP_FAILED";
case hipErrorUnmapFailed:
return "CUDA_ERROR_UNMAP_FAILED";
case hipErrorArrayIsMapped:
return "CUDA_ERROR_ARRAY_IS_MAPPED";
case hipErrorAlreadyMapped:
return "CUDA_ERROR_ALREADY_MAPPED";
case hipErrorNoBinaryForGpu:
return "CUDA_ERROR_NO_BINARY_FOR_GPU";
case hipErrorAlreadyAcquired:
return "CUDA_ERROR_ALREADY_ACQUIRED";
case hipErrorNotMapped:
return "CUDA_ERROR_NOT_MAPPED";
case hipErrorNotMappedAsArray:
return "CUDA_ERROR_NOT_MAPPED_AS_ARRAY";
case hipErrorNotMappedAsPointer:
return "CUDA_ERROR_NOT_MAPPED_AS_POINTER";
case hipErrorECCNotCorrectable:
return "CUDA_ERROR_ECC_UNCORRECTABLE";
case hipErrorUnsupportedLimit:
return "CUDA_ERROR_UNSUPPORTED_LIMIT";
case hipErrorContextAlreadyInUse:
return "CUDA_ERROR_CONTEXT_ALREADY_IN_USE";
case hipErrorPeerAccessUnsupported:
return "CUDA_ERROR_PEER_ACCESS_UNSUPPORTED";
case hipErrorInvalidKernelFile:
return "CUDA_ERROR_INVALID_PTX";
case hipErrorInvalidGraphicsContext:
return "CUDA_ERROR_INVALID_GRAPHICS_CONTEXT";
case hipErrorInvalidSource:
return "CUDA_ERROR_INVALID_SOURCE";
case hipErrorFileNotFound:
return "CUDA_ERROR_FILE_NOT_FOUND";
case hipErrorSharedObjectSymbolNotFound:
return "CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND";
case hipErrorSharedObjectInitFailed:
return "CUDA_ERROR_SHARED_OBJECT_INIT_FAILED";
case hipErrorOperatingSystem:
return "CUDA_ERROR_OPERATING_SYSTEM";
case hipErrorInvalidHandle:
return "CUDA_ERROR_INVALID_HANDLE";
case hipErrorIllegalState:
return "CUDA_ERROR_ILLEGAL_STATE";
case hipErrorNotFound:
return "CUDA_ERROR_NOT_FOUND";
case hipErrorNotReady:
return "CUDA_ERROR_NOT_READY";
case hipErrorIllegalAddress:
return "CUDA_ERROR_ILLEGAL_ADDRESS";
case hipErrorLaunchOutOfResources:
return "CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES";
case hipErrorLaunchTimeOut:
return "CUDA_ERROR_LAUNCH_TIMEOUT";
case hipErrorPeerAccessAlreadyEnabled:
return "CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED";
case hipErrorPeerAccessNotEnabled:
return "CUDA_ERROR_PEER_ACCESS_NOT_ENABLED";
case hipErrorSetOnActiveProcess:
return "CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE";
case hipErrorContextIsDestroyed:
return "CUDA_ERROR_CONTEXT_IS_DESTROYED";
case hipErrorAssert:
return "CUDA_ERROR_ASSERT";
case hipErrorHostMemoryAlreadyRegistered:
return "CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED";
case hipErrorHostMemoryNotRegistered:
return "CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED";
case hipErrorLaunchFailure:
return "CUDA_ERROR_LAUNCH_FAILED";
case hipErrorNotSupported:
return "CUDA_ERROR_NOT_SUPPORTED";
case hipErrorUnknown:
return "CUDA_ERROR_UNKNOWN";
case hipErrorRuntimeMemory:
return "CUDA_ERROR_UNKNOWN";
case hipErrorRuntimeOther:
return "CUDA_ERROR_UNKNOWN";
case hipErrorCooperativeLaunchTooLarge:
return "CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE";
case hipErrorStreamCaptureUnsupported:
return "CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED";
case hipErrorStreamCaptureInvalidated:
return "CUDA_ERROR_STREAM_CAPTURE_INVALIDATED";
case hipErrorStreamCaptureMerge:
return "CUDA_ERROR_STREAM_CAPTURE_MERGE";
case hipErrorStreamCaptureUnmatched:
return "CUDA_ERROR_STREAM_CAPTURE_UNMATCHED";
case hipErrorStreamCaptureUnjoined:
return "CUDA_ERROR_STREAM_CAPTURE_UNJOINED";
case hipErrorStreamCaptureIsolation:
return "CUDA_ERROR_STREAM_CAPTURE_ISOLATION";
case hipErrorStreamCaptureImplicit:
return "CUDA_ERROR_STREAM_CAPTURE_IMPLICIT";
case hipErrorCapturedEvent:
return "CUDA_ERROR_CAPTURED_EVENT";
case hipErrorStreamCaptureWrongThread:
return "CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD";
case hipErrorGraphExecUpdateFailure:
return "CUDA_ERROR_GRAPH_EXEC_UPDATE_FAILURE";
default:
return "CUDA_ERROR_UNKNOWN";
#endif
}
}
// Functional test case
// Test case to verify the returned error name is same as generated error name.
TEST_CASE("Unit_hipDrvGetErrorName_Functional") {
const char* error_string = nullptr;
hipError_t error_ret;
const auto enumerator =
GENERATE(from_range(std::begin(kErrorEnumerators),
std::end(kErrorEnumerators)));
error_ret = hipDrvGetErrorName(enumerator, &error_string);
REQUIRE(error_string != nullptr);
REQUIRE(strcmp(error_string, ErrorName(enumerator)) == 0);
REQUIRE(error_ret == hipSuccess);
}
// Negative test cases.
TEST_CASE("Unit_hipDrvGetErrorName_Negative") {
const char* error_string = nullptr;
SECTION("pass unknown value to hipError") {
REQUIRE((hipDrvGetErrorName(static_cast<hipError_t>(-1), &error_string))
== hipErrorInvalidValue);
}
#if HT_AMD
SECTION("pass nullptr to error string") {
REQUIRE((hipDrvGetErrorString(static_cast<hipError_t>(0), nullptr))
== hipErrorInvalidValue);
}
#endif
}
@@ -0,0 +1,263 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
#include <hip_test_common.hh>
#include "errorEnumerators.h"
// Local Function to return the error string.
static const char *ErrorString(hipError_t enumerator) {
switch (enumerator) {
case hipSuccess:
return "no error";
case hipErrorInvalidValue:
return "invalid argument";
case hipErrorOutOfMemory:
return "out of memory";
case hipErrorNotInitialized:
return "initialization error";
case hipErrorDeinitialized:
return "driver shutting down";
case hipErrorProfilerDisabled:
return "profiler disabled while using external profiling tool";
case hipErrorProfilerNotInitialized:
#if HT_AMD
return "profiler is not initialized";
#elif HT_NVIDIA
return "profiler not initialized: call cudaProfilerInitialize()";
#endif
case hipErrorProfilerAlreadyStarted:
return "profiler already started";
case hipErrorProfilerAlreadyStopped:
return "profiler already stopped";
#if HT_AMD
case hipErrorInvalidConfiguration:
return "invalid configuration argument";
#elif HT_NVIDIA
return "unknown error";
#endif
#if HT_AMD
case hipErrorInvalidPitchValue:
return "invalid pitch argument";
#elif HT_NVIDIA
return "unknown error";
#endif
#if HT_AMD
case hipErrorInvalidSymbol:
return "invalid device symbol";
#elif HT_NVIDIA
return "unknown error";
#endif
#if HT_AMD
case hipErrorInvalidDevicePointer:
return "invalid device pointer";
#elif HT_NVIDIA
return "unknown error";
#endif
#if HT_AMD
case hipErrorInvalidMemcpyDirection:
return "invalid copy direction for memcpy";
#elif HT_NVIDIA
return "unknown error";
#endif
#if HT_AMD
case hipErrorInsufficientDriver:
return "driver version is insufficient for runtime version";
#elif HT_NVIDIA
return "unknown error";
#endif
#if HT_AMD
case hipErrorMissingConfiguration:
return "__global__ function call is not configured";
#elif HT_NVIDIA
return "unknown error";
#endif
#if HT_AMD
case hipErrorPriorLaunchFailure:
return "unspecified launch failure in prior launch";
#elif HT_NVIDIA
return "unknown error";
#endif
#if HT_AMD
case hipErrorInvalidDeviceFunction:
return "invalid device function";
#elif HT_NVIDIA
return "unknown error";
#endif
case hipErrorNoDevice:
#if HT_AMD
return "no ROCm-capable device is detected";
#elif HT_NVIDIA
return "no CUDA-capable device is detected";
#endif
case hipErrorInvalidDevice:
return "invalid device ordinal";
case hipErrorInvalidImage:
return "device kernel image is invalid";
case hipErrorInvalidContext:
return "invalid device context";
case hipErrorContextAlreadyCurrent:
#if HT_AMD
return "context is already current context";
#elif HT_NVIDIA
return "context already current";
#endif
case hipErrorMapFailed:
return "mapping of buffer object failed";
case hipErrorUnmapFailed:
return "unmapping of buffer object failed";
case hipErrorArrayIsMapped:
return "array is mapped";
case hipErrorAlreadyMapped:
return "resource already mapped";
case hipErrorNoBinaryForGpu:
return "no kernel image is available for execution on the device";
case hipErrorAlreadyAcquired:
return "resource already acquired";
case hipErrorNotMapped:
return "resource not mapped";
case hipErrorNotMappedAsArray:
return "resource not mapped as array";
case hipErrorNotMappedAsPointer:
return "resource not mapped as pointer";
case hipErrorECCNotCorrectable:
return "uncorrectable ECC error encountered";
case hipErrorUnsupportedLimit:
return "limit is not supported on this architecture";
case hipErrorContextAlreadyInUse:
return "exclusive-thread device already in use by a different thread";
case hipErrorPeerAccessUnsupported:
return "peer access is not supported between these two devices";
case hipErrorInvalidKernelFile:
#if HT_AMD
return "invalid kernel file";
#elif HT_NVIDIA
return "a PTX JIT compilation failed";
#endif
case hipErrorInvalidGraphicsContext:
return "invalid OpenGL or DirectX context";
case hipErrorInvalidSource:
return "device kernel image is invalid";
case hipErrorFileNotFound:
return "file not found";
case hipErrorSharedObjectSymbolNotFound:
return "shared object symbol not found";
case hipErrorSharedObjectInitFailed:
return "shared object initialization failed";
case hipErrorOperatingSystem:
return "OS call failed or operation not supported on this OS";
case hipErrorInvalidHandle:
return "invalid resource handle";
case hipErrorIllegalState:
return "the operation cannot be performed in the present state";
case hipErrorNotFound:
return "named symbol not found";
case hipErrorNotReady:
return "device not ready";
case hipErrorIllegalAddress:
return "an illegal memory access was encountered";
case hipErrorLaunchOutOfResources:
return "too many resources requested for launch";
case hipErrorLaunchTimeOut:
return "the launch timed out and was terminated";
case hipErrorPeerAccessAlreadyEnabled:
return "peer access is already enabled";
case hipErrorPeerAccessNotEnabled:
return "peer access has not been enabled";
case hipErrorSetOnActiveProcess:
return "cannot set while device is active in this process";
case hipErrorContextIsDestroyed:
return "context is destroyed";
case hipErrorAssert:
return "device-side assert triggered";
case hipErrorHostMemoryAlreadyRegistered:
return "part or all of the requested memory range is already mapped";
case hipErrorHostMemoryNotRegistered:
return "pointer does not correspond to a registered memory region";
case hipErrorLaunchFailure:
return "unspecified launch failure";
case hipErrorCooperativeLaunchTooLarge:
return "too many blocks in cooperative launch";
case hipErrorNotSupported:
return "operation not supported";
case hipErrorStreamCaptureUnsupported:
return "operation not permitted when stream is capturing";
case hipErrorStreamCaptureInvalidated:
return "operation failed due to a previous error during capture";
case hipErrorStreamCaptureMerge:
return "operation would result in a merge of separate capture sequences";
case hipErrorStreamCaptureUnmatched:
return "capture was not ended in the same stream as it began";
case hipErrorStreamCaptureUnjoined:
return "capturing stream has unjoined work";
case hipErrorStreamCaptureIsolation:
return "dependency created on uncaptured work in another stream";
case hipErrorStreamCaptureImplicit:
return "operation would make the legacy stream depend on a capturing blocking stream"; //NOLINT
case hipErrorCapturedEvent:
return "operation not permitted on an event last recorded in a capturing stream"; //NOLINT
case hipErrorStreamCaptureWrongThread:
return "attempt to terminate a thread-local capture sequence from another thread"; //NOLINT
case hipErrorGraphExecUpdateFailure:
return "the graph update was not performed because it included changes which violated constraints specific to instantiated graph update"; //NOLINT
case hipErrorRuntimeMemory:
return "runtime memory call returned error";
case hipErrorRuntimeOther:
return "runtime call other than memory returned error";
case hipErrorUnknown:
default:
#if HT_AMD
return "unknown error";
#elif HT_NVIDIA
return "unknown error";
#endif
}
}
// Test case to verify the returned error string is
// same as generated error string.
TEST_CASE("Unit_hipDrvGetErrorString_Functional") {
const char* error_string = nullptr;
const auto enumerator =
GENERATE(from_range(std::begin(kErrorEnumerators),
std::end(kErrorEnumerators)));
hipError_t error_ret = hipDrvGetErrorString(enumerator, &error_string);
REQUIRE(error_string != nullptr);
REQUIRE(strcmp(error_string, ErrorString(enumerator)) == 0);
REQUIRE(error_ret == hipSuccess);
}
// Negative test cases.
TEST_CASE("Unit_hipDrvGetErrorString_Negative") {
const char* error_string = nullptr;
SECTION("pass unknown value to hipError") {
REQUIRE((hipDrvGetErrorString(static_cast<hipError_t>(-1), &error_string))
== hipErrorInvalidValue);
}
#if HT_AMD
SECTION("pass nullptr to error string") {
REQUIRE((hipDrvGetErrorString(static_cast<hipError_t>(0), nullptr))
== hipErrorInvalidValue);
}
#endif
}
@@ -96,6 +96,10 @@ set(TEST_SRC
hipLaunchHostFunc.cc
hipStreamGetCaptureInfo_v2_old.cc
hipUserObjectCreate.cc
hipUserObjectRelease.cc
hipUserObjectRetain.cc
hipGraphReleaseUserObject.cc
hipGraphRetainUserObject.cc
hipGraphDebugDotPrint.cc
hipGraphCloneComplx.cc
hipGraphNodeSetEnabled.cc
@@ -105,8 +109,21 @@ set(TEST_SRC
hipGraphExecDestroy.cc
hipGraphUpload.cc
hipGraphKernelNodeCopyAttributes.cc
hipGraphKernelNodeGetAttribute.cc
)
if(HIP_PLATFORM MATCHES "amd")
set(AMD_SRC
hipStreamCaptureExtModuleLaunchKernel.cc
)
set(TEST_SRC ${TEST_SRC} ${AMD_SRC})
endif()
hip_add_exe_to_target(NAME GraphsTest
TEST_SRC ${TEST_SRC}
TEST_TARGET_NAME build_tests)
if(HIP_PLATFORM MATCHES "amd")
add_custom_target(hipMatMul COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/hipMatMul.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/graph/hipMatMul.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH})
add_dependencies(build_tests hipMatMul)
endif()
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
@@ -21,10 +21,18 @@ THE SOFTWARE.
Testcase Scenarios :
1) Create and add empty node to graph and verify addition is successful.
2) Negative Scenarios
3) Functional Test to use empty node as barrier to wait for multiple nodes.
In a graph add 3 independent memcpy_h2d nodes, add an empty node with
dependencies on the 3 memcpy_h2d nodes, add 3 independent kernel nodes,
add another empty node with dependencies on the 3 kernel nodes and
add 3 independent memcpy_d2h nodes with dependencies on previous empty
node. Execute the graph and validate the results.
*/
#include <hip_test_checkers.hh>
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#define TEST_LOOP_SIZE 50
/**
* Functional Test to add empty node with dependencies
*/
@@ -90,7 +98,7 @@ TEST_CASE("Unit_hipGraphAddEmptyNode_NegTest") {
dependencies.data(), dependencies.size()));
}
// pDependencies is nullptr
SECTION("Null Graph") {
SECTION("Dependencies is null") {
REQUIRE(hipErrorInvalidValue == hipGraphAddEmptyNode(&emptyNode, graph,
nullptr, dependencies.size()));
}
@@ -98,3 +106,148 @@ TEST_CASE("Unit_hipGraphAddEmptyNode_NegTest") {
HIP_CHECK(hipFree(pOutBuff_d));
HIP_CHECK(hipGraphDestroy(graph));
}
// Function to fill input data
static void fillRandInpData(int *A1_h, int *A2_h, int *A3_h, size_t N) {
unsigned int seed = time(nullptr);
for (size_t i = 0; i < N; i++) {
A1_h[i] = (HipTest::RAND_R(&seed) & 0xFF);
A2_h[i] = (HipTest::RAND_R(&seed) & 0xFF);
A3_h[i] = (HipTest::RAND_R(&seed) & 0xFF);
}
}
// Function to validate result
static void validateOutData(int *A1_h, int *A2_h, size_t N) {
for (size_t i = 0; i < N; i++) {
int result = (A1_h[i]*A1_h[i]);
REQUIRE(result == A2_h[i]);
}
}
/**
* Functional Test to use empty node as barrier to wait for multiple nodes.
*/
TEST_CASE("Unit_hipGraphAddEmptyNode_BarrierFunc") {
size_t size = 1024;
constexpr auto blocksPerCU = 6;
constexpr auto threadsPerBlock = 256;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU,
threadsPerBlock, size);
hipGraph_t graph;
std::vector<hipGraphNode_t> nodeDependencies;
HIP_CHECK(hipGraphCreate(&graph, 0));
int *inputVec_d1{nullptr}, *inputVec_h1{nullptr}, *outputVec_h1{nullptr},
*outputVec_d1{nullptr};
int *inputVec_d2{nullptr}, *inputVec_h2{nullptr}, *outputVec_h2{nullptr},
*outputVec_d2{nullptr};
int *inputVec_d3{nullptr}, *inputVec_h3{nullptr}, *outputVec_h3{nullptr},
*outputVec_d3{nullptr};
// host and device allocation
HipTest::initArrays<int>(&inputVec_d1, &outputVec_d1, nullptr,
&inputVec_h1, &outputVec_h1, nullptr, size, false);
HipTest::initArrays<int>(&inputVec_d2, &outputVec_d2, nullptr,
&inputVec_h2, &outputVec_h2, nullptr, size, false);
HipTest::initArrays<int>(&inputVec_d3, &outputVec_d3, nullptr,
&inputVec_h3, &outputVec_h3, nullptr, size, false);
// add nodes to graph
hipGraphNode_t memcpyH2D_1, memcpyH2D_2, memcpyH2D_3;
hipGraphNode_t vecSqr1, vecSqr2, vecSqr3;
hipGraphNode_t memcpyD2H_1, memcpyD2H_2, memcpyD2H_3;
hipGraphNode_t emptyNode1, emptyNode2;
// Create memcpy h2d nodes
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_1, graph, nullptr,
0, inputVec_d1, inputVec_h1, (sizeof(int)*size), hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_2, graph, nullptr,
0, inputVec_d2, inputVec_h2, (sizeof(int)*size), hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_3, graph, nullptr,
0, inputVec_d3, inputVec_h3, (sizeof(int)*size), hipMemcpyHostToDevice));
// Create dependency list
nodeDependencies.push_back(memcpyH2D_1);
nodeDependencies.push_back(memcpyH2D_2);
nodeDependencies.push_back(memcpyH2D_3);
// Create emptyNode and add it to graph with dependency
HIP_CHECK(hipGraphAddEmptyNode(&emptyNode1, graph, nodeDependencies.data(),
nodeDependencies.size()));
nodeDependencies.clear();
nodeDependencies.push_back(emptyNode1);
// Creating kernel nodes
hipKernelNodeParams kerNodeParams1{}, kerNodeParams2{}, kerNodeParams3{};
void* kernelArgs1[] = {reinterpret_cast<void*>(&inputVec_d1),
reinterpret_cast<void*>(&outputVec_d1),
reinterpret_cast<void*>(&size)};
kerNodeParams1.func = reinterpret_cast<void*>(HipTest::vector_square<int>);
kerNodeParams1.gridDim = dim3(blocks);
kerNodeParams1.blockDim = dim3(threadsPerBlock);
kerNodeParams1.sharedMemBytes = 0;
kerNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs1);
kerNodeParams1.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&vecSqr1, graph, nodeDependencies.data(),
nodeDependencies.size(), &kerNodeParams1));
void* kernelArgs2[] = {reinterpret_cast<void*>(&inputVec_d2),
reinterpret_cast<void*>(&outputVec_d2),
reinterpret_cast<void*>(&size)};
kerNodeParams2.func = reinterpret_cast<void*>(HipTest::vector_square<int>);
kerNodeParams2.gridDim = dim3(blocks);
kerNodeParams2.blockDim = dim3(threadsPerBlock);
kerNodeParams2.sharedMemBytes = 0;
kerNodeParams2.kernelParams = reinterpret_cast<void**>(kernelArgs2);
kerNodeParams2.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&vecSqr2, graph, nodeDependencies.data(),
nodeDependencies.size(), &kerNodeParams2));
void* kernelArgs3[] = {reinterpret_cast<void*>(&inputVec_d3),
reinterpret_cast<void*>(&outputVec_d3),
reinterpret_cast<void*>(&size)};
kerNodeParams3.func = reinterpret_cast<void*>(HipTest::vector_square<int>);
kerNodeParams3.gridDim = dim3(blocks);
kerNodeParams3.blockDim = dim3(threadsPerBlock);
kerNodeParams3.sharedMemBytes = 0;
kerNodeParams3.kernelParams = reinterpret_cast<void**>(kernelArgs3);
kerNodeParams3.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&vecSqr3, graph, nodeDependencies.data(),
nodeDependencies.size(), &kerNodeParams3));
nodeDependencies.clear();
nodeDependencies.push_back(vecSqr1);
nodeDependencies.push_back(vecSqr2);
nodeDependencies.push_back(vecSqr3);
// Create emptyNode and add it to graph with dependency
HIP_CHECK(hipGraphAddEmptyNode(&emptyNode2, graph, nodeDependencies.data(),
nodeDependencies.size()));
nodeDependencies.clear();
nodeDependencies.push_back(emptyNode2);
// Create memcpy d2h nodes
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_1, graph,
nodeDependencies.data(), nodeDependencies.size(), outputVec_h1,
outputVec_d1, (sizeof(int)*size), hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_2, graph,
nodeDependencies.data(), nodeDependencies.size(), outputVec_h2,
outputVec_d2, (sizeof(int)*size), hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_3, graph,
nodeDependencies.data(), nodeDependencies.size(), outputVec_h3,
outputVec_d3, (sizeof(int)*size), hipMemcpyDeviceToHost));
nodeDependencies.clear();
// Create executable graph
hipStream_t streamForGraph;
hipGraphExec_t graphExec{nullptr};
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
nullptr, 0));
// Execute graph
for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) {
fillRandInpData(inputVec_h1, inputVec_h2, inputVec_h3, size);
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
validateOutData(inputVec_h1, outputVec_h1, size);
validateOutData(inputVec_h2, outputVec_h2, size);
validateOutData(inputVec_h3, outputVec_h3, size);
}
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipStreamDestroy(streamForGraph));
// Free
HipTest::freeArrays<int>(inputVec_d1, outputVec_d1, nullptr,
inputVec_h1, outputVec_h1, nullptr, false);
HipTest::freeArrays<int>(inputVec_d2, outputVec_d2, nullptr,
inputVec_h2, outputVec_h2, nullptr, false);
HipTest::freeArrays<int>(inputVec_d3, outputVec_d3, nullptr,
inputVec_h3, outputVec_h3, nullptr, false);
HIP_CHECK(hipGraphDestroy(graph));
}
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2023 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
@@ -48,10 +48,15 @@ Testcase Scenarios : Functional
Memcpy nodes are added and assigned to Peer device.
4) Perform memcpy operation for 1D, 2D and 3D arrays on Peer device and
verify the results.
5) Create two host pointers, copy the data between them by the api
hipGraphAddMemcpyNode with data transfer kind hipMemcpyHostToHost.
Validate the output.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <vector>
#include <numeric>
#define ZSIZE 32
#define YSIZE 32
@@ -517,4 +522,49 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_PeerAccessFunctional") {
validateMemcpyNode1DArray(true);
}
}
/*
* Create two host pointers, copy the data between them by the api
* hipGraphAddMemcpyNode with data transfer kind hipMemcpyHostToHost.
* Validate the output.
*/
TEST_CASE("Unit_hipGraphAddMemcpyNode_HostToHost") {
constexpr size_t size = 1024;
size_t numW = size * sizeof(int);
// Host Vectors
std::vector<int> A_h(numW);
std::vector<int> B_h(numW);
// Initialization
std::iota(A_h.begin(), A_h.end(), 0);
std::fill_n(B_h.begin(), size, 0);
hipGraph_t graph;
hipStream_t streamForGraph;
hipGraphExec_t graphExec;
hipGraphNode_t memcpyH2H;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&streamForGraph));
hipMemcpy3DParms myparms{};
myparms.srcPos = make_hipPos(0, 0, 0);
myparms.dstPos = make_hipPos(0, 0, 0);
myparms.srcPtr = make_hipPitchedPtr(A_h.data(), numW, numW, 1);
myparms.dstPtr = make_hipPitchedPtr(B_h.data(), numW, numW, 1);
myparms.extent = make_hipExtent(numW, 1, 1);
myparms.kind = hipMemcpyHostToHost;
// Host to Host
HIP_CHECK(hipGraphAddMemcpyNode(&memcpyH2H, graph, nullptr,
0, &myparms));
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
// Validation
REQUIRE(memcmp(A_h.data(), B_h.data(), numW) == 0);
}
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2023 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
@@ -24,6 +24,9 @@ Functional -
Memcpy nodes are added and assigned to default device.
2) Allocate memory on default device(Dev 0), Perform memcpy operation for 1D arrays on Peer device(Dev 1) and
verify the results.
3) Create two host pointers, copy the data between them by the api hipGraphAddMemcpyNode1D with data transfer
kind hipMemcpyHostToHost. Validate the output.
Negative -
1) Pass pGraphNode as nullptr and check if api returns error.
2) When graph is un-initialized argument(skipping graph creation), api should return error code.
@@ -38,7 +41,8 @@ Negative -
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <vector>
#include <numeric>
static void validateMemcpyNode1DArray(bool peerAccess) {
constexpr int SIZE{32};
@@ -197,46 +201,42 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Negative") {
HIP_CHECK(hipFree(A_h));
HIP_CHECK(hipGraphDestroy(graph));
}
/*
* Create two host pointers, copy the data between them by the api
* hipGraphAddMemcpyNode1D with data transfer kind hipMemcpyHostToHost.
* Validate the output.
*/
TEST_CASE("Unit_hipGraphAddMemcpyNode1D_HostToHost") {
constexpr size_t size = 1024;
size_t numBytes{size * sizeof(int)};
TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Negative_Basic") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h;
// Host Vectors
std::vector<int> A_h(size);
std::vector<int> B_h(size);
// Initialization
std::iota(A_h.begin(), A_h.end(), 0);
std::fill_n(B_h.begin(), size, 0);
hipGraphNode_t memcpy_A, memcpy_B, memcpy_C;
hipError_t ret;
hipGraph_t graph;
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
hipStream_t streamForGraph;
hipGraphExec_t graphExec;
hipGraphNode_t memcpyH2H;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipStreamCreate(&streamForGraph));
// Pass memcpy direction as hipMemcpyDeviceToHost and
// source pointer as host and destination pointer as device pointer
ret = hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyDeviceToHost);
REQUIRE(hipErrorInvalidValue == ret);
// Host to Host
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2H, graph, nullptr, 0,
B_h.data(), A_h.data(), numBytes, hipMemcpyHostToHost));
// Pass memcpy direction as hipMemcpyHostToDevice and
// source pointer as device and destination pointer as host pointer
ret = hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Pass memcpy direction as hipMemcpyDeviceToDevice and
// pass source pointer as device and destination pointer as host pointer
ret = hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
// Pass memcpy direction as hipMemcpyDeviceToDevice and
// pass source pointer as host and destination pointer as device pointer
ret = hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_d, C_h,
Nbytes, hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
// Validation
REQUIRE(std::equal(A_h.begin(), A_h.end(), B_h.begin(), B_h.end()));
}
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2022-2023 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
@@ -17,6 +17,14 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
* @addtogroup hipGraphDebugDotPrint hipGraphDebugDotPrint
* @{
* @ingroup GraphTest
* `hipGraphDebugDotPrint(hipGraph_t graph, const char* path, unsigned int flags)` -
* Write a DOT file describing graph structure.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
@@ -30,6 +38,13 @@ THE SOFTWARE.
#include <algorithm>
#include <string>
__device__ int globalIn[N];
static void callbackfunc(void *A_h) {
int *A = reinterpret_cast<int *>(A_h);
std::iota(A, A+N, 0);
}
static void deleteFile(const char* fName) {
if ( remove(fName) != 0 ) {
INFO("Error in deleting file -" << fName);
@@ -42,8 +57,9 @@ static bool checkFileExists(const char* fName) {
return (access(fName, F_OK) != -1);
}
static int countSubstr(std::string input_str, std::string substr) {
int count = 0;
static unsigned countSubstr(const std::string &input_str,
const std::string &substr) {
unsigned count = 0;
std::string::size_type srch_pos = 0, cur_pos = 0;
while ((cur_pos = input_str.find(substr, srch_pos)) != std::string::npos) {
count++;
@@ -53,35 +69,65 @@ static int countSubstr(std::string input_str, std::string substr) {
}
static bool validateDotFile(const char* fName,
std::map<std::string, int> *graphData) {
bool isTestPassed = true;
const std::map<std::string, unsigned> &graphData) {
std::ifstream infile(fName);
std::stringstream buffer;
buffer << infile.rdbuf();
std::map<std::string, int>::iterator it;
for (it = (*graphData).begin(); it != (*graphData).end(); it++) {
if (it->second != countSubstr(buffer.str(), it->first)) {
isTestPassed = false;
break;
const std::string buffer_str = buffer.str();
for (auto it = graphData.begin(); it != graphData.end(); it++) {
unsigned count = countSubstr(buffer_str, it->first);
if (it->second != count) {
INFO("validateDotFile: Failed for key :: " << it->first <<
" : " << count << " Expected : " << it->second);
return false;
}
}
return isTestPassed;
return true;
}
/**
* Test Description
* ------------------------
*  - Functional Test for API - hipGraphDebugDotPrint
* Call hipGraphDebugDotPrint and provice path where to write the DOT file.
* Verify that DOT file get created or not for each flag passed.
* 1) Add MemcpyNode node to graph & validate its DebugDotPrint descriptions
* 2) Add kernel node to graph & validate its DebugDotPrint descriptions
* 3) Add memset node to graph & validate its DebugDotPrint descriptions
* 4) Add emptyNode to graph & validate its DebugDotPrint descriptions
* 5) Add childGraphNode to graph & validate its DebugDotPrint descriptions
* 6) Add eventRecord to graph & validate its DebugDotPrint descriptions
* 7) Add eventWait to graph & validate its DebugDotPrint descriptions
* 8) Add hostNode to graph & validate its DebugDotPrint descriptions
* 9) Add mecpyNode1D to graph & validate its DebugDotPrint descriptions
* 10) Add mecpyNode3D to graph & validate its DebugDotPrint descriptions
* 11) Add MemcpyNodeToSymbol to graph & validate its DebugDotPrint descriptions
* 12) Add MemcpyNodeFromSymbol to graph & validate its DebugDotPrint descriptions
* 13) Add Dependencies to graph & validate its DebugDotPrint descriptions
* Test source
* ------------------------
*  - /unit/graph/hipGraphDebugDotPrint.cc
* Test requirements
* ------------------------
*  - HIP_VERSION >= 5.6
*/
static void hipGraphDebugDotPrint_Functional(const char* fName,
unsigned int flag = 0) {
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipGraph_t graph, childGraph;
hipStream_t stream;
hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, kNodeAdd, memsetNode;
hipGraphNode_t emptyNode, childGraphNode, eventWait, eventRecord, hostNode;
hipKernelNodeParams kNodeParams{};
int *A_d, *B_d, *C_d, *mem_d;
int *A_h, *B_h, *C_h;
int *A_h, *B_h, *C_h, *mem_h;
hipGraphExec_t graphExec;
size_t NElem{N};
mem_h = reinterpret_cast<int*>(malloc(Nbytes));
HIP_CHECK(hipMalloc(&mem_d, Nbytes));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
@@ -100,15 +146,88 @@ static void hipGraphDebugDotPrint_Functional(const char* fName,
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
// Add Kernel node to graph & validate its DebugDotPrint descriptions
HIP_CHECK(hipGraphAddKernelNode(&kNodeAdd, graph, nullptr, 0, &kNodeParams));
// Add MemCpy node to graph & validate its DebugDotPrint descriptions
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
// Add Dependencies to graph & validate its DebugDotPrint descriptions
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kNodeAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kNodeAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kNodeAdd, &memcpy_C, 1));
// Add emptyNode to graph & validate its DebugDotPrint descriptions
HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, NULL, 0));
// Add hostNode to graph & validate its DebugDotPrint descriptions
hipHostNodeParams hostParams = {0, 0};
hostParams.fn = callbackfunc;
hostParams.userData = mem_h;
HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams));
hipEvent_t event;
HIP_CHECK(hipEventCreate(&event));
// Add eventRecord to graph & validate its DebugDotPrint descriptions
HIP_CHECK(hipGraphAddEventRecordNode(&eventRecord, graph, nullptr,
0, event));
// Add eventWait to graph & validate its DebugDotPrint descriptions
HIP_CHECK(hipGraphAddEventWaitNode(&eventWait, graph, nullptr, 0, event));
HIP_CHECK(hipGraphCreate(&childGraph, 0));
// Add emcpyNode3D to graph & validate its DebugDotPrint descriptions
constexpr int width{10}, height{10}, depth{10};
hipArray *devArray1;
hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
hipMemcpy3DParms myparams;
uint32_t size = width * height * depth * sizeof(int);
hipGraphNode_t mcpyNode3D;
int *hData = reinterpret_cast<int*>(malloc(size));
REQUIRE(hData != nullptr);
// Initialize host buffer
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
hData[i*width*height + j*width + k] = i*width*height + j*width + k;
}
}
}
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
0, 0, 0, formatKind);
HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc,
make_hipExtent(width, height, depth), hipArrayDefault));
memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
myparams.srcPos = make_hipPos(0, 0, 0);
myparams.dstPos = make_hipPos(0, 0, 0);
myparams.extent = make_hipExtent(width , height, depth);
myparams.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int),
width, height);
myparams.dstArray = devArray1;
myparams.kind = hipMemcpyHostToDevice;
HIP_CHECK(hipGraphAddMemcpyNode(&mcpyNode3D, childGraph,
nullptr, 0, &myparams));
// Add MemcpyNodeToSymbol to graph & validate its DebugDotPrint description
hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode;
HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, childGraph,
nullptr, 0, HIP_SYMBOL(globalIn),
B_h, Nbytes, 0, hipMemcpyHostToDevice));
// Add MemcpyNodeFromSymbol to graph & validate its DebugDotPrint description
HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, childGraph,
nullptr, 0, B_h, HIP_SYMBOL(globalIn),
Nbytes, 0, hipMemcpyDeviceToHost));
// Add memset node to graph & validate its DebugDotPrint descriptions
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(mem_d);
@@ -117,20 +236,42 @@ static void hipGraphDebugDotPrint_Functional(const char* fName,
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, childGraph, nullptr, 0,
&memsetParams));
std::map<std::string, int> graphData;
graphData["->"] = 3; // number of edges
graphData["MEMCPY"] = 3;
graphData["HtoD"] = 2;
graphData["DtoH"] = 1;
// Add childGraphNode to graph & validate its DebugDotPrint descriptions
HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph,
nullptr, 0, childGraph));
std::map<std::string, unsigned> graphData;
graphData["->"] = 3; // number of edges
graphData["MEMCPY"] = 6;
graphData["HtoA"] = 1;
graphData["HtoD"] = 3;
graphData["DtoH"] = 2;
graphData["MEMSET"] = 1;
if ( flag == hipGraphDebugDotFlagsVerbose ) graphData["KERNEL"] = 1;
graphData["EMPTY"] = 1;
graphData["EVENT_WAIT"] = 1;
graphData["EVENT_RECORD"] = 1;
graphData["subgraph"] = 2;
graphData["HOST"] = 1;
#if HT_NVIDIA
if ( flag == hipGraphDebugDotFlagsVerbose ||
flag == hipGraphDebugDotFlagsMemcpyNodeParams ) {
graphData["HOST"] = 7;
}
#endif
if ( flag == hipGraphDebugDotFlagsVerbose ||
flag == hipGraphDebugDotFlagsKernelNodeAttributes
) {
graphData["KERNEL"] = 1;
}
HIP_CHECK(hipGraphDebugDotPrint(graph, fName, flag));
REQUIRE(true == checkFileExists(fName));
REQUIRE(true == validateDotFile(fName, &graphData));
REQUIRE(true == validateDotFile(fName, graphData));
deleteFile(fName);
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
@@ -140,10 +281,12 @@ static void hipGraphDebugDotPrint_Functional(const char* fName,
// Verify graph execution result
HipTest::checkVectorADD<int>(A_h, B_h, C_h, N);
free(mem_h);
HIP_CHECK(hipFree(mem_d));
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipGraphDestroy(childGraph));
HIP_CHECK(hipStreamDestroy(stream));
}
@@ -189,17 +332,24 @@ TEST_CASE("Unit_hipGraphDebugDotPrint_Functional") {
hipGraphDebugDotFlagsKernelNodeAttributes);
}
}
#endif // __linux__
/**
* Negative Test for API - hipGraphDebugDotPrint Argument Check
1) Pass graph as nullptr.
2) Pass graph as uninitialize structure
3) Pass path for dot file to store as nullptr
4) Pass path for dot file to store as empth path
5) Pass flag as hipGraphDebugDotFlags MIN - 1
6) Pass flag as hipGraphDebugDotFlags MAX + 1
7) Pass flag as INT_MAX
* Test Description
* ------------------------
*  - Negative Test for API - hipGraphDebugDotPrint Argument Check
* 1) Pass graph as nullptr
* 2) Pass graph as uninitialize structure
* 3) Pass path for dot file to store as nullptr
* 4) Pass path for dot file to store as empth path
* 5) Pass flag as hipGraphDebugDotFlags MIN - 1
* 6) Pass flag as hipGraphDebugDotFlags MAX + 1
* 7) Pass flag as INT_MAX
* Test source
* ------------------------
*  - /unit/graph/hipGraphDebugDotPrint.cc
* Test requirements
* ------------------------
*  - HIP_VERSION >= 5.6
*/
#define DOT_FILE_PATH_NEG "./graphDotFileNeg.dot"
@@ -241,6 +391,7 @@ TEST_CASE("Unit_hipGraphDebugDotPrint_Argument_Check") {
ret = hipGraphDebugDotPrint(graph, DOT_FILE_PATH_NEG, INT_MAX);
REQUIRE(hipSuccess == ret);
}
deleteFile(DOT_FILE_PATH_NEG);
HIP_CHECK(hipGraphDestroy(graph));
}
#endif // __linux__
@@ -184,17 +184,15 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative_Parameters")
}
#endif
SECTION("Changing dst allocation device") {
if (HipTest::getDeviceCount() < 2) {
HipTest::HIP_SKIP_TEST("Test requires two connected GPUs");
return;
}
if (HipTest::getDeviceCount() >= 2) {
SECTION("Changing dst allocation device") {
HIP_CHECK(hipSetDevice(1));
LinearAllocGuard<int> new_var(LinearAllocs::hipMalloc, sizeof(int));
HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsFromSymbol(
graph_exec, node, new_var.ptr(), SYMBOL(int_device_var),
sizeof(*new_var.ptr()), 0, static_cast<hipMemcpyKind>(-1)),
sizeof(*new_var.ptr()), 0, hipMemcpyDeviceToDevice),
hipErrorInvalidValue);
}
}
HIP_CHECK(hipGraphExecDestroy(graph_exec));
@@ -186,7 +186,7 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Negative_Parameters") {
}
HIP_CHECK(hipSetDevice(1));
LinearAllocGuard<int> new_var(LinearAllocs::hipMalloc, sizeof(int));
HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsFromSymbol(
HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsToSymbol(
graph_exec, node, SYMBOL(int_device_var), new_var.ptr(),
sizeof(*new_var.ptr()), 0, static_cast<hipMemcpyKind>(-1)),
hipErrorInvalidValue);
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2022-2023 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
@@ -22,10 +22,42 @@ THE SOFTWARE.
#include <hip_test_kernels.hh>
/**
* Functional Test for API - hipGraphKernelNodeCopyAttributes
- Create graph with 2 kernel node.
Copy the attribute from 1st kernel node to 2nd kernel node.
*/
* @addtogroup hipGraphKernelNodeCopyAttributes hipGraphKernelNodeCopyAttributes
* @{
* @ingroup GraphTest
* `hipGraphKernelNodeCopyAttributes(hipGraphNode_t hSrc, hipGraphNode_t hDst)` -
* Copies attributes from source node to destination node.
*/
/**
* Test Description
* ------------------------
*  - Functional Test for API - hipGraphKernelNodeCopyAttributes
* Create graph with kernel node and add its- functionality.
* 1) Copy kernelNodeAttribute to same graph kernel node.
* 2) Copy kernelNodeAttribute to different graph kernel node.
* 3) Copy kernelNodeAttribute to cloned graph kernel node.
* 4) Copy kernelNodeAttribute to child graph kernel node.
* Test source
* ------------------------
*  - /unit/graph/hipGraphKernelNodeCopyAttributes.cc
* Test requirements
* ------------------------
*  - HIP_VERSION >= 5.6
*/
static bool validateKernelNodeAttrValue(hipKernelNodeAttrValue in,
hipKernelNodeAttrValue out) {
if ((in.accessPolicyWindow.base_ptr != out.accessPolicyWindow.base_ptr) ||
(in.accessPolicyWindow.hitProp != out.accessPolicyWindow.hitProp) ||
(in.accessPolicyWindow.hitRatio != out.accessPolicyWindow.hitRatio) ||
(in.accessPolicyWindow.missProp != out.accessPolicyWindow.missProp) ||
(in.accessPolicyWindow.num_bytes != out.accessPolicyWindow.num_bytes) ||
(in.cooperative != out.cooperative)) {
return false;
}
return true;
}
TEST_CASE("Unit_hipGraphKernelNodeCopyAttributes_Functional") {
constexpr size_t N = 1024;
@@ -33,26 +65,25 @@ TEST_CASE("Unit_hipGraphKernelNodeCopyAttributes_Functional") {
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipGraphNode_t memcpyNode, kNode;
hipGraphExec_t graphExec;
hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, kernel_vecAdd;
hipKernelNodeParams kNodeParams{};
hipStream_t stream;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
std::vector<hipGraphNode_t> dependencies;
hipGraphExec_t graphExec;
size_t NElem{N};
HIP_CHECK(hipStreamCreate(&stream));
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(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
@@ -61,18 +92,101 @@ TEST_CASE("Unit_hipGraphKernelNodeCopyAttributes_Functional") {
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(),
dependencies.size(), &kNodeParams));
dependencies.clear();
dependencies.push_back(kNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(),
dependencies.size(), C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
&kNodeParams));
hipGraphNode_t kNode2;
// Create dependencies
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpy_C, 1));
HIP_CHECK(hipGraphAddKernelNode(&kNode2, graph, nullptr, 0, &kNodeParams));
HIP_CHECK(hipGraphKernelNodeCopyAttributes(kNode2, kNode));
hipKernelNodeAttrValue value_in, value_out;
memset(&value_in, 0, sizeof(hipKernelNodeAttrValue));
memset(&value_out, 0, sizeof(hipKernelNodeAttrValue));
SECTION("Copy kernelNodeAttribute to same graph kernel node") {
hipGraphNode_t kNode2;
HIP_CHECK(hipGraphAddKernelNode(&kNode2, graph, nullptr, 0, &kNodeParams));
memset(&value_in, 0, sizeof(hipKernelNodeAttrValue));
memset(&value_out, 0, sizeof(hipKernelNodeAttrValue));
HIP_CHECK(hipGraphKernelNodeGetAttribute(kernel_vecAdd,
hipKernelNodeAttributeAccessPolicyWindow, &value_in));
HIP_CHECK(hipGraphKernelNodeCopyAttributes(kernel_vecAdd, kNode2));
HIP_CHECK(hipGraphKernelNodeGetAttribute(kNode2,
hipKernelNodeAttributeAccessPolicyWindow, &value_out));
REQUIRE(true == validateKernelNodeAttrValue(value_in, value_out));
// copy back the node attributes for functional verification only
HIP_CHECK(hipGraphKernelNodeCopyAttributes(kNode2, kernel_vecAdd));
}
SECTION("Copy kernelNodeAttribute to different graph kernel node") {
hipGraphNode_t kNode3;
hipGraph_t graph3;
HIP_CHECK(hipGraphCreate(&graph3, 0));
HIP_CHECK(hipGraphAddKernelNode(&kNode3, graph3,
nullptr, 0, &kNodeParams));
memset(&value_in, 0, sizeof(hipKernelNodeAttrValue));
memset(&value_out, 0, sizeof(hipKernelNodeAttrValue));
HIP_CHECK(hipGraphKernelNodeGetAttribute(kernel_vecAdd,
hipKernelNodeAttributeAccessPolicyWindow, &value_in));
HIP_CHECK(hipGraphKernelNodeCopyAttributes(kernel_vecAdd, kNode3));
HIP_CHECK(hipGraphKernelNodeGetAttribute(kNode3,
hipKernelNodeAttributeAccessPolicyWindow, &value_out));
REQUIRE(true == validateKernelNodeAttrValue(value_in, value_out));
// copy back the node attributes for functional verification only
HIP_CHECK(hipGraphKernelNodeCopyAttributes(kNode3, kernel_vecAdd));
HIP_CHECK(hipGraphDestroy(graph3));
}
SECTION("Copy kernelNodeAttribute to cloned graph kernel node") {
hipGraphNode_t kNode4;
hipGraph_t clonedGraph;
HIP_CHECK(hipGraphClone(&clonedGraph, graph));
HIP_CHECK(hipGraphAddKernelNode(&kNode4, clonedGraph,
nullptr, 0, &kNodeParams));
memset(&value_in, 0, sizeof(hipKernelNodeAttrValue));
memset(&value_out, 0, sizeof(hipKernelNodeAttrValue));
HIP_CHECK(hipGraphKernelNodeGetAttribute(kernel_vecAdd,
hipKernelNodeAttributeAccessPolicyWindow, &value_in));
HIP_CHECK(hipGraphKernelNodeCopyAttributes(kernel_vecAdd, kNode4));
HIP_CHECK(hipGraphKernelNodeGetAttribute(kNode4,
hipKernelNodeAttributeAccessPolicyWindow, &value_out));
REQUIRE(true == validateKernelNodeAttrValue(value_in, value_out));
// copy back the node attributes for functional verification only
HIP_CHECK(hipGraphKernelNodeCopyAttributes(kNode4, kernel_vecAdd));
HIP_CHECK(hipGraphDestroy(clonedGraph));
}
SECTION("Copy kernelNodeAttribute to child graph kernel node") {
hipGraphNode_t kNode5, childGraphNode;
hipGraph_t childGraph;
HIP_CHECK(hipGraphCreate(&childGraph, 0));
HIP_CHECK(hipGraphAddKernelNode(&kNode5, childGraph,
nullptr, 0, &kNodeParams));
HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph,
nullptr, 0, childGraph));
memset(&value_in, 0, sizeof(hipKernelNodeAttrValue));
memset(&value_out, 0, sizeof(hipKernelNodeAttrValue));
HIP_CHECK(hipGraphKernelNodeGetAttribute(kernel_vecAdd,
hipKernelNodeAttributeAccessPolicyWindow, &value_in));
HIP_CHECK(hipGraphKernelNodeCopyAttributes(kernel_vecAdd, kNode5));
HIP_CHECK(hipGraphKernelNodeGetAttribute(kNode5,
hipKernelNodeAttributeAccessPolicyWindow, &value_out));
REQUIRE(true == validateKernelNodeAttrValue(value_in, value_out));
// copy back the node attributes for functional verification only
HIP_CHECK(hipGraphKernelNodeCopyAttributes(kNode5, kernel_vecAdd));
HIP_CHECK(hipGraphDestroy(childGraph));
}
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
@@ -89,12 +203,20 @@ TEST_CASE("Unit_hipGraphKernelNodeCopyAttributes_Functional") {
}
/**
Negative Test for API - hipGraphKernelNodeCopyAttributes
1) Pass source kernel node as nullptr for copy attribute api
2) Pass destination kernel node as nullptr for copy attribute api
3) Pass source kernel node as Uninitialize for copy attribute api
4) Pass dest kernel node as Uninitialize for copy attribute api
*/
* Test Description
* ------------------------
*  - Negative Test for API - hipGraphKernelNodeCopyAttributes
* 1) Pass source kernel node as nullptr for copy attribute api
* 2) Pass destination kernel node as nullptr for copy attribute api
* 3) Pass source kernel node as Uninitialize for copy attribute api
* 4) Pass dest kernel node as Uninitialize for copy attribute api
* Test source
* ------------------------
*  - unit/graph/hipGraphKernelNodeCopyAttributes.cc
* Test requirements
* ------------------------
*  - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipGraphKernelNodeCopyAttributes_Attribute_Negative") {
constexpr size_t N = 1024;
@@ -162,4 +284,3 @@ TEST_CASE("Unit_hipGraphKernelNodeCopyAttributes_Attribute_Negative") {
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
@@ -0,0 +1,195 @@
/*
Copyright (c) 2023 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.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
/**
* @addtogroup hipGraphKernelNodeGetAttribute hipGraphKernelNodeGetAttribute
* @{
* @ingroup GraphTest
* `hipGraphKernelNodeGetAttribute(hipGraphNode_t hNode,
* hipKernelNodeAttrID attr, hipKernelNodeAttrValue* value_out )` -
* Queries node attribute.
*/
/**
* Test Description
* ------------------------
*  - Functional Test for API - hipGraphKernelNodeGetAttribute
* 1) GetKernelAttribute for ID hipKernelNodeAttributeCooperative
* 2) GetKernelAttribute for ID hipKernelNodeAttributeAccessPolicyWindow
* Test source
* ------------------------
*  - unit/graph/hipGraphKernelNodeGetAttribute.cc
* Test requirements
* ------------------------
*  - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipGraphKernelNodeGetAttribute_Functional") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipGraphExec_t graphExec;
hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, kernel_vecAdd;
hipKernelNodeParams kNodeParams{};
hipStream_t stream;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t NElem{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(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
&kNodeParams));
// Create dependencies
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpy_C, 1));
hipKernelNodeAttrValue value_out;
memset(&value_out, 0, sizeof(hipKernelNodeAttrValue));
SECTION("GetKernelAttribute for hipKernelNodeAttributeCooperative") {
HIP_CHECK(hipGraphKernelNodeGetAttribute(kernel_vecAdd,
hipKernelNodeAttributeCooperative, &value_out));
}
SECTION("GetKernelAttribute for hipKernelNodeAttributeAccessPolicyWindow") {
HIP_CHECK(hipGraphKernelNodeGetAttribute(kernel_vecAdd,
hipKernelNodeAttributeAccessPolicyWindow, &value_out));
}
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
// Verify graph execution result
HipTest::checkVectorADD<int>(A_h, B_h, C_h, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream));
}
/**
* Test Description
* ------------------------
*  - Negative Test for API - hipGraphKernelNodeGetAttribute
* 1) Pass kernel node as nullptr for Get attribute api & verify
* 2) Pass KernelNodeAttrID as negative value for Get attribute api & verify
* 3) Pass KernelNodeAttrID as INT_MAX value for Get attribute api & verify
* 4) Pass KernelNodeAttrValue as nullptr for Get attribute api & verify
* Test source
* ------------------------
*  - unit/graph/hipGraphKernelNodeGetAttribute.cc
* Test requirements
* ------------------------
*  - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipGraphKernelNodeGetAttribute_Negative") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, kernel_vecAdd;
hipKernelNodeParams kNodeParams{};
hipStream_t stream;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t NElem{N};
hipError_t ret;
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(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
&kNodeParams));
hipKernelNodeAttrValue value_out;
memset(&value_out, 0, sizeof(hipKernelNodeAttrValue));
SECTION("Pass kernel node as nullptr for Get attribute api") {
ret = hipGraphKernelNodeGetAttribute(nullptr,
hipKernelNodeAttributeAccessPolicyWindow, &value_out);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass KernelNodeAttrID as negative value for Get attribute api") {
ret = hipGraphKernelNodeGetAttribute(kernel_vecAdd,
hipKernelNodeAttrID(-1), &value_out);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass KernelNodeAttrID as INT_MAX value for Get attribute api") {
ret = hipGraphKernelNodeGetAttribute(kernel_vecAdd,
hipKernelNodeAttrID(INT_MAX), &value_out);
REQUIRE(hipErrorInvalidValue == ret);
}
#if HT_AMD // getting SIGSEGV error in Cuda Setup
SECTION("Pass KernelNodeAttrValue as nullptr for Get attribute api") {
ret = hipGraphKernelNodeGetAttribute(kernel_vecAdd,
hipKernelNodeAttributeAccessPolicyWindow, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
}
#endif
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream));
}
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2023 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
@@ -19,7 +19,7 @@ THE SOFTWARE.
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
/* Test verifies hipGraphLaunch API
Negative scenarios -
1) Pass graphExec as nullptr and verify api returns error code.
@@ -32,8 +32,16 @@ Negative scenarios -
Functional Scenario -
1) Check basic functionality with stream as hipStreamPerThread
2) Test hipGraphLaunch call on multiple devices.
3) Create a graph with multiple nodes. Create an executable graph.
Launch the executable graph 3 times in stream simultaneously.
Wait for stream. Validate the output. No issues should be observed
4) Create a graph with multiple nodes. Create an executable graph.
Verify if an executable graph be launched on null stream.
*/
#define SIZE 1024
#define TEST_LOOP_SIZE 3
TEST_CASE("Unit_hipGraphLaunch_Negative") {
hipError_t ret;
SECTION("Pass pGraphExec as nullptr") {
@@ -305,3 +313,100 @@ TEST_CASE("Unit_hipGraphLaunch_Functional_multidevice_test") {
SUCCEED("Skipped the testcase as there is no device to test.");
}
}
// Function to fill input data
static void fillRandInpData(int *A1_h, int *A2_h, size_t N) {
unsigned int seed = time(nullptr);
for (size_t i = 0; i < N; i++) {
A1_h[i] = (HipTest::RAND_R(&seed) & 0xFF);
A2_h[i] = (HipTest::RAND_R(&seed) & 0xFF);
}
}
// Function to validate result
static void validateOutData(int *A1_h, int *A2_h, size_t N) {
for (size_t i = 0; i < N; i++) {
int result = (A1_h[i]*A1_h[i]);
REQUIRE(result == A2_h[i]);
}
}
/*
* 1.Create a graph with multiple nodes. Create an executable graph.
* Launch the executable graph 3 times in stream simultaneously.
* Wait for stream. Validate the output. No issues should be observed
* 2.Create a graph with multiple nodes. Create an executable graph.
* Verify if an executable graph be launched on null stream.
*/
TEST_CASE("Unit_hipGraphLaunch_Functional_MultipleLaunch") {
size_t memSize = SIZE;
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU,
threadsPerBlock, SIZE);
hipGraph_t graph;
std::vector<hipGraphNode_t> nodeDependencies;
HIP_CHECK(hipGraphCreate(&graph, 0));
int *A_h{nullptr}, *A_d{nullptr}, *C_d{nullptr}, *C_h{nullptr};
HipTest::initArrays<int>(&A_d, &C_d, nullptr,
&A_h, &C_h, nullptr, SIZE, false);
hipGraphNode_t memcpyH2D, memcpyD2H, kernelNode;
// Create memcpy H2D nodes
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr,
0, A_d, A_h, (sizeof(int)*SIZE), hipMemcpyHostToDevice));
nodeDependencies.push_back(memcpyH2D);
// Creating kernel node
hipKernelNodeParams kerNodeParams;
void* kernelArgs[] = {reinterpret_cast<void*>(&A_d),
reinterpret_cast<void*>(&C_d),
reinterpret_cast<void*>(&memSize)};
kerNodeParams.func = reinterpret_cast<void*>(HipTest::vector_square<int>);
kerNodeParams.gridDim = dim3(blocks);
kerNodeParams.blockDim = dim3(threadsPerBlock);
kerNodeParams.sharedMemBytes = 0;
kerNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kerNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nodeDependencies.data(),
nodeDependencies.size(), &kerNodeParams));
nodeDependencies.clear();
nodeDependencies.push_back(kernelNode);
// Create memcpy D2H nodes
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, nodeDependencies.data(),
nodeDependencies.size(), C_h, C_d, (sizeof(int)*SIZE),
hipMemcpyDeviceToHost));
nodeDependencies.clear();
// Create executable graph
hipStream_t streamForGraph;
hipGraphExec_t graphExec{nullptr};
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
nullptr, 0));
// Execute graph
SECTION("Multiple Graph Launch") {
for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) {
fillRandInpData(A_h, C_h, SIZE);
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
validateOutData(A_h, C_h, SIZE);
}
}
SECTION("Graph launch on Null stream") {
for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) {
fillRandInpData(A_h, C_h, SIZE);
HIP_CHECK(hipGraphLaunch(graphExec, 0));
HIP_CHECK(hipStreamSynchronize(0));
validateOutData(A_h, C_h, SIZE);
}
}
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipStreamDestroy(streamForGraph));
// Free
HipTest::freeArrays<int>(A_d, C_d, nullptr, A_h, C_h, nullptr, false);
}
@@ -0,0 +1,62 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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.
*/
#include <hip_test_common.hh>
#include "user_object_common.hh"
/**
* Negative Test for API - hipGraphReleaseUserObject
1) Pass graph as nullptr
2) Pass User Object as nullptr
3) Pass initialRefcount as 0
4) Pass initialRefcount as INT_MAX
*/
TEST_CASE("Unit_hipGraphReleaseUserObject_Negative") {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
float* object = new float();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(
hipUserObjectCreate(&hObject, object, destroyFloatObj, 1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 1, hipGraphUserObjectMove));
SECTION("Pass graph as nullptr") {
HIP_CHECK_ERROR(hipGraphReleaseUserObject(nullptr, hObject, 1), hipErrorInvalidValue);
}
SECTION("Pass User Object as nullptr") {
HIP_CHECK_ERROR(hipGraphReleaseUserObject(graph, nullptr, 1), hipErrorInvalidValue);
}
SECTION("Pass initialRefcount as 0") {
HIP_CHECK_ERROR(hipGraphReleaseUserObject(graph, hObject, 0), hipErrorInvalidValue);
}
SECTION("Pass initialRefcount as INT_MAX") {
HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, INT_MAX));
}
HIP_CHECK(hipUserObjectRelease(hObject, 1));
HIP_CHECK(hipGraphDestroy(graph));
}
@@ -0,0 +1,243 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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.
*/
#include <hip_test_checkers.hh>
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include "user_object_common.hh"
/**
* Functional Test for API - hipGraphRetainUserObject
*/
/* 1) Create GraphUserObject and retain it by calling hipGraphRetainUserObject
and release it by calling hipGraphReleaseUserObject. */
static void hipGraphRetainUserObject_Functional_1(void* object, void destroyObj(void*)) {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object, destroyObj, 1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 1, hipGraphUserObjectMove));
HIP_CHECK(hipGraphReleaseUserObject(graph, hObject));
HIP_CHECK(hipUserObjectRelease(hObject));
HIP_CHECK(hipGraphDestroy(graph));
}
TEST_CASE("Unit_hipGraphRetainUserObject_Functional_1") {
SECTION("Called with int Object") {
int* object = new int();
REQUIRE(object != nullptr);
hipGraphRetainUserObject_Functional_1(object, destroyIntObj);
}
SECTION("Called with float Object") {
float* object = new float();
REQUIRE(object != nullptr);
hipGraphRetainUserObject_Functional_1(object, destroyFloatObj);
}
SECTION("Called with Class Object") {
BoxClass* object = new BoxClass();
REQUIRE(object != nullptr);
hipGraphRetainUserObject_Functional_1(object, destroyClassObj);
}
SECTION("Called with Struct Object") {
BoxStruct* object = new BoxStruct();
REQUIRE(object != nullptr);
hipGraphRetainUserObject_Functional_1(object, destroyStructObj);
}
}
/* 2) Create UserObject and GraphUserObject and retain using custom reference
count and release it by calling hipGraphReleaseUserObject with count. */
TEST_CASE("Unit_hipGraphRetainUserObject_Functional_2") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipGraphNode_t memcpyNode, kNode;
hipKernelNodeParams kNodeParams{};
hipStream_t streamForGraph;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
std::vector<hipGraphNode_t> dependencies;
hipGraphExec_t graphExec;
size_t NElem{N};
HIP_CHECK(hipStreamCreate(&streamForGraph));
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(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, Nbytes,
hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h, Nbytes,
hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void*>(&NElem)};
kNodeParams.func = reinterpret_cast<void*>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
HIP_CHECK(
hipGraphAddKernelNode(&kNode, graph, dependencies.data(), dependencies.size(), &kNodeParams));
dependencies.clear();
dependencies.push_back(kNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), dependencies.size(),
C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
int refCount = 2;
int refCountRetain = 3;
float* object = new float();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object, destroyFloatObj, refCount,
hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
HIP_CHECK(hipUserObjectRetain(hObject, refCountRetain));
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, refCountRetain, hipGraphUserObjectMove));
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verify result
HipTest::checkVectorADD<int>(A_h, B_h, C_h, N);
HIP_CHECK(hipUserObjectRelease(hObject, refCount + refCountRetain));
HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, refCountRetain));
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
/**
* Negative Test for API - hipGraphRetainUserObject
1) Pass graph as nullptr
2) Pass User Object as nullptr
3) Pass initialRefcount as 0
4) Pass initialRefcount as INT_MAX
5) Pass flag as 0
6) Pass flag as INT_MAX
*/
TEST_CASE("Unit_hipGraphRetainUserObject_Negative") {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
float* object = new float();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(
hipUserObjectCreate(&hObject, object, destroyFloatObj, 1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
SECTION("Pass graph as nullptr") {
HIP_CHECK_ERROR(hipGraphRetainUserObject(nullptr, hObject, 1, hipGraphUserObjectMove),
hipErrorInvalidValue);
}
SECTION("Pass User Object as nullptr") {
HIP_CHECK_ERROR(hipGraphRetainUserObject(graph, nullptr, 1, hipGraphUserObjectMove),
hipErrorInvalidValue);
}
SECTION("Pass initialRefcount as 0") {
HIP_CHECK_ERROR(hipGraphRetainUserObject(graph, hObject, 0, hipGraphUserObjectMove),
hipErrorInvalidValue);
}
SECTION("Pass initialRefcount as INT_MAX") {
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, INT_MAX, hipGraphUserObjectMove));
}
SECTION("Pass flag as 0") { HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 1, 0)); }
SECTION("Pass flag as INT_MAX") {
HIP_CHECK_ERROR(hipGraphRetainUserObject(graph, hObject, 1, INT_MAX), hipErrorInvalidValue);
}
HIP_CHECK(hipUserObjectRelease(hObject, 1));
HIP_CHECK(hipGraphDestroy(graph));
}
TEST_CASE("Unit_hipGraphRetainUserObject_Negative_Basic") {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
float* object = new float();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(
hipUserObjectCreate(&hObject, object, destroyFloatObj, 1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
// Retain graph object with reference count 2
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 2, hipGraphUserObjectMove));
// Release graph object with reference count more than 2
HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, 4));
// Again Retain graph object with reference count 8
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 8, hipGraphUserObjectMove));
// Release graph object with reference count 1
HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, 1));
HIP_CHECK(hipUserObjectRelease(hObject, 1));
HIP_CHECK(hipGraphDestroy(graph));
}
TEST_CASE("Unit_hipGraphRetainUserObject_Negative_Null_Object") {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
float* object = nullptr; // this is used for Null_Object test
hipUserObject_t hObject;
HIP_CHECK(
hipUserObjectCreate(&hObject, object, destroyFloatObj, 1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
// Retain graph object with reference count 2
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 2, hipGraphUserObjectMove));
// Release graph object with reference count more than 2
HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, 4));
// Again Retain graph object with reference count 8
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 8, 0));
// Release graph object with reference count 1
HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, 1));
HIP_CHECK(hipUserObjectRelease(hObject, 1));
HIP_CHECK(hipGraphDestroy(graph));
}
@@ -0,0 +1,58 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/*
This code object should be automatically built via "make build_tests".
In case it's missing, please type the following to generate it,
/opt/rocm/hip/bin/hipcc --genco hipMatMul.cc -o hipMatMul.code
*/
#include"hip/hip_runtime.h"
__device__ int deviceGlobal = 1;
extern "C" __global__ void matmulK(int* A, int* B, int* C,
int N) {
int ROW = blockIdx.y*blockDim.y+threadIdx.y;
int COL = blockIdx.x*blockDim.x+threadIdx.x;
int tmpSum = 0;
if ((ROW < N) && (COL < N)) {
// each thread computes one element of the block sub-matrix
for (int i = 0; i < N; i++) {
tmpSum += A[ROW * N + i] * B[i * N + COL];
}
C[ROW * N + COL] = tmpSum;
}
}
extern "C" __global__ void KernelandExtraParams(int* A, int* B, int* C,
int *D, int N) {
int ROW = blockIdx.y*blockDim.y+threadIdx.y;
int COL = blockIdx.x*blockDim.x+threadIdx.x;
int tmpSum = 0;
if (ROW < N && COL < N) {
// each thread computes one element of the block sub-matrix
for (int i = 0; i < N; i++) {
tmpSum += A[ROW * N + i] * B[i * N + COL];
}
}
C[ROW * N + COL] = tmpSum;
D[ROW * N + COL] = tmpSum;
}
extern "C" __global__ void dummyKernel() {
}
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2022 - 2023 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
@@ -20,8 +20,7 @@ THE SOFTWARE.
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_defgroups.hh>
#include "stream_capture_common.hh"
#include "stream_capture_common.hh" // NOLINT
/**
* @addtogroup hipStreamBeginCapture hipStreamBeginCapture
@@ -56,7 +55,8 @@ static void hostNodeCallback(void* data) {
}
template <typename T, typename F>
void captureStreamAndLaunchGraph(F graphFunc, hipStreamCaptureMode mode, hipStream_t stream) {
void captureStreamAndLaunchGraph(F graphFunc, hipStreamCaptureMode mode,
hipStream_t stream) {
constexpr size_t N = 1000000;
size_t Nbytes = N * sizeof(T);
@@ -88,7 +88,8 @@ void captureStreamAndLaunchGraph(F graphFunc, hipStreamCaptureMode mode, hipStre
std::fill_n(A_h.host_ptr(), N, static_cast<float>(i));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
ArrayFindIfNot(B_h.host_ptr(), static_cast<float>(i) * static_cast<float>(i), N);
ArrayFindIfNot(B_h.host_ptr(),
static_cast<float>(i) * static_cast<float>(i), N);
}
HIP_CHECK(hipGraphExecDestroy(graphExec))
@@ -115,15 +116,16 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Functional") {
StreamGuard stream_guard(stream_type);
hipStream_t stream = stream_guard.stream();
const hipStreamCaptureMode captureMode = GENERATE(
hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal,
hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
EventsGuard events_guard(3);
StreamsGuard streams_guard(2);
SECTION("Linear graph capture") {
captureStreamAndLaunchGraph<float>(
[](float* A_h, float* A_d, float* B_h, float* B_d, size_t N, hipStream_t stream) {
[](float* A_h, float* A_d, float* B_h, float* B_d, size_t N,
hipStream_t stream) {
return captureSequenceLinear(A_h, A_d, B_h, B_d, N, stream);
},
captureMode, stream);
@@ -131,10 +133,10 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Functional") {
SECTION("Branched graph capture") {
captureStreamAndLaunchGraph<float>(
[&streams_guard, &events_guard](float* A_h, float* A_d, float* B_h, float* B_d, size_t N,
hipStream_t stream) {
captureSequenceBranched(A_h, A_d, B_h, B_d, N, stream, streams_guard.stream_list(),
events_guard.event_list());
[&streams_guard, &events_guard](float* A_h, float* A_d, float* B_h,
float* B_d, size_t N, hipStream_t stream) {
captureSequenceBranched(A_h, A_d, B_h, B_d, N, stream,
streams_guard.stream_list(), events_guard.event_list());
},
captureMode, stream);
}
@@ -170,7 +172,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_Parameters") {
hipErrorIllegalState);
}
SECTION("Creating hipStream with invalid mode") {
HIP_CHECK_ERROR(hipStreamBeginCapture(stream, hipStreamCaptureMode(-1)), hipErrorInvalidValue);
HIP_CHECK_ERROR(hipStreamBeginCapture(stream, hipStreamCaptureMode(-1)),
hipErrorInvalidValue);
}
#if HT_NVIDIA // EXSWHTEC-216
SECTION("Stream capture on uninitialized stream returns error code.") {
@@ -178,7 +181,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_Parameters") {
StreamGuard sg(Streams::created);
return sg.stream();
};
HIP_CHECK_ERROR(hipStreamBeginCapture(InvalidStream(), hipStreamCaptureModeGlobal),
HIP_CHECK_ERROR(hipStreamBeginCapture(InvalidStream(),
hipStreamCaptureModeGlobal),
hipErrorContextIsDestroyed);
}
#endif
@@ -202,8 +206,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Basic") {
StreamGuard stream_guard(stream_type);
hipStream_t s = stream_guard.stream();
const hipStreamCaptureMode captureMode = GENERATE(
hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal,
hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
HIP_CHECK(hipStreamBeginCapture(s, captureMode));
@@ -213,7 +217,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Basic") {
/* Local function for inter stream event synchronization
*/
static void interStrmEventSyncCapture(const hipStream_t& stream1, const hipStream_t& stream2) {
static void interStrmEventSyncCapture(const hipStream_t& stream1,
const hipStream_t& stream2) {
hipGraph_t graph1{nullptr}, graph2{nullptr};
hipGraphExec_t graphExec1{nullptr}, graphExec2{nullptr};
@@ -260,7 +265,8 @@ static void interStrmEventSyncCapture(const hipStream_t& stream1, const hipStrea
/* Local function for colligated stream capture
*/
static void colligatedStrmCapture(const hipStream_t& stream1, const hipStream_t& stream2) {
static void colligatedStrmCapture(const hipStream_t& stream1,
const hipStream_t& stream2) {
hipGraph_t graph1{nullptr}, graph2{nullptr};
hipGraphExec_t graphExec1{nullptr}, graphExec2{nullptr};
@@ -303,7 +309,8 @@ static void colligatedStrmCapture(const hipStream_t& stream1, const hipStream_t&
/* Local function for colligated stream capture functionality
*/
static void colligatedStrmCaptureFunc(const hipStream_t& stream1, const hipStream_t& stream2) {
static void colligatedStrmCaptureFunc(const hipStream_t& stream1,
const hipStream_t& stream2) {
constexpr size_t N = 1000000;
size_t Nbytes = N * sizeof(int);
@@ -323,8 +330,10 @@ static void colligatedStrmCaptureFunc(const hipStream_t& stream1, const hipStrea
// Capture 2 streams
HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
HIP_CHECK(hipStreamBeginCapture(stream2, hipStreamCaptureModeGlobal));
captureSequenceLinear(A_h.host_ptr(), A_d.ptr(), B_h.host_ptr(), B_d.ptr(), N, stream1);
captureSequenceLinear(C_h.host_ptr(), C_d.ptr(), D_h.host_ptr(), D_d.ptr(), N, stream2);
captureSequenceLinear(A_h.host_ptr(), A_d.ptr(), B_h.host_ptr(), B_d.ptr(),
N, stream1);
captureSequenceLinear(C_h.host_ptr(), C_d.ptr(), D_h.host_ptr(), D_d.ptr(),
N, stream2);
captureSequenceCompute(A_d.ptr(), B_h.host_ptr(), B_d.ptr(), N, stream1);
captureSequenceCompute(C_d.ptr(), D_h.host_ptr(), D_d.ptr(), N, stream2);
HIP_CHECK(hipStreamEndCapture(stream1, &graph1));
@@ -360,8 +369,9 @@ static void colligatedStrmCaptureFunc(const hipStream_t& stream1, const hipStrea
/* Stream Capture thread function
*/
static void threadStrmCaptureFunc(hipStream_t stream, int* A_h, int* A_d, int* B_h, int* B_d,
hipGraph_t* graph, size_t N, hipStreamCaptureMode mode) {
static void threadStrmCaptureFunc(hipStream_t stream, int* A_h, int* A_d,
int* B_h, int* B_d, hipGraph_t* graph,
size_t N, hipStreamCaptureMode mode) {
// Capture stream
HIP_CHECK(hipStreamBeginCapture(stream, mode));
captureSequenceLinear(A_h, A_d, B_h, B_d, N, stream);
@@ -393,10 +403,10 @@ static void multithreadedTest(hipStreamCaptureMode mode) {
LinearAllocGuard<int> D_d(LinearAllocs::hipMalloc, Nbytes);
// Launch 2 threads to capture the 2 streams into graphs
std::thread t1(threadStrmCaptureFunc, stream1, A_h.host_ptr(), A_d.ptr(), B_h.host_ptr(),
B_d.ptr(), &graph1, N, mode);
std::thread t2(threadStrmCaptureFunc, stream2, C_h.host_ptr(), C_d.ptr(), D_h.host_ptr(),
D_d.ptr(), &graph2, N, mode);
std::thread t1(threadStrmCaptureFunc, stream1, A_h.host_ptr(), A_d.ptr(),
B_h.host_ptr(), B_d.ptr(), &graph1, N, mode);
std::thread t2(threadStrmCaptureFunc, stream2, C_h.host_ptr(), C_d.ptr(),
D_h.host_ptr(), D_d.ptr(), &graph2, N, mode);
t1.join();
t2.join();
@@ -469,9 +479,11 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_InterStrmEventSync_Flags") {
TEST_CASE("Unit_hipStreamBeginCapture_Positive_InterStrmEventSync_Priority") {
int minPriority = 0, maxPriority = 0;
HIP_CHECK(hipDeviceGetStreamPriorityRange(&minPriority, &maxPriority));
StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault, minPriority);
StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault,
minPriority);
hipStream_t stream1 = stream_guard1.stream();
StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault, maxPriority);
StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault,
maxPriority);
hipStream_t stream2 = stream_guard2.stream();
interStrmEventSyncCapture(stream1, stream2);
}
@@ -517,12 +529,14 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCapture_Flags") {
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCapture_Priority") {
TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCapture_Prio") {
int minPriority = 0, maxPriority = 0;
HIP_CHECK(hipDeviceGetStreamPriorityRange(&minPriority, &maxPriority));
StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault, minPriority);
StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault,
minPriority);
hipStream_t stream1 = stream_guard1.stream();
StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault, maxPriority);
StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault,
maxPriority);
hipStream_t stream2 = stream_guard2.stream();
colligatedStrmCapture(stream1, stream2);
}
@@ -563,8 +577,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCaptureFunc") {
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamBeginCapture_Positive_Multithreaded") {
const hipStreamCaptureMode captureMode = GENERATE(
hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal,
hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
multithreadedTest(captureMode);
}
@@ -693,7 +707,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_CapturingFromWithinStrms") {
HIP_CHECK(hipEventRecord(events[2], streams[2]));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[1], 0));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[2], 0));
HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault, streams[0]));
HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault,
streams[0]));
HIP_CHECK(hipStreamEndCapture(streams[0], &graph)); // End Capture
// Reset device memory
HIP_CHECK(hipMemset(devMem, 0, sizeof(int)));
@@ -735,7 +750,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_DetectingInvalidCapture") {
dummyKernel<<<1, 1, 0, streams[0]>>>();
// Since stream[1] is already in capture mode due to event wait
// hipStreamBeginCapture on stream[1] is expected to return error.
HIP_CHECK_ERROR(hipStreamBeginCapture(streams[1], hipStreamCaptureModeGlobal),
HIP_CHECK_ERROR(hipStreamBeginCapture(streams[1],
hipStreamCaptureModeGlobal),
hipErrorIllegalState);
}
@@ -768,7 +784,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_CapturingMultGraphsFrom1Strm") {
for (int i = 0; i < 3; i++) {
HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
for (int j = 0; j <= i; j++) incrementKernel<<<1, 1, 0, stream1>>>(devMem);
HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault, stream1));
HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault,
stream1));
HIP_CHECK(hipStreamEndCapture(stream1, &graphs[i]));
}
// Instantiate and execute all graphs
@@ -784,7 +801,6 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_CapturingMultGraphsFrom1Strm") {
}
}
#if HT_NVIDIA
/**
* Test Description
* ------------------------
@@ -808,29 +824,35 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_CheckingSyncDuringCapture") {
EventsGuard events_guard(1);
hipEvent_t e = events_guard[0];
const hipStreamCaptureMode captureMode = GENERATE(
hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal,
hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
HIP_CHECK(hipStreamBeginCapture(stream, captureMode));
SECTION("Synchronize stream during capture") {
HIP_CHECK_ERROR(hipStreamSynchronize(stream), hipErrorStreamCaptureUnsupported);
HIP_CHECK_ERROR(hipStreamSynchronize(stream),
hipErrorStreamCaptureUnsupported);
}
SECTION("Query stream during capture") {
HIP_CHECK_ERROR(hipStreamQuery(stream),
hipErrorStreamCaptureUnsupported);
}
#if HT_NVIDIA
SECTION("Synchronize device during capture") {
HIP_CHECK_ERROR(hipDeviceSynchronize(), hipErrorStreamCaptureUnsupported);
HIP_CHECK_ERROR(hipDeviceSynchronize(),
hipErrorStreamCaptureUnsupported);
}
SECTION("Synchronize event during capture") {
HIP_CHECK(hipEventRecord(e, stream));
HIP_CHECK_ERROR(hipEventSynchronize(e), hipErrorCapturedEvent);
}
SECTION("Query stream during capture") {
HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorStreamCaptureUnsupported);
}
SECTION("Query for an event during capture") {
HIP_CHECK(hipEventRecord(e, stream));
HIP_CHECK_ERROR(hipEventQuery(e), hipErrorCapturedEvent);
}
#endif
}
#if HT_NVIDIA
/**
* Test Description
* ------------------------
@@ -861,14 +883,17 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_UnsafeCallsDuringCapture") {
HIP_CHECK(hipStreamBeginCapture(stream, captureMode));
SECTION("hipMalloc during capture") {
HIP_CHECK_ERROR(hipMalloc(&devMem2, sizeof(int)), hipErrorStreamCaptureUnsupported);
HIP_CHECK_ERROR(hipMalloc(&devMem2, sizeof(int)),
hipErrorStreamCaptureUnsupported);
}
SECTION("hipMemcpy during capture") {
HIP_CHECK_ERROR(hipMemcpy(devMem.ptr(), hostMem.host_ptr(), sizeof(int), hipMemcpyHostToDevice),
HIP_CHECK_ERROR(hipMemcpy(devMem.ptr(), hostMem.host_ptr(), sizeof(int),
hipMemcpyHostToDevice),
hipErrorStreamCaptureImplicit);
}
SECTION("hipMemset during capture") {
HIP_CHECK_ERROR(hipMemset(devMem.ptr(), 0, sizeof(int)), hipErrorStreamCaptureImplicit);
HIP_CHECK_ERROR(hipMemset(devMem.ptr(), 0, sizeof(int)),
hipErrorStreamCaptureImplicit);
}
}
#endif
@@ -889,7 +914,7 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_UnsafeCallsDuringCapture") {
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapturewhenCaptureInProgress") {
TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapwhenCapInProg") {
hipGraph_t graph{nullptr};
StreamsGuard streams_guard(2);
@@ -905,7 +930,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapturewhenCaptureInProgres
HIP_CHECK(hipEventRecord(e, stream1));
HIP_CHECK(hipStreamWaitEvent(stream2, e, 0));
dummyKernel<<<1, 1, 0, stream2>>>();
HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph), hipErrorStreamCaptureUnjoined);
HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph),
hipErrorStreamCaptureUnjoined);
}
SECTION("End strm capture when forked strm still has operations") {
EventsGuard events_guard(2);
@@ -919,7 +945,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapturewhenCaptureInProgres
HIP_CHECK(hipEventRecord(e2, stream2));
HIP_CHECK(hipStreamWaitEvent(stream1, e2, 0));
dummyKernel<<<1, 1, 0, stream2>>>();
HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph), hipErrorStreamCaptureUnjoined);
HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph),
hipErrorStreamCaptureUnjoined);
}
}
/**
@@ -942,17 +969,19 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_MultiGPU") {
SUCCEED("skipping the testcases as numDevices < 2");
return;
}
hipStream_t* stream = reinterpret_cast<hipStream_t*>(malloc(devcount * sizeof(hipStream_t)));
hipStream_t* stream = reinterpret_cast<hipStream_t*>
(malloc(devcount * sizeof(hipStream_t)));
REQUIRE(stream != nullptr);
hipGraph_t* graph = reinterpret_cast<hipGraph_t*>(malloc(devcount * sizeof(hipGraph_t)));
hipGraph_t* graph = reinterpret_cast<hipGraph_t*>
(malloc(devcount * sizeof(hipGraph_t)));
REQUIRE(graph != nullptr);
int **devMem{nullptr}, **hostMem{nullptr};
hostMem = reinterpret_cast<int**>(malloc(sizeof(int*) * devcount));
REQUIRE(hostMem != nullptr);
devMem = reinterpret_cast<int**>(malloc(sizeof(int*) * devcount));
REQUIRE(devMem != nullptr);
hipGraphExec_t* graphExec =
reinterpret_cast<hipGraphExec_t*>(malloc(devcount * sizeof(hipGraphExec_t)));
hipGraphExec_t* graphExec = reinterpret_cast<hipGraphExec_t*>
(malloc(devcount * sizeof(hipGraphExec_t)));
// Capture stream in each device
for (int dev = 0; dev < devcount; dev++) {
HIP_CHECK(hipSetDevice(dev));
@@ -964,14 +993,15 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_MultiGPU") {
for (int i = 0; i < (dev + 1); i++) {
incrementKernel<<<1, 1, 0, stream[dev]>>>(devMem[dev]);
}
HIP_CHECK(
hipMemcpyAsync(hostMem[dev], devMem[dev], sizeof(int), hipMemcpyDefault, stream[dev]));
HIP_CHECK(hipMemcpyAsync(hostMem[dev], devMem[dev], sizeof(int),
hipMemcpyDefault, stream[dev]));
HIP_CHECK(hipStreamEndCapture(stream[dev], &graph[dev]));
}
// Launch the captured graphs in the respective device
for (int dev = 0; dev < devcount; dev++) {
HIP_CHECK(hipSetDevice(dev));
HIP_CHECK(hipGraphInstantiate(&graphExec[dev], graph[dev], nullptr, nullptr, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec[dev], graph[dev], nullptr,
nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec[dev], stream[dev]));
}
// Validate output
@@ -1038,8 +1068,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_nestedStreamCapture") {
HIP_CHECK(hipEventRecord(events[3], streams[2]));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[3], 0));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[2], 0));
HIP_CHECK(hipMemcpyAsync(hostMem_g.host_ptr(), devMem_g.ptr(), sizeof(int), hipMemcpyDefault,
streams[0]));
HIP_CHECK(hipMemcpyAsync(hostMem_g.host_ptr(), devMem_g.ptr(), sizeof(int),
hipMemcpyDefault, streams[0]));
HIP_CHECK(hipStreamEndCapture(streams[0], &graph)); // End Capture
// Reset device memory
HIP_CHECK(hipMemset(devMem_g.ptr(), 0, sizeof(int)));
@@ -1077,15 +1107,23 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_streamReuse") {
hipGraph_t graphs[3];
StreamsGuard streams(3);
EventsGuard events(4);
LinearAllocGuard<int> hostMem_g1 = LinearAllocGuard<int>(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> hostMem_g2 = LinearAllocGuard<int>(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> hostMem_g3 = LinearAllocGuard<int>(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> devMem_g1 = LinearAllocGuard<int>(LinearAllocs::hipMalloc, sizeof(int));
LinearAllocGuard<int> devMem_g2 = LinearAllocGuard<int>(LinearAllocs::hipMalloc, sizeof(int));
LinearAllocGuard<int> devMem_g3 = LinearAllocGuard<int>(LinearAllocs::hipMalloc, sizeof(int));
LinearAllocGuard<int> hostMem_g1 = LinearAllocGuard<int>
(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> hostMem_g2 = LinearAllocGuard<int>
(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> hostMem_g3 = LinearAllocGuard<int>
(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> devMem_g1 = LinearAllocGuard<int>
(LinearAllocs::hipMalloc, sizeof(int));
LinearAllocGuard<int> devMem_g2 = LinearAllocGuard<int>
(LinearAllocs::hipMalloc, sizeof(int));
LinearAllocGuard<int> devMem_g3 = LinearAllocGuard<int>
(LinearAllocs::hipMalloc, sizeof(int));
std::vector<int*> hostMem = {hostMem_g1.host_ptr(), hostMem_g2.host_ptr(), hostMem_g3.host_ptr()};
std::vector<int*> devMem = {devMem_g1.ptr(), devMem_g2.ptr(), devMem_g3.ptr()};
std::vector<int*> hostMem = {hostMem_g1.host_ptr(), hostMem_g2.host_ptr(),
hostMem_g3.host_ptr()};
std::vector<int*> devMem = {devMem_g1.ptr(), devMem_g2.ptr(),
devMem_g3.ptr()};
// Create a device memory of size int and initialize it to 0
for (int i = 0; i < 3; i++) {
memset(hostMem[i], 0, sizeof(int));
@@ -1109,14 +1147,16 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_streamReuse") {
HIP_CHECK(hipEventRecord(events[3], streams[2]));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[3], 0));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[2], 0));
HIP_CHECK(hipMemcpyAsync(hostMem[0], devMem[0], sizeof(int), hipMemcpyDefault, streams[0]));
HIP_CHECK(hipMemcpyAsync(hostMem[0], devMem[0], sizeof(int),
hipMemcpyDefault, streams[0]));
HIP_CHECK(hipStreamEndCapture(streams[0], &graphs[0])); // End Capture
// Start capturing graph2 from stream 2
HIP_CHECK(hipStreamBeginCapture(streams[1], hipStreamCaptureModeGlobal));
incrementKernel<<<1, 1, 0, streams[1]>>>(devMem[1]);
incrementKernel<<<1, 1, 0, streams[1]>>>(devMem[1]);
incrementKernel<<<1, 1, 0, streams[1]>>>(devMem[1]);
HIP_CHECK(hipMemcpyAsync(hostMem[1], devMem[1], sizeof(int), hipMemcpyDefault, streams[1]));
HIP_CHECK(hipMemcpyAsync(hostMem[1], devMem[1], sizeof(int),
hipMemcpyDefault, streams[1]));
HIP_CHECK(hipStreamEndCapture(streams[1], &graphs[1])); // End Capture
// Start capturing graph3 from stream 3
HIP_CHECK(hipStreamBeginCapture(streams[2], hipStreamCaptureModeGlobal));
@@ -1125,7 +1165,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_streamReuse") {
incrementKernel<<<1, 1, 0, streams[2]>>>(devMem[2]);
incrementKernel<<<1, 1, 0, streams[2]>>>(devMem[2]);
incrementKernel<<<1, 1, 0, streams[2]>>>(devMem[2]);
HIP_CHECK(hipMemcpyAsync(hostMem[2], devMem[2], sizeof(int), hipMemcpyDefault, streams[2]));
HIP_CHECK(hipMemcpyAsync(hostMem[2], devMem[2], sizeof(int),
hipMemcpyDefault, streams[2]));
HIP_CHECK(hipStreamEndCapture(streams[2], &graphs[2])); // End Capture
// Reset device memory
HIP_CHECK(hipMemset(devMem[0], 0, sizeof(int)));
@@ -1169,32 +1210,40 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureComplexGraph") {
EventsGuard events(7);
// Allocate Device memory and Host memory
size_t N = GRIDSIZE * BLOCKSIZE;
LinearAllocGuard<int> Ah = LinearAllocGuard<int>(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Bh = LinearAllocGuard<int>(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Ch = LinearAllocGuard<int>(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Ad = LinearAllocGuard<int>(LinearAllocs::hipMalloc, N * sizeof(int));
LinearAllocGuard<int> Bd = LinearAllocGuard<int>(LinearAllocs::hipMalloc, N * sizeof(int));
LinearAllocGuard<int> Ah = LinearAllocGuard<int>
(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Bh = LinearAllocGuard<int>
(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Ch = LinearAllocGuard<int>
(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Ad = LinearAllocGuard<int>
(LinearAllocs::hipMalloc, N * sizeof(int));
LinearAllocGuard<int> Bd = LinearAllocGuard<int>
(LinearAllocs::hipMalloc, N * sizeof(int));
// Capture streams into graph
HIP_CHECK(hipStreamBeginCapture(streams[0], hipStreamCaptureModeGlobal));
HIP_CHECK(hipEventRecord(events[0], streams[0]));
HIP_CHECK(hipStreamWaitEvent(streams[3], events[0], 0));
HIP_CHECK(hipStreamWaitEvent(streams[4], events[0], 0));
HIP_CHECK(
hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), (N * sizeof(int)), hipMemcpyDefault, streams[0]));
HIP_CHECK(
hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), (N * sizeof(int)), hipMemcpyDefault, streams[4]));
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), (N * sizeof(int)),
hipMemcpyDefault, streams[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), (N * sizeof(int)),
hipMemcpyDefault, streams[4]));
hipHostFn_t fn = hostNodeCallback;
HIPCHECK(hipLaunchHostFunc(streams[3], fn, nullptr));
HIP_CHECK(hipEventRecord(events[1], streams[0]));
HIP_CHECK(hipStreamWaitEvent(streams[1], events[1], 0));
int* Ad_2nd_half = Ad.ptr() + N / 2;
int* Ad_1st_half = Ad.ptr();
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[0]>>>(Ad_2nd_half, CONST_KER2_VAL);
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[1]>>>(Ad_1st_half, CONST_KER1_VAL);
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[0]>>>(Ad_2nd_half,
CONST_KER2_VAL);
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[1]>>>(Ad_1st_half,
CONST_KER1_VAL);
HIP_CHECK(hipEventRecord(events[2], streams[1]));
HIP_CHECK(hipStreamWaitEvent(streams[2], events[2], 0));
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[1]>>>(Ad_1st_half, CONST_KER3_VAL);
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[1]>>>(Ad_1st_half,
CONST_KER3_VAL);
HIPCHECK(hipLaunchHostFunc(streams[2], fn, nullptr));
HIP_CHECK(hipEventRecord(events[6], streams[1]));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[6], 0));
@@ -1205,8 +1254,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureComplexGraph") {
HIP_CHECK(hipStreamWaitEvent(streams[0], events[3], 0));
HIP_CHECK(hipEventRecord(events[4], streams[3]));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[4], 0));
HIP_CHECK(
hipMemcpyAsync(Ch.host_ptr(), Ad.ptr(), (N * sizeof(int)), hipMemcpyDefault, streams[0]));
HIP_CHECK(hipMemcpyAsync(Ch.host_ptr(), Ad.ptr(), (N * sizeof(int)),
hipMemcpyDefault, streams[0]));
HIP_CHECK(hipStreamEndCapture(streams[0], &graph)); // End Capture
// Execute and test the graph
hipGraphExec_t graphExec{nullptr};
@@ -1219,10 +1268,11 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureComplexGraph") {
HIP_CHECK(hipStreamSynchronize(streams[0]));
for (size_t i = 0; i < N; i++) {
if (i > (N / 2 - 1)) {
REQUIRE(Ch.host_ptr()[i] == (Bh.host_ptr()[i] + Ah.host_ptr()[i] * CONST_KER2_VAL));
REQUIRE(Ch.host_ptr()[i] == (Bh.host_ptr()[i] +
Ah.host_ptr()[i] * CONST_KER2_VAL));
} else {
REQUIRE(Ch.host_ptr()[i] ==
(Bh.host_ptr()[i] + Ah.host_ptr()[i] * CONST_KER1_VAL * CONST_KER3_VAL));
REQUIRE(Ch.host_ptr()[i] == (Bh.host_ptr()[i] +
Ah.host_ptr()[i] * CONST_KER1_VAL * CONST_KER3_VAL));
}
}
}
@@ -1267,3 +1317,226 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureEmptyStreams") {
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Test Description
* ------------------------
* - Test to verify hipStreamSynchronize on a stream works when stream capture
* on another stream is ongoing.
* Test source
* ------------------------
* - catch\unit\graph\hipStreamBeginCapture.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipStreamBeginCapture_StreamSync_OngoingCapture") {
hipStreamCaptureMode flag = hipStreamCaptureModeRelaxed;
constexpr int GRIDSIZE = 1;
constexpr int BLOCKSIZE = 512;
constexpr int VALUE1 = 7, VALUE2 = 11;
hipGraph_t graph{nullptr};
hipGraphExec_t graphExec{nullptr};
// Allocate device memory
LinearAllocGuard<int> Ah = LinearAllocGuard<int>(LinearAllocs::malloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Ad = LinearAllocGuard<int>(LinearAllocs::hipMalloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Bh = LinearAllocGuard<int>(LinearAllocs::malloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Bd = LinearAllocGuard<int>(LinearAllocs::hipMalloc,
BLOCKSIZE * sizeof(int));
// Fill input data
std::fill_n(Ah.host_ptr(), BLOCKSIZE, VALUE1);
std::fill_n(Bh.host_ptr(), BLOCKSIZE, VALUE2);
// Stream create
StreamsGuard stream0(1);
// Capture streams into graph
SECTION("Stream Creation Before Capture") {
StreamsGuard stream1(1);
HIP_CHECK(hipStreamBeginCapture(stream0[0], flag));
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipStreamSynchronize(stream1[0]));
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream0[0]>>>(Ad.ptr(), Bd.ptr());
HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture
}
SECTION("Synchronizing multiple streams during Capture") {
StreamsGuard stream1(1), stream2(1);
HIP_CHECK(hipStreamBeginCapture(stream0[0], flag));
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream2[0]));
HIP_CHECK(hipStreamSynchronize(stream1[0]));
HIP_CHECK(hipStreamSynchronize(stream2[0]));
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream0[0]>>>(Ad.ptr(), Bd.ptr());
HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture
}
SECTION("Stream Creation After Capture") {
HIP_CHECK(hipStreamBeginCapture(stream0[0], flag));
StreamsGuard stream1(1);
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipStreamSynchronize(stream1[0]));
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream0[0]>>>(Ad.ptr(), Bd.ptr());
HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture
}
SECTION("Stream Synchronize Before Capture") {
StreamsGuard stream1(1);
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipStreamSynchronize(stream1[0]));
HIP_CHECK(hipStreamBeginCapture(stream0[0], flag));
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream0[0]>>>(Ad.ptr(), Bd.ptr());
HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture
}
SECTION("Stream Synchronize After Capture") {
HIP_CHECK(hipStreamBeginCapture(stream0[0], flag));
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream0[0]>>>(Ad.ptr(), Bd.ptr());
HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture
StreamsGuard stream1(1);
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipStreamSynchronize(stream1[0]));
}
// Execute and test the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, stream0[0]));
HIP_CHECK(hipStreamSynchronize(stream0[0]));
// Check output
HIP_CHECK(hipMemcpy(Ah.host_ptr(), Ad.ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDeviceToHost));
for (int idx = 0; idx < BLOCKSIZE; idx++) {
REQUIRE(Ah.host_ptr()[idx] == (VALUE1 + VALUE2));
}
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Test Description
* ------------------------
* - Test to verify hipStreamSynchronize on a stream behavior when stream capture
* on another stream is ongoing in another thread.
* Test source
* ------------------------
* - catch\unit\graph\hipStreamBeginCapture.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
// Local function executed as thread
static void strmSyncThread(int *Ah, int *Ad, int *Bh, int *Bd,
int BLOCKSIZE, hipError_t *error) {
StreamsGuard stream(1);
HIP_CHECK(hipMemcpyAsync(Ad, Ah, BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream[0]));
HIP_CHECK(hipMemcpyAsync(Bd, Bh, BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream[0]));
*error = hipStreamSynchronize(stream[0]);
}
// Local function executed as thread
static void captureStrmThread(hipGraph_t *graph, int *Ah, int *Ad,
int *Bh, int *Bd, int BLOCKSIZE, int GRIDSIZE,
hipStreamCaptureMode flag, hipError_t *error) {
StreamsGuard stream(1);
// Capture streams into graph
HIP_CHECK(hipStreamBeginCapture(stream[0], flag));
std::thread t1(strmSyncThread, Ah, Ad, Bh, Bd, BLOCKSIZE, error);
t1.join();
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream[0]>>>(Ad, Bd);
HIP_CHECK(hipStreamEndCapture(stream[0], graph)); // End Capture
}
TEST_CASE("Unit_hipStreamBeginCapture_StreamSync_OngoingCapture_MThread") {
constexpr int GRIDSIZE = 1;
constexpr int BLOCKSIZE = 512;
constexpr int VALUE1 = 7, VALUE2 = 11;
hipGraph_t graph{nullptr};
// Allocate device memory
LinearAllocGuard<int> Ah = LinearAllocGuard<int>(LinearAllocs::malloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Ad = LinearAllocGuard<int>(LinearAllocs::hipMalloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Bh = LinearAllocGuard<int>(LinearAllocs::malloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Bd = LinearAllocGuard<int>(LinearAllocs::hipMalloc,
BLOCKSIZE * sizeof(int));
// Fill input data
std::fill_n(Ah.host_ptr(), BLOCKSIZE, VALUE1);
std::fill_n(Bh.host_ptr(), BLOCKSIZE, VALUE2);
// Stream create
hipError_t error = hipSuccess;
SECTION("Capture Flag = hipStreamCaptureModeGlobal Single Threaded") {
StreamsGuard stream(2);
// Capture streams into graph
HIP_CHECK(hipStreamBeginCapture(stream[0], hipStreamCaptureModeGlobal));
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(),
BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(),
BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1]));
error = hipStreamSynchronize(stream[1]);
REQUIRE(error == hipErrorStreamCaptureUnsupported);
}
#if HT_NVIDIA
SECTION("Capture Flag = hipStreamCaptureModeThreadLocal Single Threaded") {
StreamsGuard stream(2);
// Capture streams into graph
HIP_CHECK(hipStreamBeginCapture(stream[0],
hipStreamCaptureModeThreadLocal));
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(),
BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(),
BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1]));
error = hipStreamSynchronize(stream[1]);
REQUIRE(error == hipErrorStreamCaptureUnsupported);
}
#endif
#if HT_AMD
SECTION("Capture Flag = hipStreamCaptureModeGlobal Multithreaded") {
captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(),
Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE,
hipStreamCaptureModeGlobal, &error);
REQUIRE(error == hipErrorStreamCaptureUnsupported);
}
#endif
SECTION("Capture Flag = hipStreamCaptureModeThreadLocal Multithreaded") {
captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(),
Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE,
hipStreamCaptureModeThreadLocal, &error);
REQUIRE(error == hipSuccess);
}
SECTION("Capture Flag = hipStreamCaptureModeRelaxed Multithreaded") {
captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(),
Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE,
hipStreamCaptureModeRelaxed, &error);
REQUIRE(error == hipSuccess);
}
if (graph != nullptr) {
hipGraphExec_t graphExec{nullptr};
StreamsGuard stream(1);
// Execute and test the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, stream[0]));
HIP_CHECK(hipStreamSynchronize(stream[0]));
// Check output
HIP_CHECK(hipMemcpy(Ah.host_ptr(), Ad.ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDeviceToHost));
for (int idx = 0; idx < BLOCKSIZE; idx++) {
REQUIRE(Ah.host_ptr()[idx] == (VALUE1 + VALUE2));
}
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
}
}
@@ -0,0 +1,215 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <math.h>
#include "hip/hip_ext.h"
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#include <sys/stat.h>
#if !defined(S_IFREG) && defined(_S_IFREG)
#define S_IFREG _S_IFREG
#endif
struct gridblockDim {
unsigned int gridX;
unsigned int gridY;
unsigned int gridZ;
unsigned int blockX;
unsigned int blockY;
unsigned int blockZ;
};
class GraphModuleLaunchKernel {
int N = 64;
int SIZE = N*N;
int *A, *B, *C;
hipDeviceptr_t *Ad, *Bd;
hipStream_t stream1, stream2;
hipModule_t module;
hipFunction_t multKernel;
struct {
void* _Ad;
void* _Bd;
void* _Cd;
int _n;
} args1, args2;
size_t size1, size2;
static constexpr char matmulK[] = "matmulK";
public :
GraphModuleLaunchKernel() {
allocateMemory();
moduleLoad();
}
~GraphModuleLaunchKernel() {
deAllocateMemory();
}
void allocateMemory();
void deAllocateMemory();
void moduleLoad();
bool extModuleKernelExecutionMatmul();
bool extModuleKernelExecutionMatmulwithStreamCapture(bool LaunchByDifferentStream = false);
static constexpr char fileName[] = "hipMatMul.code";
};
void GraphModuleLaunchKernel::allocateMemory() {
A = new int[N*N*sizeof(int)];
B = new int[N*N*sizeof(int)];
for (int i=0; i < N; i++) {
for (int j=0; j < N; j++) {
A[i*N +j] = 1;
B[i*N +j] = 1;
}
}
HIPCHECK(hipStreamCreate(&stream1));
HIPCHECK(hipStreamCreate(&stream2));
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Ad),
SIZE*sizeof(int)));
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Bd),
SIZE*sizeof(int)));
HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(&C), SIZE*sizeof(int)));
HIPCHECK(hipMemcpy(Ad, A, SIZE*sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(Bd, B, SIZE*sizeof(int), hipMemcpyHostToDevice));
args1._Ad = Ad;
args1._Bd = Bd;
args1._Cd = C;
args1._n = N;
args2._Ad = NULL;
args2._Bd = NULL;
args2._Cd = NULL;
args2._n = 0;
size1 = sizeof(args1);
size2 = sizeof(args2);
}
void GraphModuleLaunchKernel::moduleLoad() {
HIPCHECK(hipModuleLoad(&module, fileName));
HIPCHECK(hipModuleGetFunction(&multKernel, module, matmulK));
}
void GraphModuleLaunchKernel::deAllocateMemory() {
HIPCHECK(hipStreamDestroy(stream1));
HIPCHECK(hipStreamDestroy(stream2));
delete[] A;
delete[] B;
HIPCHECK(hipFree(Ad));
HIPCHECK(hipFree(Bd));
HIPCHECK(hipHostFree(C));
HIPCHECK(hipModuleUnload(module));
}
bool GraphModuleLaunchKernel::extModuleKernelExecutionMatmul() {
bool testStatus = true;
int mismatch = 0;
void* config1[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args1,
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size1,
HIP_LAUNCH_PARAM_END};
HIPCHECK(hipExtModuleLaunchKernel(multKernel, N, N, 1, 32, 32 , 1, 0,
stream1, NULL,
reinterpret_cast<void**>(&config1),
NULL, NULL, 0));
HIPCHECK(hipStreamSynchronize(stream1));
for (int i = 0; i < N; i++) {
for (int j = 0; j < N; j++) {
if (C[i*N + j] != N)
mismatch++;
}
}
if (mismatch) {
printf("Test failed: the result of matrix multiplications incorrect.\n");
testStatus = false;
}
return testStatus;
}
bool GraphModuleLaunchKernel::extModuleKernelExecutionMatmulwithStreamCapture(bool LaunchByDifferentStream) {
bool testStatus = true;
int mismatch = 0;
hipGraph_t graph{nullptr};
hipGraphExec_t graphExec{nullptr};
HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
void* config1[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args1,
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size1,
HIP_LAUNCH_PARAM_END};
HIPCHECK(hipExtModuleLaunchKernel(multKernel, N, N, 1, 32, 32 , 1, 0,
stream1, NULL,
reinterpret_cast<void**>(&config1),
NULL, NULL, 0));
HIP_CHECK(hipStreamEndCapture(stream1, &graph));
// Validate end capture is successful
REQUIRE(graph != nullptr);
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
REQUIRE(graphExec != nullptr);
// Replay the recorded sequence
HIP_CHECK(hipGraphLaunch(graphExec, LaunchByDifferentStream ? stream2 : stream1));
HIP_CHECK(hipStreamSynchronize(LaunchByDifferentStream ? stream2 : stream1));
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
for (int i = 0; i < N; i++) {
for (int j = 0; j < N; j++) {
if (C[i*N + j] != N)
mismatch++;
}
}
if (mismatch) {
printf("Test failed: the result of matrix multiplications incorrect.\n");
testStatus = false;
}
return testStatus;
}
TEST_CASE("Unit_hipStreamCapture_ExtModuleLaunchKernel") {
struct stat fileStat;
if (stat(GraphModuleLaunchKernel::fileName, &fileStat)
|| !(fileStat.st_mode & S_IFREG)) {
FAIL("module file " << GraphModuleLaunchKernel::fileName
<< " doesn't exist! aborted! \n"
<< "To generate the file, type\n"
<< "/opt/rocm/hip/bin/hipcc --genco hipMatMul.cc -o hipMatMul.code");
return;
}
HIPCHECK(hipSetDevice(0));
GraphModuleLaunchKernel kernelLaunch;
SECTION("extModuleKernelExecutionMatmul") {
REQUIRE(kernelLaunch.extModuleKernelExecutionMatmul());
}
SECTION("extModuleKernelExecutionMatmul_withStreamCapture") {
REQUIRE(kernelLaunch.extModuleKernelExecutionMatmulwithStreamCapture());
}
SECTION("extModuleKernelExecutionMatmul_withStreamCapture_launchByDifferentStream") {
REQUIRE(kernelLaunch.extModuleKernelExecutionMatmulwithStreamCapture(true));
}
}
@@ -6,8 +6,10 @@ 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
@@ -18,9 +20,8 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#include "user_object_common.hh"
/**
* Functional Test for API - hipUserObjectCreate
@@ -34,76 +35,33 @@ THE SOFTWARE.
hipUserObjectRelease with count as X+Y.
*/
struct BoxStruct {
int count;
BoxStruct() {
INFO("Constructor called for Struct!\n");
}
~BoxStruct() {
INFO("Destructor called for Struct!\n");
}
};
class BoxClass {
public:
BoxClass() {
INFO("Constructor called for Class!\n");
}
~BoxClass() {
INFO("Destructor called for Class!\n");
}
};
static void destroyStructObj(void *ptr) {
BoxStruct *ptr1 = reinterpret_cast<BoxStruct *>(ptr);
delete ptr1;
}
static void destroyClassObj(void *ptr) {
BoxClass *ptr2 = reinterpret_cast<BoxClass *>(ptr);
delete ptr2;
}
static void destroyIntObj(void *ptr) {
int *ptr2 = reinterpret_cast<int *>(ptr);
delete ptr2;
}
static void destroyFloatObj(void *ptr) {
float *ptr2 = reinterpret_cast<float *>(ptr);
delete ptr2;
}
/* 1) Call hipUserObjectCreate once and release it by
calling hipUserObjectRelease */
static void hipUserObjectCreate_Functional_1(void *object,
void destroyObj(void *)) {
static void hipUserObjectCreate_Functional_1(void* object, void destroyObj(void*)) {
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyObj, 1,
hipUserObjectNoDestructorSync));
HIP_CHECK(hipUserObjectCreate(&hObject, object, destroyObj, 1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
HIP_CHECK(hipUserObjectRelease(hObject));
}
TEST_CASE("Unit_hipUserObjectCreate_Functional_1") {
SECTION("Called with int Object") {
int *object = new int();
int* object = new int();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_1(object, destroyIntObj);
}
SECTION("Called with float Object") {
float *object = new float();
float* object = new float();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_1(object, destroyFloatObj);
}
SECTION("Called with Class Object") {
BoxClass *object = new BoxClass();
BoxClass* object = new BoxClass();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_1(object, destroyClassObj);
}
SECTION("Called with Struct Object") {
BoxStruct *object = new BoxStruct();
BoxStruct* object = new BoxStruct();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_1(object, destroyStructObj);
}
@@ -111,35 +69,33 @@ TEST_CASE("Unit_hipUserObjectCreate_Functional_1") {
/* 2) Call hipUserObjectCreate refCount as X and release it by
calling hipUserObjectRelease with same refCount. */
static void hipUserObjectCreate_Functional_2(void *object,
void destroyObj(void *)) {
static void hipUserObjectCreate_Functional_2(void* object, void destroyObj(void*)) {
int refCount = 5;
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyObj,
refCount, hipUserObjectNoDestructorSync));
HIP_CHECK(
hipUserObjectCreate(&hObject, object, destroyObj, refCount, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
HIP_CHECK(hipUserObjectRelease(hObject, refCount));
}
TEST_CASE("Unit_hipUserObjectCreate_Functional_2") {
SECTION("Called with int Object") {
int *object = new int();
int* object = new int();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_2(object, destroyIntObj);
}
SECTION("Called with float Object") {
float *object = new float();
float* object = new float();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_2(object, destroyFloatObj);
}
SECTION("Called with Class Object") {
BoxClass *object = new BoxClass();
BoxClass* object = new BoxClass();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_2(object, destroyClassObj);
}
SECTION("Called with Struct Object") {
BoxStruct *object = new BoxStruct();
BoxStruct* object = new BoxStruct();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_2(object, destroyStructObj);
}
@@ -147,12 +103,9 @@ TEST_CASE("Unit_hipUserObjectCreate_Functional_2") {
/* 3) Call hipUserObjectCreate, retain it by calling hipUserObjectRetain
and release it by calling hipUserObjectRelease twice. */
static void hipUserObjectCreate_Functional_3(void *object,
void destroyObj(void *)) {
static void hipUserObjectCreate_Functional_3(void* object, void destroyObj(void*)) {
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyObj,
1, hipUserObjectNoDestructorSync));
HIP_CHECK(hipUserObjectCreate(&hObject, object, destroyObj, 1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
HIP_CHECK(hipUserObjectRetain(hObject));
HIP_CHECK(hipUserObjectRelease(hObject));
@@ -161,22 +114,22 @@ static void hipUserObjectCreate_Functional_3(void *object,
TEST_CASE("Unit_hipUserObjectCreate_Functional_3") {
SECTION("Called with int Object") {
int *object = new int();
int* object = new int();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_3(object, destroyIntObj);
}
SECTION("Called with float Object") {
float *object = new float();
float* object = new float();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_3(object, destroyFloatObj);
}
SECTION("Called with Class Object") {
BoxClass *object = new BoxClass();
BoxClass* object = new BoxClass();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_3(object, destroyClassObj);
}
SECTION("Called with Struct Object") {
BoxStruct *object = new BoxStruct();
BoxStruct* object = new BoxStruct();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_3(object, destroyStructObj);
}
@@ -185,43 +138,40 @@ TEST_CASE("Unit_hipUserObjectCreate_Functional_3") {
/* 4) Call hipUserObjectCreate with refCount as X, retain it by calling
hipUserObjectRetain with count as Y and release it by calling
hipUserObjectRelease with count as X+Y. */
static void hipUserObjectCreate_Functional_4(void *object,
void destroyObj(void *)) {
static void hipUserObjectCreate_Functional_4(void* object, void destroyObj(void*)) {
int refCount = 5;
int refCountRetain = 8;
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyObj,
refCount, hipUserObjectNoDestructorSync));
HIP_CHECK(
hipUserObjectCreate(&hObject, object, destroyObj, refCount, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
HIP_CHECK(hipUserObjectRetain(hObject, refCountRetain));
HIP_CHECK(hipUserObjectRelease(hObject, refCount+refCountRetain));
HIP_CHECK(hipUserObjectRelease(hObject, refCount + refCountRetain));
}
TEST_CASE("Unit_hipUserObjectCreate_Functional_4") {
SECTION("Called with int Object") {
int *object = new int();
int* object = new int();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_4(object, destroyIntObj);
}
SECTION("Called with float Object") {
float *object = new float();
float* object = new float();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_4(object, destroyFloatObj);
}
SECTION("Called with Class Object") {
BoxClass *object = new BoxClass();
BoxClass* object = new BoxClass();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_4(object, destroyClassObj);
}
SECTION("Called with Struct Object") {
BoxStruct *object = new BoxStruct();
BoxStruct* object = new BoxStruct();
REQUIRE(object != nullptr);
hipUserObjectCreate_Functional_4(object, destroyStructObj);
}
}
/**
* Negative Test for API - hipUserObjectCreate
1) Pass User Object as nullptr
@@ -231,411 +181,53 @@ TEST_CASE("Unit_hipUserObjectCreate_Functional_4") {
5) Pass initialRefcount as INT_MAX
6) Pass flag other than hipUserObjectNoDestructorSync
*/
TEST_CASE("Unit_hipUserObjectCreate_Negative") {
hipError_t ret;
int *object = new int();
int* object = new int();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
SECTION("Pass User Object as nullptr") {
ret = hipUserObjectCreate(nullptr, object, destroyIntObj,
1, hipUserObjectNoDestructorSync);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(
hipUserObjectCreate(nullptr, object, destroyIntObj, 1, hipUserObjectNoDestructorSync),
hipErrorInvalidValue);
}
SECTION("Pass object as nullptr") {
ret = hipUserObjectCreate(&hObject, nullptr, destroyIntObj,
1, hipUserObjectNoDestructorSync);
REQUIRE(hipSuccess == ret);
HIP_CHECK(
hipUserObjectCreate(&hObject, nullptr, destroyIntObj, 1, hipUserObjectNoDestructorSync));
}
SECTION("Pass Callback function as nullptr") {
ret = hipUserObjectCreate(&hObject, object, nullptr,
1, hipUserObjectNoDestructorSync);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(
hipUserObjectCreate(&hObject, object, nullptr, 1, hipUserObjectNoDestructorSync),
hipErrorInvalidValue);
}
SECTION("Pass initialRefcount as 0") {
ret = hipUserObjectCreate(&hObject, object, destroyIntObj,
0, hipUserObjectNoDestructorSync);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(
hipUserObjectCreate(&hObject, object, destroyIntObj, 0, hipUserObjectNoDestructorSync),
hipErrorInvalidValue);
}
SECTION("Pass initialRefcount as INT_MAX") {
ret = hipUserObjectCreate(&hObject, object, destroyIntObj,
INT_MAX, hipUserObjectNoDestructorSync);
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipUserObjectCreate(&hObject, object, destroyIntObj, INT_MAX,
hipUserObjectNoDestructorSync));
}
SECTION("Pass flag other than hipUserObjectNoDestructorSync") {
ret = hipUserObjectCreate(&hObject, object, destroyIntObj,
1, hipUserObjectFlags(9));
REQUIRE(hipErrorInvalidValue == ret);
}
}
/**
* Negative Test for API - hipUserObjectRelease
1) Pass User Object as nullptr
2) Pass initialRefcount as 0
3) Pass initialRefcount as INT_MAX
*/
TEST_CASE("Unit_hipUserObjectRelease_Negative") {
hipError_t ret;
int *object = new int();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyIntObj,
1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
SECTION("Pass User Object as nullptr") {
ret = hipUserObjectRelease(nullptr, 1);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass initialRefcount as 0") {
ret = hipUserObjectRelease(hObject, 0);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass initialRefcount as INT_MAX") {
ret = hipUserObjectRelease(hObject, INT_MAX);
REQUIRE(hipSuccess == ret);
}
}
/**
* Negative Test for API - hipUserObjectRetain
1) Pass User Object as nullptr
2) Pass initialRefcount as 0
3) Pass initialRefcount as INT_MAX
*/
TEST_CASE("Unit_hipUserObjectRetain_Negative") {
hipError_t ret;
int *object = new int();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyIntObj,
1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
SECTION("Pass User Object as nullptr") {
ret = hipUserObjectRetain(nullptr, 1);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass initialRefcount as 0") {
ret = hipUserObjectRetain(hObject, 0);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass initialRefcount as INT_MAX") {
ret = hipUserObjectRetain(hObject, INT_MAX);
REQUIRE(hipSuccess == ret);
HIP_CHECK_ERROR(hipUserObjectCreate(&hObject, object, destroyIntObj, 1, hipUserObjectFlags(9)),
hipErrorInvalidValue);
}
}
TEST_CASE("Unit_hipUserObj_Negative_Test") {
hipError_t ret;
int *object = new int();
int* object = new int();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
// Create a new hObject with 2 reference
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyIntObj,
2, hipUserObjectNoDestructorSync));
HIP_CHECK(hipUserObjectCreate(&hObject, object, destroyIntObj, 2, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
// Release more than created.
ret = hipUserObjectRelease(hObject, 4);
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipUserObjectRelease(hObject, 4));
// Retain reference to a removed user object
ret = hipUserObjectRetain(hObject, 1);
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipUserObjectRetain(hObject, 1));
}
/**
* Functional Test for API - hipGraphRetainUserObject
*/
/* 1) Create GraphUserObject and retain it by calling hipGraphRetainUserObject
and release it by calling hipGraphReleaseUserObject. */
static void hipGraphRetainUserObject_Functional_1(void *object,
void destroyObj(void *)) {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyObj,
1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 1,
hipGraphUserObjectMove));
HIP_CHECK(hipGraphReleaseUserObject(graph, hObject));
HIP_CHECK(hipUserObjectRelease(hObject));
HIP_CHECK(hipGraphDestroy(graph));
}
TEST_CASE("Unit_hipGraphRetainUserObject_Functional_1") {
SECTION("Called with int Object") {
int *object = new int();
REQUIRE(object != nullptr);
hipGraphRetainUserObject_Functional_1(object, destroyIntObj);
}
SECTION("Called with float Object") {
float *object = new float();
REQUIRE(object != nullptr);
hipGraphRetainUserObject_Functional_1(object, destroyFloatObj);
}
SECTION("Called with Class Object") {
BoxClass *object = new BoxClass();
REQUIRE(object != nullptr);
hipGraphRetainUserObject_Functional_1(object, destroyClassObj);
}
SECTION("Called with Struct Object") {
BoxStruct *object = new BoxStruct();
REQUIRE(object != nullptr);
hipGraphRetainUserObject_Functional_1(object, destroyStructObj);
}
}
/* 2) Create UserObject and GraphUserObject and retain using custom reference
count and release it by calling hipGraphReleaseUserObject with count. */
TEST_CASE("Unit_hipGraphRetainUserObject_Functional_2") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipGraphNode_t memcpyNode, kNode;
hipKernelNodeParams kNodeParams{};
hipStream_t streamForGraph;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
std::vector<hipGraphNode_t> dependencies;
hipGraphExec_t graphExec;
size_t NElem{N};
HIP_CHECK(hipStreamCreate(&streamForGraph));
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(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, Nbytes,
hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h, Nbytes,
hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void*>(&NElem)};
kNodeParams.func = reinterpret_cast<void*>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(
hipGraphAddKernelNode(&kNode, graph, dependencies.data(), dependencies.size(), &kNodeParams));
dependencies.clear();
dependencies.push_back(kNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), dependencies.size(),
C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
float* object = new float();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(
hipUserObjectCreate(&hObject, object, destroyFloatObj, 1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 1,
hipGraphUserObjectMove)); // Pass ownership to hipGraph
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verify result
HipTest::checkVectorADD<int>(A_h, B_h, C_h, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
/**
* Negative Test for API - hipGraphRetainUserObject
1) Pass graph as nullptr
2) Pass User Object as nullptr
3) Pass initialRefcount as 0
4) Pass initialRefcount as INT_MAX
5) Pass flag as 0
6) Pass flag as INT_MAX
*/
TEST_CASE("Unit_hipGraphRetainUserObject_Negative") {
hipError_t ret;
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
float *object = new float();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyFloatObj,
1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
SECTION("Pass graph as nullptr") {
ret = hipGraphRetainUserObject(nullptr, hObject, 1,
hipGraphUserObjectMove);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass User Object as nullptr") {
ret = hipGraphRetainUserObject(graph, nullptr, 1,
hipGraphUserObjectMove);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass initialRefcount as 0") {
ret = hipGraphRetainUserObject(graph, hObject, 0,
hipGraphUserObjectMove);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass initialRefcount as INT_MAX") {
ret = hipGraphRetainUserObject(graph, hObject, INT_MAX,
hipGraphUserObjectMove);
REQUIRE(hipSuccess == ret);
}
SECTION("Pass flag as 0") {
ret = hipGraphRetainUserObject(graph, hObject, 1, 0);
REQUIRE(hipSuccess == ret);
}
SECTION("Pass flag as INT_MAX") {
ret = hipGraphRetainUserObject(graph, hObject, 1, INT_MAX);
REQUIRE(hipErrorInvalidValue == ret);
}
HIP_CHECK(hipUserObjectRelease(hObject, 1));
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Negative Test for API - hipGraphReleaseUserObject
1) Pass graph as nullptr
2) Pass User Object as nullptr
3) Pass initialRefcount as 0
4) Pass initialRefcount as INT_MAX
*/
TEST_CASE("Unit_hipGraphReleaseUserObject_Negative") {
hipError_t ret;
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
float *object = new float();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyFloatObj,
1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 1,
hipGraphUserObjectMove));
SECTION("Pass graph as nullptr") {
ret = hipGraphReleaseUserObject(nullptr, hObject, 1);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass User Object as nullptr") {
ret = hipGraphReleaseUserObject(graph, nullptr, 1);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass initialRefcount as 0") {
ret = hipGraphReleaseUserObject(graph, hObject, 0);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass initialRefcount as INT_MAX") {
ret = hipGraphReleaseUserObject(graph, hObject, INT_MAX);
REQUIRE(hipSuccess == ret);
}
HIP_CHECK(hipUserObjectRelease(hObject, 1));
HIP_CHECK(hipGraphDestroy(graph));
}
TEST_CASE("Unit_hipGraphRetainUserObject_Negative_Basic") {
hipError_t ret;
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
float *object = new float();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyFloatObj,
1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
// Retain graph object with reference count 2
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 2,
hipGraphUserObjectMove));
// Release graph object with reference count more than 2
ret = hipGraphReleaseUserObject(graph, hObject, 4);
REQUIRE(hipSuccess == ret);
// Again Retain graph object with reference count 8
ret = hipGraphRetainUserObject(graph, hObject, 8,
hipGraphUserObjectMove);
REQUIRE(hipSuccess == ret);
// Release graph object with reference count 1
ret = hipGraphReleaseUserObject(graph, hObject, 1);
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipUserObjectRelease(hObject, 1));
HIP_CHECK(hipGraphDestroy(graph));
}
TEST_CASE("Unit_hipGraphRetainUserObject_Negative_Null_Object") {
hipError_t ret;
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
float *object = nullptr; // this is used for Null_Object test
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object,
destroyFloatObj,
1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
// Retain graph object with reference count 2
HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 2,
hipGraphUserObjectMove));
// Release graph object with reference count more than 2
ret = hipGraphReleaseUserObject(graph, hObject, 4);
REQUIRE(hipSuccess == ret);
// Again Retain graph object with reference count 8
ret = hipGraphRetainUserObject(graph, hObject, 8, 0);
REQUIRE(hipSuccess == ret);
// Release graph object with reference count 1
ret = hipGraphReleaseUserObject(graph, hObject, 1);
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipUserObjectRelease(hObject, 1));
HIP_CHECK(hipGraphDestroy(graph));
}
@@ -0,0 +1,48 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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.
*/
#include <hip_test_common.hh>
#include "user_object_common.hh"
/**
* Negative Test for API - hipUserObjectRelease
1) Pass User Object as nullptr
2) Pass initialRefcount as 0
3) Pass initialRefcount as INT_MAX
*/
TEST_CASE("Unit_hipUserObjectRelease_Negative") {
int* object = new int();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object, destroyIntObj, 1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
SECTION("Pass User Object as nullptr") {
HIP_CHECK_ERROR(hipUserObjectRelease(nullptr, 1), hipErrorInvalidValue);
}
SECTION("Pass initialRefcount as 0") {
HIP_CHECK_ERROR(hipUserObjectRelease(hObject, 0), hipErrorInvalidValue);
}
SECTION("Pass initialRefcount as INT_MAX") { HIP_CHECK(hipUserObjectRelease(hObject, INT_MAX)); }
}
@@ -0,0 +1,48 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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.
*/
#include <hip_test_common.hh>
#include "user_object_common.hh"
/**
* Negative Test for API - hipUserObjectRetain
1) Pass User Object as nullptr
2) Pass initialRefcount as 0
3) Pass initialRefcount as INT_MAX
*/
TEST_CASE("Unit_hipUserObjectRetain_Negative") {
int* object = new int();
REQUIRE(object != nullptr);
hipUserObject_t hObject;
HIP_CHECK(hipUserObjectCreate(&hObject, object, destroyIntObj, 1, hipUserObjectNoDestructorSync));
REQUIRE(hObject != nullptr);
SECTION("Pass User Object as nullptr") {
HIP_CHECK_ERROR(hipUserObjectRetain(nullptr, 1), hipErrorInvalidValue);
}
SECTION("Pass initialRefcount as 0") {
HIP_CHECK_ERROR(hipUserObjectRetain(hObject, 0), hipErrorInvalidValue);
}
SECTION("Pass initialRefcount as INT_MAX") { HIP_CHECK(hipUserObjectRetain(hObject, INT_MAX)); }
}
@@ -0,0 +1,61 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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
#include <hip_test_common.hh>
struct BoxStruct {
int count;
BoxStruct() { INFO("Constructor called for Struct!\n"); }
~BoxStruct() { INFO("Destructor called for Struct!\n"); }
};
class BoxClass {
public:
BoxClass() { INFO("Constructor called for Class!\n"); }
~BoxClass() { INFO("Destructor called for Class!\n"); }
};
namespace {
void destroyStructObj(void* ptr) {
BoxStruct* ptr1 = reinterpret_cast<BoxStruct*>(ptr);
delete ptr1;
}
void destroyClassObj(void* ptr) {
BoxClass* ptr2 = reinterpret_cast<BoxClass*>(ptr);
delete ptr2;
}
void destroyIntObj(void* ptr) {
int* ptr2 = reinterpret_cast<int*>(ptr);
delete ptr2;
}
void destroyFloatObj(void* ptr) {
float* ptr2 = reinterpret_cast<float*>(ptr);
delete ptr2;
}
} // anonymous namespace
@@ -140,6 +140,14 @@ static void DisplayHmmFlgs(int *Signal) {
}
TEST_CASE("Unit_HMM_OverSubscriptionTst") {
hipDeviceProp_t prop;
HIP_CHECK(hipGetDeviceProperties(&prop, 0));
char *p = nullptr;
p = strstr(prop.gcnArchName, "xnack+");
if (p == nullptr) {
INFO("Skipped due current device is non xnack device.");
return;
}
int HmmEnabled = 0;
// The following Shared Mem is to get Max GPU Mem
// The size requested is for three ints
@@ -55,8 +55,6 @@ static void Malloc3DArray_DiffSizes(int gpu) {
size_t width{size}, height{size}, depth{size};
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
std::array<hipArray_t, ARRAY_LOOP> arr;
size_t pavail, avail;
HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
for (int i = 0; i < ARRAY_LOOP; i++) {
HIP_CHECK_THREAD(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, height, depth),
@@ -65,9 +63,6 @@ static void Malloc3DArray_DiffSizes(int gpu) {
for (int i = 0; i < ARRAY_LOOP; i++) {
HIP_CHECK_THREAD(hipFreeArray(arr[i]));
}
HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr));
REQUIRE_THREAD(pavail == avail);
}
}
@@ -85,7 +80,6 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") {
std::vector<std::thread> threadlist;
int devCnt = 0;
devCnt = HipTest::getDeviceCount();
const auto pavail = getFreeMem();
for (int i = 0; i < devCnt; i++) {
threadlist.push_back(std::thread(Malloc3DArray_DiffSizes, i));
}
@@ -94,12 +88,6 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") {
t.join();
}
HIP_CHECK_THREAD_FINALIZE();
const auto avail = getFreeMem();
if (pavail != avail) {
WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
REQUIRE(false);
}
}
namespace {
@@ -57,9 +57,6 @@ static void MallocArray_DiffSizes(int gpu) {
for (const auto& size : runs) {
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
std::array<hipArray_t, ARRAY_LOOP> A_d;
size_t pavail, avail;
HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
for (int i = 0; i < ARRAY_LOOP; i++) {
HIP_CHECK_THREAD(
hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault));
@@ -67,9 +64,6 @@ static void MallocArray_DiffSizes(int gpu) {
for (int i = 0; i < ARRAY_LOOP; i++) {
HIP_CHECK_THREAD(hipFreeArray(A_d[i]));
}
HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr));
REQUIRE_THREAD(pavail == avail);
}
}
@@ -87,7 +81,6 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") {
std::vector<std::thread> threadlist;
int devCnt = 0;
devCnt = HipTest::getDeviceCount();
const auto pavail = getFreeMem();
for (int i = 0; i < devCnt; i++) {
threadlist.push_back(std::thread(MallocArray_DiffSizes, i));
}
@@ -96,12 +89,6 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") {
t.join();
}
HIP_CHECK_THREAD_FINALIZE();
const auto avail = getFreeMem();
if (pavail != avail) {
WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
REQUIRE(false);
}
}
// Kernels ///////////////////////////////////////
@@ -85,7 +85,6 @@ static std::atomic<bool> g_thTestPassed{true};
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
@@ -93,16 +92,7 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
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);
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;
}
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
@@ -121,16 +111,7 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
TestPassed = false;
}
HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
if (!concurOnOneGPU && (curAvl < prevAvl || 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;
}
@@ -140,8 +121,7 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
* Regress memory allocation and free in loop
*/
static bool regressAllocInLoop(int gpu) {
bool TestPassed = true;
size_t tot, avail, ptot, pavail, numBytes;
size_t numBytes;
int i = 0;
int* ptr;
@@ -150,41 +130,18 @@ static bool regressAllocInLoop(int gpu) {
// 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;
return true;
}
/**
@@ -194,23 +151,13 @@ static bool regressAllocInLoop(int gpu) {
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);
HIPCHECK(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;
}
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
@@ -229,24 +176,7 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) {
TestPassed = false;
}
HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot));
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIPCHECK(hipMemGetInfo(&curAvl, &curTot));
if (!concurOnOneGPU && (curAvl < prevAvl || 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;
}
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;
}
@@ -256,8 +186,7 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) {
* In Multithreaded Environment
*/
static bool regressAllocInLoopMthread(int gpu) {
bool TestPassed = true;
size_t tot, avail, ptot, pavail, numBytes;
size_t numBytes;
int i = 0;
int* ptr;
@@ -266,41 +195,18 @@ static bool regressAllocInLoopMthread(int gpu) {
// 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;
return true;
}
/*
@@ -359,7 +265,7 @@ TEST_CASE("Unit_hipMalloc_LoopRegressionAllocFreeCycles") {
* of time.
*/
TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
size_t avail{0}, tot{0}, pavail{0}, ptot{0};
size_t avail{0}, tot{0};
bool ret{false};
hipError_t err{};
std::vector<int*> ptrlist{};
@@ -370,8 +276,6 @@ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
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) {
@@ -392,13 +296,8 @@ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
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);
}
@@ -346,8 +346,6 @@ static void MemoryAllocDiffSizes(int gpu) {
} else {
width = LARGECHUNK_NUMW * sizeof(T);
}
size_t tot, avail, ptot, pavail;
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
for (int i = 0; i < CHUNK_LOOP; i++) {
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d[i]),
&pitch_A, width, sizes));
@@ -355,10 +353,6 @@ static void MemoryAllocDiffSizes(int gpu) {
for (int i = 0; i < CHUNK_LOOP; i++) {
HIP_CHECK(hipFree(A_d[i]));
}
HIP_CHECK(hipMemGetInfo(&avail, &tot));
if (pavail != avail) {
HIPASSERT(false);
}
}
}
@@ -487,9 +481,6 @@ TEST_CASE("Unit_hipMallocPitch_MultiThread", "") {
int devCnt = 0;
devCnt = HipTest::getDeviceCount();
size_t tot, avail, ptot, pavail;
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
for (int i = 0; i < devCnt; i++) {
threadlist.push_back(std::thread(threadFunc, i));
}
@@ -497,12 +488,6 @@ TEST_CASE("Unit_hipMallocPitch_MultiThread", "") {
for (auto &t : threadlist) {
t.join();
}
HIP_CHECK(hipMemGetInfo(&avail, &tot));
if (pavail != avail) {
WARN("Memory leak of hipMallocPitch API in multithreaded scenario");
REQUIRE(false);
}
}
/*
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2022 - 2023 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
@@ -211,13 +211,11 @@ TEST_CASE("Unit_hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_NegTst") {
blockSizeToDynamicSMemSize, 0, 0);
REQUIRE(ret == hipErrorInvalidValue);
}
#if HT_NVIDIA
SECTION("invalid flag") {
ret = hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(&minGridSize,
&blockSize, f1, blockSizeToDynamicSMemSize, 0, 0xffff);
REQUIRE(ret == hipErrorInvalidValue);
}
#endif
}
/**
@@ -3,6 +3,7 @@ set(TEST_SRC
saxpy.cc
warpsize.cc
hipRtcFunctional.cc
hipStreamCaptureRtc.cc
)
# AMD only tests
@@ -0,0 +1,136 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip/hiprtc.h>
#include <math.h>
#include <vector>
#include <hip_test_common.hh>
static constexpr auto kernel_src {
R"_KERN_EMBED_(
extern "C" __global__ void kernel_func(float* f)
{
f[0] = 1.0;
}
)_KERN_EMBED_"
};
TEST_CASE("Unit_hipStreamCaptureRtc") {
hipStream_t stream = nullptr;
hipGraph_t graph = nullptr;
hipGraphExec_t graph_exec = nullptr;
float data_h = 0.0;
float* data_d = nullptr;
// Init data
HIPCHECK(hipMalloc(&data_d, sizeof(float)));
HIPCHECK(hipMemcpy(data_d, &data_h, sizeof(float), hipMemcpyHostToDevice));
// Compile kernel
std::vector<char> code;
hiprtcProgram prog;
HIPRTC_CHECK(hiprtcCreateProgram(&prog, kernel_src, "hipStreamCaptureRtc.cu", 0, nullptr, nullptr));
hipDeviceProp_t props;
int device = 0;
HIP_CHECK(hipSetDevice(device));
HIP_CHECK(hipGetDeviceProperties(&props, device));
#ifdef __HIP_PLATFORM_AMD__
std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
#else
std::string sarg = std::string("--fmad=false");
#endif
std::vector<const char*> options = { sarg.c_str() };
auto compileResult = hiprtcCompileProgram(prog, options.size(), options.data());
if (compileResult != HIPRTC_SUCCESS) {
size_t logSize = 0;
hiprtcGetProgramLogSize(prog, &logSize);
if (logSize) {
std::vector<char> log(logSize, '\0');
if (hiprtcGetProgramLog(prog, log.data()) == HIPRTC_SUCCESS) {
FAIL("hiprtcCompileProgram failed with log" << log.data());
return;
}
}
FAIL("hiprtcCompileProgram failed without log");
return;
}
size_t codeSize = 0;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
code.resize(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module = nullptr;
hipFunction_t kernel = nullptr;
#if HT_NVIDIA
HIPCHECK(hipInit(0));
hipCtx_t ctx;
HIPCHECK(hipCtxCreate(&ctx, 0, device));
#endif
HIPCHECK(hipModuleLoadData(&module, code.data()));
HIPCHECK(hipModuleGetFunction(&kernel, module, "kernel_func"));
// Start capture
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
// Launch kernel
auto size = sizeof(float*);
void *config[] = { HIP_LAUNCH_PARAM_BUFFER_POINTER, &data_d,
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END };
HIPCHECK(hipModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, nullptr, config));
HIPCHECK(hipStreamEndCapture(stream, &graph));
size_t numNodes = 0;
HIPCHECK(hipGraphGetNodes(graph, nullptr, &numNodes));
INFO("Num of nodes returned by GetNodes : " << numNodes);
REQUIRE(numNodes == 1);
// Ensure that no actual work has been done for the captured
// stream before graph execution
float tmp = 2.0;
HIPCHECK(hipMemcpy(&tmp, data_d, sizeof(float), hipMemcpyDeviceToHost));
REQUIRE(tmp == 0.0);
HIPCHECK(hipGraphInstantiate(&graph_exec, graph, NULL, NULL, 0));
HIPCHECK(hipGraphDestroy(graph));
HIPCHECK(hipGraphLaunch(graph_exec, stream));
HIPCHECK(hipStreamSynchronize(stream));
HIPCHECK(hipStreamDestroy(stream));
// Check that the work was done
HIPCHECK(hipMemcpy(&tmp, data_d, sizeof(float), hipMemcpyDeviceToHost));
HIPCHECK(hipFree(data_d));
REQUIRE(tmp == 1.0);
#if HT_NVIDIA
HIPCHECK(hipCtxDestroy(ctx));
#endif
}
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2022-2023 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
@@ -18,38 +18,20 @@ THE SOFTWARE.
*/
/**
Testcase Scenarios :
* @addtogroup hipStreamCreateWithPriority hipStreamCreateWithPriority
* @{
* @ingroup StreamTest
* `hipStreamCreateWithPriority (hipStream_t *stream, unsigned int flags, int priority)` -
* begins graph capture on a stream
*/
1)Create streams with default flag for all available priority levels and
queue tasks in each of these streams, perform device synchronize and validate behavior.
2)Create streams with non-blocking flag for all available priority levels and
queue tasks in each of these streams, perform stream synchronize and validate behavior.
3)Create streams with default flag for all available priority levels and
queue tasks in each of these streams, perform stream synchronize and validate behavior.
4)Create streams with non-blocking flag for all available priority levels and
queue tasks in each of these streams, perform device synchronize and validate behavior.
5)Create a stream for each priority level with default flag, Launch memcpy and kernel
tasks on these streams from multiple threads. Validate all the results.
6)Create a stream for each priority level with non-blocking flag, Launch memcpy and
kernel tasks on these streams from multiple threads. Validate all the results.
7) Validate negative scenarios for hipStreamCreateWithPriority api.
8) Validate stream priorities with event after classifying them as low, medium, high.
*/
#include "streamCommon.hh"
#include <hip_test_kernels.hh>
#include <atomic>
#include <vector>
#include "streamCommon.hh" // NOLINT
#define MEMCPYSIZE 64*1024*1024
#define MEMCPYSIZE2 1024*1024
#define MEMCPYSIZE1 (64*1024*1024)
#define MEMCPYSIZE2 (1024*1024)
#define NUMITERS 2
#define GRIDSIZE 1024
#define BLOCKSIZE 256
@@ -87,7 +69,7 @@ void funcTestsForAllPriorityLevelsWrtNullStrm(unsigned int flags,
int priority;
int priority_low{};
int priority_high{};
size_t size = MEMCPYSIZE2*sizeof(int);
size_t size = MEMCPYSIZE2 * sizeof(int);
// Test is to get the Stream Priority Range
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
@@ -194,7 +176,7 @@ void funcTestsForAllPriorityLevelsWrtNullStrm(unsigned int flags,
*/
void queueTasksInStreams(std::vector<hipStream_t> *stream,
const int arrsize) {
size_t size = MEMCPYSIZE2*sizeof(int);
size_t size = MEMCPYSIZE2 * sizeof(int);
// Allocate memory
int **A_d = reinterpret_cast<int**>(malloc(arrsize*sizeof(int *)));
int **C_d = reinterpret_cast<int**>(malloc(arrsize*sizeof(int *)));
@@ -211,8 +193,8 @@ void queueTasksInStreams(std::vector<hipStream_t> *stream,
HIPASSERT(A_h[idx] != nullptr);
C_h[idx] = reinterpret_cast<int*>(malloc(size));
HIPASSERT(C_h[idx] != nullptr);
HIPCHECK(hipMalloc(&A_d[idx], size));
HIPCHECK(hipMalloc(&C_d[idx], size));
HIP_CHECK(hipMalloc(&A_d[idx], size));
HIP_CHECK(hipMalloc(&C_d[idx], size));
}
// Initialize host memory
constexpr int initVal = 2;
@@ -223,20 +205,20 @@ void queueTasksInStreams(std::vector<hipStream_t> *stream,
}
// Launch task on each stream
for (int idx = 0; idx < arrsize; idx++) {
HIPCHECK(hipMemcpyAsync(A_d[idx], A_h[idx], size,
HIP_CHECK(hipMemcpyAsync(A_d[idx], A_h[idx], size,
hipMemcpyHostToDevice, (*stream)[idx]));
hipLaunchKernelGGL((HipTest::vector_square), dim3(GRIDSIZE),
dim3(BLOCKSIZE), 0, (*stream)[idx], A_d[idx],
C_d[idx], MEMCPYSIZE2);
HIP_CHECK(hipGetLastError());
HIPCHECK(hipMemcpyAsync(C_h[idx], C_d[idx], size,
HIP_CHECK(hipMemcpyAsync(C_h[idx], C_d[idx], size,
hipMemcpyDeviceToHost, (*stream)[idx]));
}
bool isPassed = true;
// Validate the output of each queue
for (int idx = 0; idx < arrsize; idx++) {
HIPCHECK(hipStreamSynchronize((*stream)[idx]));
HIP_CHECK(hipStreamSynchronize((*stream)[idx]));
for (int idy = 0; idy < MEMCPYSIZE2; idy++) {
if (C_h[idx][idy] != A_h[idx][idy] * A_h[idx][idy]) {
UNSCOPED_INFO("Data mismatch at idx:" << idx << " idy:" << idy);
@@ -248,8 +230,8 @@ void queueTasksInStreams(std::vector<hipStream_t> *stream,
}
// Deallocate memory
for (int idx = 0; idx < arrsize; idx++) {
HIPCHECK(hipFree(reinterpret_cast<void*>(C_d[idx])));
HIPCHECK(hipFree(reinterpret_cast<void*>(A_d[idx])));
HIP_CHECK(hipFree(reinterpret_cast<void*>(C_d[idx])));
HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d[idx])));
free(C_h[idx]);
free(A_h[idx]);
}
@@ -322,7 +304,7 @@ bool runFuncTestsForAllPriorityLevelsMultThread(unsigned int flags) {
template <typename T>
bool validateStreamPrioritiesWithEvents() {
size_t size = NUMITERS*MEMCPYSIZE;
size_t size = NUMITERS * MEMCPYSIZE1;
// get the range of priorities available
#define OP(x) \
@@ -421,13 +403,13 @@ bool validateStreamPrioritiesWithEvents() {
#undef OP
// launch kernels repeatedly on each of the prioritiy streams
for (int i = 0; i < static_cast<int>(size); i += MEMCPYSIZE) {
for (int i = 0; i < static_cast<int>(size); i += MEMCPYSIZE1) {
int j = i / sizeof(T);
#define OP(x) \
if (enable_priority_##x) { \
hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), \
dim3(BLOCKSIZE), 0, stream_##x, dst_d_##x + j, src_d_##x + j, \
(MEMCPYSIZE / sizeof(T))); \
(MEMCPYSIZE1 / sizeof(T))); \
HIP_CHECK(hipGetLastError()); \
}
OP(low)
@@ -522,19 +504,330 @@ bool validateStreamPrioritiesWithEvents() {
return true;
}
#define LOW_PRIORITY_STREAMCOUNT 2
#define HIGH_PRIORITY_STREAMCOUNT 2
#define NORMAL_PRIORITY_STREAMCOUNT 2
template <typename T>
void TestForMultipleStreamWithPriority(void) {
size_t size = NUMITERS * MEMCPYSIZE1;
// get the range of priorities available
#define OP(x) \
int priority_##x; \
bool enable_priority_##x = false;
OP(low)
OP(normal)
OP(high)
#undef OP
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
INFO("HIP stream priority range - low: " << priority_low << ",high: "
<< priority_high << ",normal: "
<< (priority_low + priority_high)/2);
// Check if priorities are indeed supported
// Enable/disable priorities based on number of available priority levels
enable_priority_low = true;
enable_priority_high = true;
if ((priority_low - priority_high) > 1) {
enable_priority_normal = true;
}
if (enable_priority_normal) {
priority_normal = ((priority_low + priority_high) / 2);
}
// create streams with low priority
hipStream_t stream_low[LOW_PRIORITY_STREAMCOUNT];
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_low) {
HIP_CHECK(hipStreamCreateWithPriority(&stream_low[i],
hipStreamDefault, priority_low));
}
}
// create streams with normal priority
hipStream_t stream_normal[NORMAL_PRIORITY_STREAMCOUNT];
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_normal) {
HIP_CHECK(hipStreamCreateWithPriority(&stream_normal[i],
hipStreamDefault, priority_normal));
}
}
// create streams with high priority
hipStream_t stream_high[HIGH_PRIORITY_STREAMCOUNT];
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_high) {
HIP_CHECK(hipStreamCreateWithPriority(&stream_high[i],
hipStreamDefault, priority_high));
}
}
// allocate and initialise host source and destination buffers for
// low streams
T* src_h_low[LOW_PRIORITY_STREAMCOUNT];
T* dst_h_low[LOW_PRIORITY_STREAMCOUNT];
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_low) {
src_h_low[i] = reinterpret_cast<T*>(malloc(size));
REQUIRE(src_h_low[i] != nullptr);
mem_init<T>(src_h_low[i], (size / sizeof(T)));
dst_h_low[i] = reinterpret_cast<T*>(malloc(size));
REQUIRE(dst_h_low[i] != nullptr);
memset(dst_h_low[i], 0, size);
}
}
// allocate and initialise host source and destination buffers for
// normal streams
T* src_h_normal[NORMAL_PRIORITY_STREAMCOUNT];
T* dst_h_normal[NORMAL_PRIORITY_STREAMCOUNT];
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_normal) {
src_h_normal[i] = reinterpret_cast<T*>(malloc(size));
REQUIRE(src_h_normal[i] != nullptr);
mem_init<T>(src_h_normal[i], (size / sizeof(T)));
dst_h_normal[i] = reinterpret_cast<T*>(malloc(size));
REQUIRE(dst_h_normal[i] != nullptr);
memset(dst_h_normal[i], 0, size);
}
}
// allocate and initialise host source and destination buffers for
// high streams
T* src_h_high[HIGH_PRIORITY_STREAMCOUNT];
T* dst_h_high[HIGH_PRIORITY_STREAMCOUNT];
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_high) {
src_h_high[i] = reinterpret_cast<T*>(malloc(size));
REQUIRE(src_h_high[i] != nullptr);
mem_init<T>(src_h_high[i], (size / sizeof(T)));
dst_h_high[i] = reinterpret_cast<T*>(malloc(size));
REQUIRE(dst_h_high[i] != nullptr);
memset(dst_h_high[i], 0, size);
}
}
// allocate and initialize device source and destination buffers for
// low streams
T* src_d_low[LOW_PRIORITY_STREAMCOUNT];
T* dst_d_low[LOW_PRIORITY_STREAMCOUNT];
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_low) {
HIP_CHECK(hipMalloc(&src_d_low[i], size));
HIP_CHECK(hipMemcpy(src_d_low[i], src_h_low[i], size,
hipMemcpyHostToDevice));
HIP_CHECK(hipMalloc(&dst_d_low[i], size));
}
}
// allocate and initialize device source and destination buffers for
// normal streams
T* src_d_normal[NORMAL_PRIORITY_STREAMCOUNT];
T* dst_d_normal[NORMAL_PRIORITY_STREAMCOUNT];
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_normal) {
HIP_CHECK(hipMalloc(&src_d_normal[i], size));
HIP_CHECK(hipMemcpy(src_d_normal[i], src_h_normal[i], size,
hipMemcpyHostToDevice));
HIP_CHECK(hipMalloc(&dst_d_normal[i], size));
}
}
// allocate and initialize device source and destination buffers for
// high streams
T* src_d_high[HIGH_PRIORITY_STREAMCOUNT];
T* dst_d_high[HIGH_PRIORITY_STREAMCOUNT];
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_high) {
HIP_CHECK(hipMalloc(&src_d_high[i], size));
HIP_CHECK(hipMemcpy(src_d_high[i], src_h_high[i], size,
hipMemcpyHostToDevice));
HIP_CHECK(hipMalloc(&dst_d_high[i], size));
}
}
hipEvent_t event_start_low[LOW_PRIORITY_STREAMCOUNT];
hipEvent_t event_end_low[LOW_PRIORITY_STREAMCOUNT];
// create events for measuring time spent in kernel execution for low
// priority streams
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_low) {
HIP_CHECK(hipEventCreate(&event_start_low[i]));
HIP_CHECK(hipEventCreate(&event_end_low[i]));
}
}
hipEvent_t event_start_normal[NORMAL_PRIORITY_STREAMCOUNT];
hipEvent_t event_end_normal[NORMAL_PRIORITY_STREAMCOUNT];
// create events for measuring time spent in kernel execution for
// normal priority streams
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_normal) {
HIP_CHECK(hipEventCreate(&event_start_normal[i]));
HIP_CHECK(hipEventCreate(&event_end_normal[i]));
}
}
hipEvent_t event_start_high[HIGH_PRIORITY_STREAMCOUNT];
hipEvent_t event_end_high[HIGH_PRIORITY_STREAMCOUNT];
// create events for measuring time spent in kernel execution for
// high priority streams
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_high) {
HIP_CHECK(hipEventCreate(&event_start_high[i]));
HIP_CHECK(hipEventCreate(&event_end_high[i]));
}
}
// record start events for each of the low priority streams
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_low) {
HIP_CHECK(hipEventRecord(event_start_low[i], stream_low[i]));
}
}
// record start events for each of the normal priority streams
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_normal) {
HIP_CHECK(hipEventRecord(event_start_normal[i], stream_normal[i]));
}
}
// record start events for each of the high priority streams
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_high) {
HIP_CHECK(hipEventRecord(event_start_high[i], stream_high[i]));
}
}
// launch kernels repeatedly on each of the low prioritiy stream
for (int k = 0; k < LOW_PRIORITY_STREAMCOUNT; ++k) {
for (int i = 0; i < size; i += MEMCPYSIZE1) {
int j = i / sizeof(T);
if (enable_priority_low) {
hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), dim3(BLOCKSIZE),
0, stream_low[k], dst_d_low[k] + j, src_d_low[k] + j,
(MEMCPYSIZE1 / sizeof(T)));
}
}
}
// launch kernels repeatedly on each of the normal prioritiy stream
for (int k = 0; k < NORMAL_PRIORITY_STREAMCOUNT; ++k) {
for (int i = 0; i < size; i += MEMCPYSIZE1) {
int j = i / sizeof(T);
if (enable_priority_normal) {
hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), dim3(BLOCKSIZE),
0, stream_normal[k], dst_d_normal[k] + j, src_d_normal[k] + j,
(MEMCPYSIZE1 / sizeof(T)));
}
}
}
// launch kernels repeatedly on each of the high prioritiy stream
for (int k = 0; k < HIGH_PRIORITY_STREAMCOUNT; ++k) {
for (int i = 0; i < size; i += MEMCPYSIZE1) {
int j = i / sizeof(T);
if (enable_priority_high) {
hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), dim3(BLOCKSIZE),
0, stream_high[k], dst_d_high[k] + j, src_d_high[k] + j,
(MEMCPYSIZE1 / sizeof(T)));
}
}
}
// record end events for each of the low priority streams
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_low) {
HIP_CHECK(hipEventRecord(event_end_low[i], stream_low[i]));
}
}
// record end events for each of the normal priority streams
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_normal) {
HIP_CHECK(hipEventRecord(event_end_normal[i], stream_normal[i]));
}
}
// record end events for each of the high priority streams
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_high) {
HIP_CHECK(hipEventRecord(event_end_high[i], stream_high[i]));
}
}
// synchronize events for each of the low priority streams
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_low) {
HIP_CHECK(hipEventSynchronize(event_end_low[i]));
}
}
// synchronize events for each of the normal priority streams
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_normal) {
HIP_CHECK(hipEventSynchronize(event_end_normal[i]));
}
}
// synchronize events for each of the high priority streams
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_high) {
HIP_CHECK(hipEventSynchronize(event_end_high[i]));
}
}
// compute time spent for memcpy in each low stream
float time_spent_low[LOW_PRIORITY_STREAMCOUNT];
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_low) {
HIP_CHECK(hipEventElapsedTime(&time_spent_low[i], event_start_low[i],
event_end_low[i]));
}
}
// compute time spent for memcpy in each normal stream
float time_spent_normal[NORMAL_PRIORITY_STREAMCOUNT];
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_normal) {
HIP_CHECK(hipEventElapsedTime(&time_spent_normal[i],
event_start_normal[i], event_end_normal[i]));
}
}
// compute time spent for memcpy in each high stream
float time_spent_high[HIGH_PRIORITY_STREAMCOUNT];
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_high) {
HIP_CHECK(hipEventElapsedTime(&time_spent_high[i], event_start_high[i],
event_end_high[i]));
}
}
// sanity check for low priority streams
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_low) {
HIP_CHECK(hipMemcpy(dst_h_low[i], dst_d_low[i], size,
hipMemcpyDeviceToHost));
REQUIRE(memcmp(dst_h_low[i], src_h_low[i], size) == 0);
}
}
// sanity check for normal priority streams
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_normal) {
HIP_CHECK(hipMemcpy(dst_h_normal[i], dst_d_normal[i], size,
hipMemcpyDeviceToHost));
REQUIRE(memcmp(dst_h_normal[i], src_h_normal[i], size) == 0);
}
}
// sanity check for high priority streams
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
if (enable_priority_high) {
HIP_CHECK(hipMemcpy(dst_h_high[i], dst_d_high[i], size,
hipMemcpyDeviceToHost));
REQUIRE(memcmp(dst_h_high[i], src_h_high[i], size) == 0);
}
}
}
} // namespace hipStreamCreateWithPriorityTest
/**
Tests following scenarios.
1)Create streams with default flag for all available priority levels and
queue tasks in each of these streams, perform device synchronize and validate behavior.
2)Create streams with non-blocking flag for all available priority levels and
queue tasks in each of these streams, perform stream synchronize and validate behavior.
3)Create streams with default flag for all available priority levels and
queue tasks in each of these streams, perform stream synchronize and validate behavior.
4)Create streams with non-blocking flag for all available priority levels and
queue tasks in each of these streams, perform device synchronize and validate behavior.
*/
* Test Description
* ------------------------
* - Create streams with default flag for all available priority levels and
* queue tasks in each of these streams, perform device synchronize and validate
* behavior.
* - Create streams with non-blocking flag for all available priority levels
* and queue tasks in each of these streams, perform stream synchronize and
* validate behavior.
* - Create streams with default flag for all available priority levels and
* queue tasks in each of these streams, perform stream synchronize and validate
* behavior.
* - Create streams with non-blocking flag for all available priority levels
* and queue tasks in each of these streams, perform device synchronize and validate
* behavior.
* Test source
* ------------------------
* - catch\unit\stream\hipStreamCreateWithPriority.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamCreateWithPriority_FunctionalForAllPriorities") {
SECTION("Default flag and device synchronize") {
hipStreamCreateWithPriorityTest::
@@ -558,8 +851,16 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_FunctionalForAllPriorities") {
}
/**
* Create a stream for each priority level with default flag, Launch memcpy and kernel
* tasks on these streams from multiple threads. Validate all the results.
* Test Description
* ------------------------
* - Create a stream for each priority level with default flag, Launch
* memcpy and kernel tasks on these streams from multiple threads. Validate
* all the results.
* ------------------------
* - catch\unit\stream\hipStreamCreateWithPriority.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadDefaultflag") {
bool TestPassed = true;
@@ -569,8 +870,16 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadDefaultflag") {
}
/**
* Create a stream for each priority level with non-blocking flag, Launch memcpy and
* kernel tasks on these streams from multiple threads. Validate all the results.
* Test Description
* ------------------------
* - Create a stream for each priority level with non-blocking flag, Launch
* memcpy and kernel tasks on these streams from multiple threads. Validate all
* the results.
* ------------------------
* - catch\unit\stream\hipStreamCreateWithPriority.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag") {
bool TestPassed = true;
@@ -579,13 +888,17 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag") {
REQUIRE(TestPassed);
}
/**
* Scenario1: Validates functionality of hipStreamCreateWithPriority when
stream = nullptr.
* Scenario2: Validates functionality of hipStreamCreateWithPriority when
flag = 0xffffffff.
*/
* Test Description
* ------------------------
* - Validates functionality of hipStreamCreateWithPriority when stream = nullptr
* - Validates functionality of hipStreamCreateWithPriority when flag = 0xffffffff
* ------------------------
* - catch\unit\stream\hipStreamCreateWithPriority.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") {
hipStream_t stream{nullptr};
int priority_low{0};
@@ -601,15 +914,26 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") {
SECTION("stream = nullptr test") {
REQUIRE(hipErrorInvalidValue ==
hipStreamCreateWithPriority(nullptr, hipStreamDefault, priority_low));
hipStreamCreateWithPriority(nullptr, hipStreamDefault, priority_low));
}
SECTION("flag value invalid test") {
REQUIRE(hipErrorInvalidValue == hipStreamCreateWithPriority(&stream, 0xffffffff, priority_low));
REQUIRE(hipErrorInvalidValue ==
hipStreamCreateWithPriority(&stream, 0xffffffff, priority_low));
}
}
TEST_CASE("Unit_hipStreamCreateWithPriority") {
/**
* Test Description
* ------------------------
* - Set and Get Priority Value
* ------------------------
* - catch\unit\stream\hipStreamCreateWithPriority.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamCreateWithPriority_CheckPriorityVal") {
int id = GENERATE(range(0, HipTest::getDeviceCount()));
HIP_CHECK(hipSetDevice(id));
@@ -619,45 +943,73 @@ TEST_CASE("Unit_hipStreamCreateWithPriority") {
hipStream_t stream{nullptr};
SECTION("Setting high priority") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high));
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
priority_high));
REQUIRE(stream != nullptr);
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high));
}
SECTION("Setting low priority") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low));
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
priority_low));
REQUIRE(stream != nullptr);
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_low));
}
SECTION("Setting lowest possible priority") {
HIP_CHECK(
hipStreamCreateWithPriority(&stream, hipStreamDefault, std::numeric_limits<int>::max()));
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
std::numeric_limits<int>::max()));
REQUIRE(stream != nullptr);
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_low));
}
SECTION("Setting highest possible priority") {
HIP_CHECK(
hipStreamCreateWithPriority(&stream, hipStreamDefault, std::numeric_limits<int>::min()));
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
std::numeric_limits<int>::min()));
REQUIRE(stream != nullptr);
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high));
}
SECTION("Setting flags to hipStreamNonBlocking") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high));
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
priority_high));
REQUIRE(stream != nullptr);
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high, hipStreamNonBlocking));
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high,
hipStreamNonBlocking));
}
HIP_CHECK(hipStreamDestroy(stream));
}
/**
* Validate stream priorities with event after classifying them as low, medium and high.
* Test Description
* ------------------------
* - Validate stream priorities with event after classifying them as low,
* medium and high.
* ------------------------
* - catch\unit\stream\hipStreamCreateWithPriority.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamCreateWithPriority_ValidateWithEvents") {
bool TestPassed = true;
TestPassed = hipStreamCreateWithPriorityTest::validateStreamPrioritiesWithEvents<int>();
TestPassed =
hipStreamCreateWithPriorityTest::validateStreamPrioritiesWithEvents<int>();
REQUIRE(TestPassed);
}
/**
* Test Description
* ------------------------
* - Test that create Multiple streams with priority low, normal and high
* then these streams used to launch the kernel in random seq high,normal,low.
* ------------------------
* - catch\unit\stream\hipStreamCreateWithPriority.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamCreateWithPriority_TestMultipleStreamWithPriority") {
hipStreamCreateWithPriorityTest::TestForMultipleStreamWithPriority<int>();
}
@@ -17,25 +17,68 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_kernels.hh>
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_defgroups.hh>
#define NUMBER_OF_THREADS 10
static bool thread_results[NUMBER_OF_THREADS];
/**
* @addtogroup hipStreamGetDevice hipStreamGetDevice
* @{
* @ingroup StreamTest
* `hipError_t hipStreamGetDevice(hipStream_t stream, hipDevice_t* device)` -
* Get the device assocaited with the stream.
* @}
*/
/**
* Test Description
* ------------------------
* - NegativeTest case.
* Pass device as nullptr and verify hipErrorInvalidValue
* Pass stream as nullptr and verify hipErrorInvalidValue
* Pass device as nullptr for hipStreamPerThread and verify hipErrorInvalidValue
* Test source
* ------------------------
* - catch/unit/stream/hipStreamGetDevice.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipStreamGetDevice_Negative") {
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK_ERROR(hipStreamGetDevice(nullptr, nullptr), hipErrorInvalidValue); // null stream
HIP_CHECK_ERROR(hipStreamGetDevice(nullptr, nullptr), hipErrorInvalidValue);
HIP_CHECK_ERROR(hipStreamGetDevice(hipStreamPerThread, nullptr),
hipErrorInvalidValue); // stream per thread
HIP_CHECK_ERROR(hipStreamGetDevice(stream, nullptr), hipErrorInvalidValue); // created stream
hipErrorInvalidValue);
HIP_CHECK_ERROR(hipStreamGetDevice(stream, nullptr), hipErrorInvalidValue);
HIP_CHECK(hipStreamDestroy(stream));
}
// Iterate over all devices, create stream on the device and match the device we get from stream
/**
* Test Description
* ------------------------
* - NegativeTest case.
* Test case to validate hipStreamGetDevice on user created stream,
* Null stream and hipStreamPerThread.
* Test source
* ------------------------
* - catch/unit/stream/hipStreamGetDevice.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipStreamGetDevice_Usecase") {
int device_count = 0;
HIP_CHECK(hipGetDeviceCount(&device_count));
REQUIRE(device_count != 0); // atleast 1 device
REQUIRE(device_count != 0);
SECTION("Null Stream") {
CTX_CREATE();
@@ -43,7 +86,7 @@ TEST_CASE("Unit_hipStreamGetDevice_Usecase") {
hipDevice_t device_from_stream, device_from_ordinal;
HIP_CHECK(hipStreamGetDevice(nullptr, &device_from_stream));
HIP_CHECK(hipDeviceGet(&device_from_ordinal, 0)); // default device
HIP_CHECK(hipDeviceGet(&device_from_ordinal, 0));
REQUIRE(device_from_stream == device_from_ordinal);
CTX_DESTROY();
@@ -55,7 +98,7 @@ TEST_CASE("Unit_hipStreamGetDevice_Usecase") {
hipDevice_t device_from_stream, device_from_ordinal;
HIP_CHECK(hipStreamGetDevice(hipStreamPerThread, &device_from_stream));
HIP_CHECK(hipDeviceGet(&device_from_ordinal, 0)); // default device
HIP_CHECK(hipDeviceGet(&device_from_ordinal, 0));
REQUIRE(device_from_stream == device_from_ordinal);
CTX_DESTROY();
@@ -72,7 +115,128 @@ TEST_CASE("Unit_hipStreamGetDevice_Usecase") {
HIP_CHECK(hipStreamGetDevice(stream, &device_from_stream));
HIP_CHECK(hipDeviceGet(&device_from_ordinal, i));
REQUIRE(device_from_stream == device_from_ordinal); // maybe match uuid??
REQUIRE(device_from_stream == device_from_ordinal);
}
}
}
/**
* Test Description
* ------------------------
* - NegativeTest case.
* Test case to multi-threaded scenario for hipStreamGetDevice
* Test source
* ------------------------
* - catch/unit/stream/hipStreamGetDevice.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
static bool validateStreamGetDevice() {
int gpu = 0;
hipDevice_t device_from_stream;
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamGetDevice(stream, &device_from_stream));
REQUIRE(device_from_stream == gpu);
return true;
}
static void thread_Test(int threadNum) {
thread_results[threadNum] = validateStreamGetDevice();
}
static bool test_hipStreamGetDevice_MThread() {
std::vector<std::thread> tests;
// Spawn the test threads
for (int idx = 0; idx < NUMBER_OF_THREADS; idx++) {
thread_results[idx] = false;
tests.push_back(std::thread(thread_Test, idx));
}
// Wait for all threads to complete
for (std::thread &t : tests) {
t.join();
}
// Wait for thread
bool status = true;
for (int idx = 0; idx < NUMBER_OF_THREADS; idx++) {
status = status & thread_results[idx];
}
return status;
}
TEST_CASE("Unit_hipStreamGetDevice_MThread") {
REQUIRE(true == test_hipStreamGetDevice_MThread());
}
/**
* Test Description
* ------------------------
* - NegativeTest case.
* Create a stream with gpu1, then set device to gpu0
* call hipStreamGetDevice on the stream and verify the
* returned device is same as initial set device.
* Test source
* ------------------------
* - catch/unit/stream/hipStreamGetDevice.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipStreamGetDevice_SetDiffDevice") {
hipDevice_t device_from_stream;
int device_count = 0;
HIP_CHECK(hipGetDeviceCount(&device_count));
if (device_count < 2) {
HipTest::HIP_SKIP_TEST("Skipping because devices < 2");
return;
}
for (int i = 0; i < device_count; ++i) {
HIP_CHECK(hipSetDevice(i));
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
for (int j = 0; j < device_count; ++j) {
if (i != j) {
HIP_CHECK(hipSetDevice(j));
HIP_CHECK(hipStreamGetDevice(stream, &device_from_stream));
REQUIRE(device_from_stream == i);
}
}
}
}
/**
* Test Description
* ------------------------
* - NegativeTest case.
* Set hip Devices to each available gpu and probe
* hipStreamGetDevice with null stream.
* Test source
* ------------------------
* - catch/unit/stream/hipStreamGetDevice.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
* Test to be run only on AMD machine as it's failing in CUDA.
*/
#if HT_AMD
TEST_CASE("Unit_hipStreamGetDevice_NullStream") {
int device_count = 0;
HIP_CHECK(hipGetDeviceCount(&device_count));
REQUIRE(device_count != 0);
for (int i = 0; i < device_count; i++) {
HIP_CHECK(hipSetDevice(i));
hipDevice_t device_from_stream;
HIP_CHECK(hipStreamGetDevice(0, &device_from_stream));
REQUIRE(device_from_stream == i);
}
}
#endif
@@ -0,0 +1,30 @@
# Copyright (c) 2023 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.
# Common Tests - Test independent of all platforms
set(TEST_SRC
hipSurfaceObj1D.cc
hipSurfaceObj2D.cc
hipSurfaceObj3D.cc
)
hip_add_exe_to_target(NAME SurfaceTest
TEST_SRC ${TEST_SRC}
TEST_TARGET_NAME build_tests)
@@ -0,0 +1,317 @@
/*
Copyright (c) 2023 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.
*/
#include <hip_test_common.hh>
#include <hip_array_common.hh>
#include <hip_texture_helper.hh>
template <typename T>
__global__ void
surf1DKernelR(hipSurfaceObject_t surfaceObject,
T* outputData, int width)
{
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
if (x < width) {
surf1Dread(outputData + x, surfaceObject, x * sizeof(T));
}
#endif
}
template <typename T>
__global__ void
surf1DKernelW(hipSurfaceObject_t surfaceObject,
T* inputData, int width)
{
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
if (x < width) {
surf1Dwrite(inputData[x], surfaceObject, x * sizeof(T));
}
#endif
}
template <typename T>
__global__ void
surf1DKernelRW(hipSurfaceObject_t surfaceObject,
hipSurfaceObject_t outputSurfObj, int width)
{
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
if (x < width) {
T data;
surf1Dread(&data, surfaceObject, x * sizeof(T));
surf1Dwrite(data, outputSurfObj, x * sizeof(T));
}
#endif
}
template <typename T>
static void runTestR(const int width)
{
unsigned int size = width * sizeof(T);
T *hData = (T*) malloc (size);
memset(hData, 0, size);
for (int j = 0; j < width; j++)
{
initVal(hData[j]);
}
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
hipArray *hipArray = nullptr;
HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, 0, hipArraySurfaceLoadStore));
HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Create surface object
hipSurfaceObject_t surfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
T *hOutputData = nullptr;
HIP_CHECK(hipHostMalloc((void**)&hOutputData, size));
memset(hOutputData, 0, size);
dim3 dimBlock (16, 1, 1);
dim3 dimGrid ((width + dimBlock.x - 1) / dimBlock.x, 1, 1);
surf1DKernelR<T><<<dimGrid, dimBlock>>>(surfaceObject, hOutputData, width);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
for (int j = 0; j < width; j++) {
if (!isEqual(hData[j], hOutputData[j])) {
printf("Difference [ %d ]:%s ----%s\n", j,
getString(hData[j]).c_str(), getString(hOutputData[j]).c_str());
REQUIRE(false);
}
}
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
HIP_CHECK(hipFreeArray(hipArray));
free(hData);
HIP_CHECK(hipHostFree(hOutputData));
REQUIRE(true);
}
template <typename T>
static void runTestW(const int width)
{
unsigned int size = width * sizeof(T);
T *hData = nullptr;
HIP_CHECK(hipHostMalloc((void**)&hData, size));
memset(hData, 0, size);
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
hipArray *hipArray = nullptr;
HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, 0, hipArraySurfaceLoadStore));
HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Create surface object
hipSurfaceObject_t surfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
for (int j = 0; j < width; j++)
{
initVal(hData[j]);
}
dim3 dimBlock (16, 1, 1);
dim3 dimGrid ((width + dimBlock.x - 1) / dimBlock.x, 1, 1);
surf1DKernelW<T><<<dimGrid, dimBlock>>>(surfaceObject, hData, width);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
T *hOutputData = (T*) malloc (size);
memset(hOutputData, 0, size);
HIP_CHECK(hipMemcpyFromArray(hOutputData, hipArray, 0, 0, size, hipMemcpyDeviceToHost));
for (int j = 0; j < width; j++) {
if (!isEqual(hData[j], hOutputData[j])) {
printf("Difference [ %d ]:%s ----%s\n", j,
getString(hData[j]).c_str(), getString(hOutputData[j]).c_str());
REQUIRE(false);
}
}
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
HIP_CHECK(hipFreeArray(hipArray));
HIP_CHECK(hipHostFree(hData));
free(hOutputData);
REQUIRE(true);
}
template <typename T>
static void runTestRW(const int width)
{
unsigned int size = width * sizeof(T);
T *hData = (T*) malloc (size);
memset(hData, 0, size);
for (int j = 0; j < width; j++)
{
initVal(hData[j]);
}
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
hipArray *hipArray = nullptr, *hipOutArray = nullptr;
HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, 0, hipArraySurfaceLoadStore));
HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Create surface object
hipSurfaceObject_t surfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
HIP_CHECK(hipMallocArray(&hipOutArray, &channelDesc, width, 0, hipArraySurfaceLoadStore));
hipResourceDesc resOutDesc;
memset(&resOutDesc, 0, sizeof(resOutDesc));
resOutDesc.resType = hipResourceTypeArray;
resOutDesc.res.array.array = hipOutArray;
hipSurfaceObject_t outSurfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject (&outSurfaceObject, &resOutDesc));
dim3 dimBlock (16, 1, 1);
dim3 dimGrid ((width + dimBlock.x - 1) / dimBlock.x, 1, 1);
surf1DKernelRW<T><<<dimGrid, dimBlock>>>(surfaceObject, outSurfaceObject, width);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
T *hOutputData = (T*) malloc (size);
memset(hOutputData, 0, size);
HIP_CHECK(hipMemcpyFromArray(hOutputData, hipOutArray, 0, 0, size, hipMemcpyDeviceToHost));
for (int j = 0; j < width; j++) {
if (!isEqual(hData[j], hOutputData[j])) {
printf("Difference [ %d ]:%s ----%s\n", j,
getString(hData[j]).c_str(), getString(hOutputData[j]).c_str());
REQUIRE(false);
}
}
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
HIP_CHECK(hipDestroySurfaceObject(outSurfaceObject));
HIP_CHECK(hipFreeArray(hipArray));
HIP_CHECK(hipFreeArray(hipOutArray));
free(hData);
free(hOutputData);
REQUIRE(true);
}
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj1D_type_R", "",
char, uchar, short, ushort, int, uint, float,
char1, uchar1, short1, ushort1, int1, uint1, float1,
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4)
{
CHECK_IMAGE_SUPPORT
auto err = hipGetLastError(); // reset last err due to previous negative tests
SECTION("Unit_hipSurfaceObj1D_type_R - 31") {
runTestR<TestType>(31);
}
SECTION("Unit_hipSurfaceObj1D_type_R - 67") {
runTestR<TestType>(67);
}
SECTION("Unit_hipSurfaceObj1D_type_R - 131") {
runTestR<TestType>(131);
}
SECTION("Unit_hipSurfaceObj1D_type_R - 263") {
runTestR<TestType>(263);
}
}
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj1D_type_W", "",
char, uchar, short, ushort, int, uint, float,
char1, uchar1, short1, ushort1, int1, uint1, float1,
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4)
{
CHECK_IMAGE_SUPPORT
auto err = hipGetLastError(); // reset last err due to previous negative tests
SECTION("Unit_hipSurfaceObj1D_type_W - 31") {
runTestW<TestType>(31);
}
SECTION("Unit_hipSurfaceObj1D_type_W - 63") {
runTestW<TestType>(63);
}
SECTION("Unit_hipSurfaceObj1D_type_W - 131") {
runTestW<TestType>(131);
}
SECTION("Unit_hipSurfaceObj1D_type_W - 263") {
runTestW<TestType>(263);
}
}
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj1D_type_RW", "",
char, uchar, short, ushort, int, uint, float,
char1, uchar1, short1, ushort1, int1, uint1, float1,
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4)
{
CHECK_IMAGE_SUPPORT
auto err = hipGetLastError(); // reset last err due to previous negative tests
SECTION("Unit_hipSurfaceObj1D_type_RW - 23") {
runTestRW<TestType>(23);
}
SECTION("Unit_hipSurfaceObj1D_type_RW - 67") {
runTestRW<TestType>(67);
}
SECTION("Unit_hipSurfaceObj1D_type_RW - 131") {
runTestRW<TestType>(131);
}
SECTION("Unit_hipSurfaceObj1D_type_RW - 263") {
runTestRW<TestType>(263);
}
}
@@ -0,0 +1,364 @@
/*
Copyright (c) 2023 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.
*/
#include <hip_test_common.hh>
#include <hip_array_common.hh>
#include <hip_texture_helper.hh>
#define LOG_DATA 0
template <typename T>
__global__ void
surf2DKernelR(hipSurfaceObject_t surfaceObject,
T* outputData, int width, int height)
{
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < width && y < height) {
surf2Dread<T>(outputData + y * width + x, surfaceObject, x * sizeof(T), y);
}
#endif
}
template <typename T>
__global__ void
surf2DKernelW(hipSurfaceObject_t surfaceObject,
T* inputData, int width, int height)
{
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < width && y < height) {
surf2Dwrite<T>(inputData[y * width + x], surfaceObject, x * sizeof(T), y);
}
#endif
}
template <typename T>
__global__ void
surf2DKernelRW(hipSurfaceObject_t surfaceObject,
hipSurfaceObject_t outputSurfObj, int width, int height)
{
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < width && y < height) {
T data;
surf2Dread<T>(&data, surfaceObject, x * sizeof(T), y);
surf2Dwrite<T>(data, outputSurfObj, x * sizeof(T), y);
}
#endif
}
template <typename T>
static void runTestR(const int width, const int height)
{
unsigned int size = width * height * sizeof(T);
T* hData = (T*) malloc(size);
memset(hData, 0, size);
for (int i = 0; i < height; i++)
{
for (int j = 0; j < width; j++)
{
initVal(hData[i * width + j]);
}
}
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
hipArray *hipArray = nullptr;
HIP_CHECK(hipMallocArray (&hipArray, &channelDesc, width, height,
hipArraySurfaceLoadStore));
// Need set source pitch, but we don't have any padding here
const size_t spitch = width * sizeof(T);
HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, spitch, spitch, height,
hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Create surface object
hipSurfaceObject_t surfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
T* hOutputData = nullptr;
HIP_CHECK(hipHostMalloc((void**)&hOutputData, size));
memset(hOutputData, 0, size);
dim3 dimBlock (16, 16, 1);
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1);
surf2DKernelR<T><<<dimGrid, dimBlock>>>(surfaceObject, hOutputData, width, height);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
int index = i * width + j;
if (!isEqual(hData[index], hOutputData[index])) {
printf("Difference [ %d %d ]:%s ----%s\n", i, j,
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
REQUIRE(false);
}
}
}
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
HIP_CHECK(hipFreeArray(hipArray));
free(hData);
HIP_CHECK(hipHostFree(hOutputData));
REQUIRE(true);
}
template <typename T>
static void runTestW(const int width, const int height)
{
unsigned int size = width * height * sizeof(T);
T* hData = nullptr;
HIP_CHECK(hipHostMalloc((void**)&hData, size));
memset(hData, 0, size);
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
hipArray *hipArray = nullptr;
HIP_CHECK(hipMallocArray (&hipArray, &channelDesc, width, height,
hipArraySurfaceLoadStore));
// Need set source pitch, but we don't have any padding here
const size_t spitch = width * sizeof(T);
HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, spitch, spitch, height,
hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Create surface object
hipSurfaceObject_t surfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
for (int i = 0; i < height; i++)
{
for (int j = 0; j < width; j++)
{
initVal(hData[i * width + j]);
}
}
dim3 dimBlock (16, 16, 1);
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1);
surf2DKernelW<T><<<dimGrid, dimBlock>>>(surfaceObject, hData, width, height);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
T* hOutputData = (T*) malloc(size);
memset(hOutputData, 0, size);
HIP_CHECK(hipMemcpy2DFromArray(hOutputData, spitch, hipArray, 0, 0, spitch,
height, hipMemcpyDeviceToHost));
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
int index = i * width + j;
if (!isEqual(hData[index], hOutputData[index])) {
printf("Difference [ %d %d ]:%s ----%s\n", i, j,
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
REQUIRE(false);
}
}
}
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
HIP_CHECK(hipFreeArray(hipArray));
HIP_CHECK(hipHostFree(hData));
free(hOutputData);
REQUIRE(true);
}
template <typename T>
static void runTestRW(const int width, const int height)
{
unsigned int size = width * height * sizeof(T);
T* hData = (T*) malloc(size);
memset(hData, 0, size);
for (int i = 0; i < height; i++)
{
for (int j = 0; j < width; j++)
{
initVal(hData[i * width + j]);
}
}
#if LOG_DATA
printf ("hData: ");
for (int i = 0; i < 32; i++)
{
printf ("%s ", getString(hData[i]).c_str());
}
printf ("\n");
#endif
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
hipArray *hipArray = nullptr, *hipOutArray = nullptr;
HIP_CHECK(hipMallocArray (&hipArray, &channelDesc, width, height,
hipArraySurfaceLoadStore));
// Need set source pitch, but we don't have any padding here
const size_t spitch = width * sizeof(T);
HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, spitch, spitch, height,
hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Create surface object
hipSurfaceObject_t surfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
HIP_CHECK(hipMallocArray(&hipOutArray, &channelDesc, width, height,
hipArraySurfaceLoadStore));
hipResourceDesc resOutDesc;
memset(&resOutDesc, 0, sizeof(resOutDesc));
resOutDesc.resType = hipResourceTypeArray;
resOutDesc.res.array.array = hipOutArray;
hipSurfaceObject_t outSurfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject (&outSurfaceObject, &resOutDesc));
dim3 dimBlock (16, 16, 1);
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1);
surf2DKernelRW<T><<<dimGrid, dimBlock>>>(surfaceObject, outSurfaceObject, width, height);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
T* hOutputData = (T*) malloc(size);
memset(hOutputData, 0, size);
HIP_CHECK(hipMemcpy2DFromArray(hOutputData, spitch, hipOutArray, 0, 0, spitch,
height, hipMemcpyDeviceToHost));
#if LOG_DATA
printf ("dData: ");
for (int i = 0; i < 32; i++)
{
printf ("%s ", getString(hOutputData[i]).c_str());
}
printf ("\n");
#endif
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
int index = i * width + j;
if (!isEqual(hData[index], hOutputData[index])) {
printf("Difference [ %d %d ]:%s ----%s\n", i, j,
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
REQUIRE(false);
}
}
}
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
HIP_CHECK(hipDestroySurfaceObject(outSurfaceObject));
HIP_CHECK(hipFreeArray(hipArray));
HIP_CHECK(hipFreeArray(hipOutArray));
free(hData);
free(hOutputData);
REQUIRE(true);
}
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj2D_type_R", "",
char, uchar, short, ushort, int, uint, float,
char1, uchar1, short1, ushort1, int1, uint1, float1,
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4)
{
CHECK_IMAGE_SUPPORT
auto err = hipGetLastError(); // reset last err due to previous negative tests
SECTION("Unit_hipSurfaceObj2D_type_R - 23, 67") {
runTestR<TestType>(23, 67);
}
SECTION("Unit_hipSurfaceObj2D_type_R - 67, 23") {
runTestR<TestType>(67, 23);
}
SECTION("Unit_hipSurfaceObj2D_type_R - 131, 67") {
runTestR<TestType>(131, 67);
}
SECTION("Unit_hipSurfaceObj2D_type_R - 263, 131") {
runTestR<TestType>(263, 131);
}
}
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj2D_type_W", "",
char, uchar, short, ushort, int, uint, float,
char1, uchar1, short1, ushort1, int1, uint1, float1,
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4)
{
CHECK_IMAGE_SUPPORT
auto err = hipGetLastError(); // reset last err due to previous negative tests
SECTION("Unit_hipSurfaceObj2D_type_W - 23, 67") {
runTestW<TestType>(23, 67);
}
SECTION("Unit_hipSurfaceObj2D_type_W - 67, 23") {
runTestW<TestType>(67, 23);
}
SECTION("Unit_hipSurfaceObj2D_type_W - 131, 67") {
runTestW<TestType>(131, 67);
}
SECTION("Unit_hipSurfaceObj2D_type_W - 263, 23") {
runTestW<TestType>(263, 23);
}
}
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj2D_type_RW", "",
char, uchar, short, ushort, int, uint, float,
char1, uchar1, short1, ushort1, int1, uint1, float1,
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4)
{
CHECK_IMAGE_SUPPORT
auto err = hipGetLastError(); // reset last err due to previous negative tests
SECTION("Unit_hipSurfaceObj2D_type_RW - 23, 67") {
runTestRW<TestType>(23, 67);
}
SECTION("Unit_hipSurfaceObj2D_type_RW - 67, 131") {
runTestRW<TestType>(67, 131);
}
SECTION("Unit_hipSurfaceObj2D_type_RW - 131, 263") {
runTestRW<TestType>(131, 263);
}
SECTION("Unit_hipSurfaceObj2D_type_RW - 263, 67") {
runTestRW<TestType>(263, 67);
}
}
@@ -0,0 +1,400 @@
/*
Copyright (c) 2023 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.
*/
#include <hip_test_common.hh>
#include <hip_array_common.hh>
#include <hip_texture_helper.hh>
template <typename T>
__global__ void
surf3DKernelR(hipSurfaceObject_t surfaceObject,
T* outputData, int width, int height, int depth)
{
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
if (x < width && y < height && z < depth) {
surf3Dread(outputData + z * width * height + y * width + x,
surfaceObject, x * sizeof(T), y, z);
}
#endif
}
template <typename T>
__global__ void
surf3DKernelW(hipSurfaceObject_t surfaceObject,
T* inputData, int width, int height, int depth)
{
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
if (x < width && y < height && z < depth) {
surf3Dwrite(inputData[z * width * height + y * width + x],
surfaceObject, x * sizeof(T), y, z);
}
#endif
}
template <typename T>
__global__ void
surf3DKernelRW(hipSurfaceObject_t surfaceObject,
hipSurfaceObject_t outputSurfObj, int width, int height, int depth)
{
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
if (x < width && y < height && z < depth) {
T data;
surf3Dread(&data, surfaceObject, x * sizeof(T), y, z);
surf3Dwrite(data, outputSurfObj, x * sizeof(T), y, z);
}
#endif
}
template <typename T>
static void runTestR(const int width, const int height, const int depth)
{
unsigned int size = width * height * depth * sizeof(T);
T *hData = (T*) malloc(size);
memset(hData, 0, size);
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
initVal(hData[i * width * height + j * width + k]);
}
}
}
// Allocate array and copy image data
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
hipArray *hipArray = nullptr;
HIP_CHECK(hipMalloc3DArray(&hipArray, &channelDesc, make_hipExtent(width, height, depth),
hipArraySurfaceLoadStore));
hipMemcpy3DParms myparms;
memset(&myparms, 0, sizeof(myparms));
myparms.srcPos = make_hipPos(0,0,0);
myparms.dstPos = make_hipPos(0,0,0);
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
myparms.dstArray = hipArray;
myparms.extent = make_hipExtent(width, height, depth);
myparms.kind = hipMemcpyHostToDevice;
HIP_CHECK(hipMemcpy3D(&myparms));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Create surface object
hipSurfaceObject_t surfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
T *hOutputData = nullptr;
HIP_CHECK(hipHostMalloc((void**)&hOutputData, size));
memset(hOutputData, 0, size);
dim3 dimBlock(8, 8, 8); // 512 threads
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y,
(depth + dimBlock.z - 1) / dimBlock.z);
surf3DKernelR<T><<<dimGrid, dimBlock>>>(surfaceObject, hOutputData, width, height, depth);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
int index = i * width * height + j * width + k;
if (!isEqual(hData[index], hOutputData[index])) {
printf("Difference [ %d %d %d]:%s ----%s\n", i, j, k,
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
REQUIRE(false);
}
}
}
}
HIP_CHECK(hipDestroySurfaceObject (surfaceObject));
HIP_CHECK(hipFreeArray(hipArray));
free(hData);
HIP_CHECK(hipHostFree(hOutputData));
REQUIRE(true);
}
template <typename T>
static void runTestW(const int width, const int height, const int depth)
{
unsigned int size = width * height * depth * sizeof(T);
T *hData = nullptr;
HIP_CHECK(hipHostMalloc((void**)&hData, size));
memset(hData, 0, size);
// Allocate array and copy image data
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
hipArray *hipArray = nullptr;
HIP_CHECK(hipMalloc3DArray(&hipArray, &channelDesc, make_hipExtent(width, height, depth),
hipArraySurfaceLoadStore));
hipMemcpy3DParms myparms;
memset(&myparms, 0, sizeof(myparms));
myparms.srcPos = make_hipPos(0,0,0);
myparms.dstPos = make_hipPos(0,0,0);
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
myparms.dstArray = hipArray;
myparms.extent = make_hipExtent(width, height, depth);
myparms.kind = hipMemcpyHostToDevice;
HIP_CHECK(hipMemcpy3D(&myparms));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Create surface object
hipSurfaceObject_t surfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
initVal(hData[i * width * height + j * width + k]);
}
}
}
dim3 dimBlock(8, 8, 8); // 512 threads
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y,
(depth + dimBlock.z - 1) / dimBlock.z);
surf3DKernelW<T><<<dimGrid, dimBlock>>>(surfaceObject, hData, width, height, depth);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
T *hOutputData = (T*) malloc (size);
memset(hOutputData, 0, size);
memset(&myparms, 0, sizeof(myparms));
myparms.srcPos = make_hipPos(0,0,0);
myparms.dstPos = make_hipPos(0,0,0);
myparms.srcArray= hipArray;
myparms.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(T), width, height);
myparms.extent = make_hipExtent(width, height, depth);
myparms.kind = hipMemcpyDeviceToHost;
HIP_CHECK(hipMemcpy3D(&myparms));
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
int index = i * width * height + j * width + k;
if (!isEqual(hData[index], hOutputData[index])) {
printf("Difference [ %d %d %d]:%s ----%s\n", i, j, k,
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
REQUIRE(false);
}
}
}
}
HIP_CHECK(hipDestroySurfaceObject (surfaceObject));
HIP_CHECK(hipFreeArray(hipArray));
HIP_CHECK(hipHostFree(hData));
free(hOutputData);
REQUIRE(true);
}
template <typename T>
static void runTestRW(const int width, const int height, const int depth)
{
unsigned int size = width * height * depth * sizeof(T);
T *hData = (T*) malloc(size);
memset(hData, 0, size);
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
initVal(hData[i * width * height + j * width + k]);
}
}
}
// Allocate array and copy image data
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
hipArray *hipArray = nullptr, *hipOutArray = nullptr;
HIP_CHECK(hipMalloc3DArray(&hipArray, &channelDesc, make_hipExtent(width, height, depth),
hipArraySurfaceLoadStore));
hipMemcpy3DParms myparms;
memset(&myparms, 0, sizeof(myparms));
myparms.srcPos = make_hipPos(0,0,0);
myparms.dstPos = make_hipPos(0,0,0);
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
myparms.dstArray = hipArray;
myparms.extent = make_hipExtent(width, height, depth);
myparms.kind = hipMemcpyHostToDevice;
HIP_CHECK(hipMemcpy3D(&myparms));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Create surface object
hipSurfaceObject_t surfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
HIP_CHECK(hipMalloc3DArray(&hipOutArray, &channelDesc, make_hipExtent(width, height, depth),
hipArraySurfaceLoadStore));
hipResourceDesc resOutDesc;
memset(&resOutDesc, 0, sizeof(resOutDesc));
resOutDesc.resType = hipResourceTypeArray;
resOutDesc.res.array.array = hipOutArray;
hipSurfaceObject_t outSurfaceObject = 0;
HIP_CHECK(hipCreateSurfaceObject(&outSurfaceObject, &resOutDesc));
dim3 dimBlock(8, 8, 8); // 512 threads
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y,
(depth + dimBlock.z - 1) / dimBlock.z);
surf3DKernelRW<T><<<dimGrid, dimBlock>>>(surfaceObject, outSurfaceObject, width, height, depth);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
T *hOutputData = (T*) malloc (size);
memset(hOutputData, 0, size);
memset(&myparms, 0, sizeof(myparms));
myparms.srcPos = make_hipPos(0,0,0);
myparms.dstPos = make_hipPos(0,0,0);
myparms.srcArray= hipOutArray;
myparms.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(T), width, height);
myparms.extent = make_hipExtent(width, height, depth);
myparms.kind = hipMemcpyDeviceToHost;
HIP_CHECK(hipMemcpy3D(&myparms));
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
int index = i * width * height + j * width + k;
if (!isEqual(hData[index], hOutputData[index])) {
printf("Difference [ %d %d %d]:%s ----%s\n", i, j, k,
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
REQUIRE(false);
}
}
}
}
HIP_CHECK(hipDestroySurfaceObject (surfaceObject));
HIP_CHECK(hipDestroySurfaceObject (outSurfaceObject));
HIP_CHECK(hipFreeArray(hipArray));
HIP_CHECK(hipFreeArray(hipOutArray));
free(hData);
free(hOutputData);
REQUIRE(true);
}
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj3D_type_R", "",
char, uchar, short, ushort, int, uint, float,
char1, uchar1, short1, ushort1, int1, uint1, float1,
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4)
{
CHECK_IMAGE_SUPPORT
auto err = hipGetLastError(); // reset last err due to previous negative tests
SECTION("Unit_hipSurfaceObj3D_type_R - 31, 67, 131") {
runTestR<TestType>(31, 67, 131);
}
SECTION("Unit_hipSurfaceObj3D_type_R - 67, 31, 263") {
runTestR<TestType>(67, 31, 263);
}
SECTION("Unit_hipSurfaceObj3D_type_R - 131, 131, 67") {
runTestR<TestType>(131, 131, 67);
}
SECTION("Unit_hipSurfaceObj3D_type_R - 263, 131, 263") {
runTestR<TestType>(263, 131, 263);
}
}
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj3D_type_W", "",
char, uchar, short, ushort, int, uint, float,
char1, uchar1, short1, ushort1, int1, uint1, float1,
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4)
{
CHECK_IMAGE_SUPPORT
auto err = hipGetLastError(); // reset last err due to previous negative tests
SECTION("Unit_hipSurfaceObj3D_type_W - 31, 67, 131") {
runTestW<TestType>(31, 67, 131);
}
SECTION("Unit_hipSurfaceObj3D_type_W - 67, 67, 31") {
runTestW<TestType>(67, 67, 31);
}
SECTION("Unit_hipSurfaceObj3D_type_W - 131, 131, 67") {
runTestW<TestType>(131, 131, 67);
}
SECTION("Unit_hipSurfaceObj3D_type_W - 263, 131, 263") {
runTestW<TestType>(263, 131, 263);
}
}
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj3D_type_RW", "",
char, uchar, short, ushort, int, uint, float,
char1, uchar1, short1, ushort1, int1, uint1, float1,
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4)
{
CHECK_IMAGE_SUPPORT
auto err = hipGetLastError(); // reset last err due to previous negative tests
SECTION("Unit_hipSurfaceObj3D_type_RW - 31, 31, 67") {
runTestRW<TestType>(31, 31, 67);
}
SECTION("Unit_hipSurfaceObj3D_type_RW - 67, 67, 31") {
runTestRW<TestType>(67, 67, 31);
}
SECTION("Unit_hipSurfaceObj3D_type_RW - 131, 67, 263") {
runTestRW<TestType>(131, 67, 263);
}
SECTION("Unit_hipSurfaceObj3D_type_RW - 263, 131, 263") {
runTestRW<TestType>(263, 131, 263);
}
}
@@ -18,6 +18,7 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_array_common.hh>
#include <vector>
#include <iostream>
@@ -31,47 +32,6 @@ __global__ void tex1dKernelFetch(T *val, hipTextureObject_t obj, int N) {
#endif
}
template <typename T>
static inline __host__ __device__ constexpr int rank() {
return sizeof(T) / sizeof(decltype(T::x));
}
template<typename T>
static inline T getRandom() {
double r = 0;
if (std::is_signed < T > ::value) {
r = (std::rand() - RAND_MAX / 2.0) / (RAND_MAX / 2.0 + 1.);
} else {
r = std::rand() / (RAND_MAX + 1.);
}
return static_cast<T>(std::numeric_limits < T > ::max() * r);
}
template<
typename T,
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
static inline void initVal(T &val) {
val.x = getRandom<decltype(T::x)>();
}
template<
typename T,
typename std::enable_if<rank<T>() == 2>::type* = nullptr>
static inline void initVal(T &val) {
val.x = getRandom<decltype(T::x)>();
val.y = getRandom<decltype(T::x)>();
}
template<
typename T,
typename std::enable_if<rank<T>() == 4>::type* = nullptr>
static inline void initVal(T &val) {
val.x = getRandom<decltype(T::x)>();
val.y = getRandom<decltype(T::x)>();
val.z = getRandom<decltype(T::x)>();
val.w = getRandom<decltype(T::x)>();
}
template<
typename T,
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
@@ -112,31 +72,6 @@ static inline void printVector(T &val) {
std::cout << ")";
}
template<
typename T,
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
static inline bool isEqual(const T &val0, const T &val1) {
return val0.x == val1.x;
}
template<
typename T,
typename std::enable_if<rank<T>() == 2>::type* = nullptr>
static inline bool isEqual(const T &val0, const T &val1) {
return val0.x == val1.x &&
val0.y == val1.y;
}
template<
typename T,
typename std::enable_if<rank<T>() == 4>::type* = nullptr>
static inline bool isEqual(const T &val0, const T &val1) {
return val0.x == val1.x &&
val0.y == val1.y &&
val0.z == val1.z &&
val0.w == val1.w;
}
template<typename T>
bool runTest() {
const int N = 1024;
@@ -144,7 +79,7 @@ bool runTest() {
// Allocating the required buffer on gpu device
T *texBuf, *texBufOut;
T val[N], output[N];
hipGetLastError(); // Clear err due to negative tests
memset(output, 0, sizeof(output));
std::srand(std::time(nullptr)); // use current time as seed for random generator
@@ -35,7 +35,7 @@ hipDispatchLatency.out: hipDispatchLatency.cpp
$(HIPCC) $(CXXFLAGS) hipDispatchLatency.cpp -o $@
hipDispatchEnqueueRateMT.out: hipDispatchEnqueueRateMT.cpp
$(HIPCC) $(CXXFLAGS) hipDispatchEnqueueRateMT.cpp -o $@
$(HIPCC) $(CXXFLAGS) hipDispatchEnqueueRateMT.cpp -lpthread -o $@
test_kernel.code: test_kernel.cpp
$(HIP_PATH)/bin/hipcc --genco $(GENCO_FLAGS) $^ -o $@
@@ -37,7 +37,7 @@ endif()
add_library(HipOptLibrary STATIC ${CPP_SOURCES})
# Set-up the correct flags to generate the static library.
target_link_libraries(HipOptLibrary PRIVATE --emit-static-lib)
target_link_options(HipOptLibrary PRIVATE --emit-static-lib)
target_include_directories(HipOptLibrary PRIVATE /opt/rocm/hsa/include)
# Create test executable that uses libHipOptLibrary.a
@@ -23,40 +23,47 @@
#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>
#include <iostream>
#include <stdexcept>
#define HIP_ASSERT(status) assert(status == hipSuccess)
#define HIP_ASSERT(status) \
{ \
if ((status != hipSuccess)) { \
std::cerr << "Failed in: " << __LINE__ << " on hip call: " #status << std::endl; \
throw std::runtime_error("generic failure"); \
} \
}
#define LEN 512
__global__ void copy(uint32_t* A, uint32_t* B) {
size_t tid = threadIdx.x + blockIdx.x * blockDim.x;
B[tid] = A[tid];
size_t tid = threadIdx.x + blockIdx.x * blockDim.x;
B[tid] = A[tid];
}
void run_test1() {
uint32_t *A_h, *B_h, *A_d, *B_d;
size_t valbytes = LEN * sizeof(uint32_t);
uint32_t *A_h, *B_h, *A_d, *B_d;
size_t valbytes = LEN * sizeof(uint32_t);
A_h = (uint32_t*)malloc(valbytes);
B_h = (uint32_t*)malloc(valbytes);
for (uint32_t i = 0; i < LEN; i++) {
A_h[i] = i;
B_h[i] = 0;
}
A_h = (uint32_t*)malloc(valbytes);
B_h = (uint32_t*)malloc(valbytes);
for (uint32_t i = 0; i < LEN; i++) {
A_h[i] = i;
B_h[i] = 0;
}
HIP_ASSERT(hipMalloc((void**)&A_d, valbytes));
HIP_ASSERT(hipMalloc((void**)&B_d, valbytes));
HIP_ASSERT(hipMalloc((void**)&A_d, valbytes));
HIP_ASSERT(hipMalloc((void**)&B_d, valbytes));
HIP_ASSERT(hipMemcpy(A_d, A_h, valbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(copy, dim3(LEN/64), dim3(64), 0, 0, A_d, B_d);
HIP_ASSERT(hipMemcpy(B_h, B_d, valbytes, hipMemcpyDeviceToHost));
HIP_ASSERT(hipMemcpy(A_d, A_h, valbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(copy, dim3(LEN / 64), dim3(64), 0, 0, A_d, B_d);
HIP_ASSERT(hipMemcpy(B_h, B_d, valbytes, hipMemcpyDeviceToHost));
for (uint32_t i = 0; i < LEN; i++) {
assert(A_h[i] == B_h[i]);
}
for (uint32_t i = 0; i < LEN; i++) {
assert(A_h[i] == B_h[i]);
}
HIP_ASSERT(hipFree(A_d));
HIP_ASSERT(hipFree(B_d));
free(A_h);
free(B_h);
std::cout << "Test Passed!\n";
HIP_ASSERT(hipFree(A_d));
HIP_ASSERT(hipFree(B_d));
free(A_h);
free(B_h);
std::cout << "Test Passed!\n";
}