SWDEV-491363 - [catch2][dtest] Added test cases for missed graph APIs
Change-Id: Ic740c52992b8d0a325f09b1db0adf599561e1a1d
Этот коммит содержится в:
Anavena Venkatesh
@@ -39,8 +39,6 @@
|
||||
"Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters",
|
||||
"Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters",
|
||||
"Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic",
|
||||
"Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic",
|
||||
"Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Positive_Basic",
|
||||
"Unit_hipKernelNameRef_Negative_Parameters",
|
||||
"Unit_hipMemAdvise_No_Flag_Interference",
|
||||
"NOTE: The following 2 tests are disabled due to defect - EXSWHTEC-238",
|
||||
@@ -707,9 +705,6 @@
|
||||
"Unit_hipFreeAsync_Negative_Parameters",
|
||||
"Unit_hipIpcGetMemHandle_Positive_Unique_Handles_Reused_Memory",
|
||||
"SWDEV-445928: These tests fail in PSDB stress test on 09/02/2024",
|
||||
"Unit_hipGraphAddNodeTypeMemset_Positive_Basic - uint8_t",
|
||||
"Unit_hipGraphAddNodeTypeMemset_Positive_Basic - uint16_t",
|
||||
"Unit_hipGraphAddNodeTypeMemset_Positive_Basic - uint32_t",
|
||||
"Unit_hipCreateSurfaceObject_Negative_Parameters",
|
||||
"Unit_hipDestroySurfaceObject_Negative_Parameters",
|
||||
"Unit_Device___float2half_rd_Accuracy_Limited_Positive",
|
||||
@@ -757,6 +752,7 @@
|
||||
"=== SWDEV-432554:Below test failed in stress test on 10/11/23 ===",
|
||||
"Unit_hipMemcpy3DAsync_Positive_Basic",
|
||||
"Unit_hipDrvMemcpy3DAsync_Positive_Basic",
|
||||
"Print_Out_Attributes",
|
||||
"Unit_hipExtGetLinkTypeAndHopCount_Positive_Basic",
|
||||
"Unit_hipClock64_Positive_Basic",
|
||||
"Unit_hipClock_Positive_Basic",
|
||||
|
||||
@@ -10,6 +10,12 @@ set(TEST_SRC
|
||||
)
|
||||
add_executable(hipGetLastErrorEnv_Exe EXCLUDE_FROM_ALL hipGetLastErrorEnv_Exe.cc)
|
||||
add_executable(hipPeekAtLastErrorEnv_Exe EXCLUDE_FROM_ALL hipPeekAtLastErrorEnv_Exe.cc)
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set(AMD_SRC
|
||||
hipExtGetLastError.cc
|
||||
)
|
||||
set(TEST_SRC ${TEST_SRC} ${AMD_SRC})
|
||||
endif()
|
||||
hip_add_exe_to_target(NAME ErrorHandlingTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests
|
||||
|
||||
@@ -0,0 +1,862 @@
|
||||
/*
|
||||
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
|
||||
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/hip_runtime_api.h>
|
||||
#include <threaded_zig_zag_test.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
#include <chrono> //NOLINT
|
||||
#include <thread> //NOLINT
|
||||
|
||||
static constexpr size_t WIDTH = 1024;
|
||||
static constexpr size_t HEIGHT = 1024;
|
||||
static constexpr size_t N = 1024 * 1024;
|
||||
static constexpr size_t Nbytes = N * sizeof(int);
|
||||
|
||||
/**
|
||||
* @addtogroup hipExtGetLastError hipExtGetLastError
|
||||
* @{
|
||||
* @ingroup ErrorHandlingTest
|
||||
* `hipError_t hipExtGetLastError ( void )` -
|
||||
* Returns the last error from a runtime call.
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipMalloc api invalid arg call.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_Positive_Basic") {
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMalloc(nullptr, 1), hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with ThreadedZigZagTest api call.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_Positive_Threaded") {
|
||||
class HipGetLastErrorThreadedTest : public
|
||||
ThreadedZigZagTest<HipGetLastErrorThreadedTest> {
|
||||
public:
|
||||
void TestPart2() {
|
||||
REQUIRE_THREAD(hipMalloc(nullptr, 1) == hipErrorInvalidValue);
|
||||
}
|
||||
void TestPart3() {
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
void TestPart4() {
|
||||
REQUIRE_THREAD(hipExtGetLastError() == hipErrorInvalidValue);
|
||||
}
|
||||
};
|
||||
|
||||
HipGetLastErrorThreadedTest test;
|
||||
test.run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipMemcpyPeerAsync api invalid arg call
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpyPeerAsync") {
|
||||
const auto device_count = HipTest::getDeviceCount();
|
||||
if (device_count < 2) {
|
||||
HipTest::HIP_SKIP_TEST("Skipping because devices < 2");
|
||||
return;
|
||||
}
|
||||
|
||||
int can_access_peer = 0;
|
||||
const auto src_device = 0;
|
||||
const auto dst_device = 1;
|
||||
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
|
||||
if (can_access_peer) {
|
||||
HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0));
|
||||
|
||||
int *A_d, *B_d;
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
REQUIRE(A_d != nullptr);
|
||||
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
HIP_CHECK(hipMalloc(&B_d, Nbytes));
|
||||
REQUIRE(B_d != nullptr);
|
||||
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMemcpyPeerAsync(B_d, dst_device, A_d, src_device,
|
||||
Nbytes * 2, stream),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
|
||||
HIP_CHECK(hipDeviceDisablePeerAccess(dst_device));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
HIP_CHECK(hipFree(B_d));
|
||||
} else {
|
||||
INFO("Peer access cannot be enabled between devices " << src_device <<
|
||||
" and devices " << dst_device);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipMemcpyDtoHAsync api invalid arg call
|
||||
* Verify hipExtGetLastError status with hipMemcpyDtoDAsync api invalid arg call
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpyDtoHAsync") {
|
||||
int *A_d, *B_d, *A_h;
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
HipTest::initArrays<int>(&A_d, &B_d, nullptr,
|
||||
&A_h, nullptr, nullptr, N, false);
|
||||
|
||||
SECTION("Verify with hipMemcpyDtoHAsync api invalid arg call") {
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMemcpyDtoHAsync(A_h, (hipDeviceptr_t)A_d,
|
||||
Nbytes * 2, stream),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
SECTION("Verify with hipMemcpyDtoDAsync api invalid arg call") {
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMemcpyDtoDAsync((hipDeviceptr_t)A_d,
|
||||
(hipDeviceptr_t)B_d, Nbytes * 2, stream),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HipTest::freeArrays<int>(A_d, B_d, nullptr,
|
||||
A_h, nullptr, nullptr, false);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipMemcpyParam2DAsync api invalid arg
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpyParam2DAsync") {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
|
||||
float* A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, * A_d{nullptr};
|
||||
size_t pitch_A;
|
||||
size_t width{WIDTH * sizeof(float)};
|
||||
constexpr auto memsetval{100};
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
// Allocating and Initializing the data
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d),
|
||||
&pitch_A, width, HEIGHT));
|
||||
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &B_h, &C_h, width*HEIGHT, false);
|
||||
HipTest::setDefaultData<float>(WIDTH*HEIGHT, A_h, B_h, C_h);
|
||||
HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, WIDTH, HEIGHT));
|
||||
|
||||
// Device to Host
|
||||
hip_Memcpy2D desc = {};
|
||||
desc.srcMemoryType = hipMemoryTypeDevice;
|
||||
desc.dstMemoryType = hipMemoryTypeHost;
|
||||
desc.srcHost = A_d;
|
||||
desc.srcDevice = hipDeviceptr_t(A_d);
|
||||
desc.srcPitch = pitch_A;
|
||||
desc.dstHost = A_h;
|
||||
desc.dstDevice = hipDeviceptr_t(A_h);
|
||||
desc.dstPitch = width;
|
||||
desc.WidthInBytes = WIDTH;
|
||||
desc.Height = HEIGHT;
|
||||
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
desc.WidthInBytes = pitch_A+1;
|
||||
HIP_CHECK_ERROR(hipMemcpyParam2DAsync(&desc, stream),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
|
||||
// DeAllocating the memory
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
|
||||
A_h, B_h, C_h, false);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipDrvMemcpy3DAsync api invalid arg
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipDrvMemcpy3DAsync") {
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipDrvMemcpy3DAsync(nullptr, hipStreamPerThread),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipMemcpy3DAsync api invalid arg call
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpy3DAsync") {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
|
||||
constexpr int width{10}, height{10}, depth{10};
|
||||
auto size = width * height * depth * sizeof(int);
|
||||
hipArray_t devArray;
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
int *hData = reinterpret_cast<int*>(malloc(size));
|
||||
REQUIRE(hData != nullptr);
|
||||
memset(hData, 0, size);
|
||||
|
||||
// Initialize host buffer
|
||||
HipTest::setDefaultData<int>(width*height*depth, hData, nullptr, nullptr);
|
||||
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
|
||||
0, 0, 0, hipChannelFormatKindSigned);
|
||||
HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc,
|
||||
make_hipExtent(width, height, 2), hipArrayDefault));
|
||||
|
||||
hipMemcpy3DParms myparams;
|
||||
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 = devArray;
|
||||
myparams.kind = hipMemcpyHostToDevice;
|
||||
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMemcpy3DAsync(&myparams, stream),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
|
||||
// DeAllocating the memory
|
||||
HIP_CHECK(hipFreeArray(devArray));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
free(hData);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipMemcpy2DToArrayAsync api invalid arg
|
||||
* Verify hipExtGetLastError status with hipMemcpy2DFromArrayAsync api invalid arg
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpy2D_To_From_ArrayAsync") {
|
||||
int *hData = reinterpret_cast<int*>(malloc(WIDTH));
|
||||
REQUIRE(hData != nullptr);
|
||||
memset(hData, 0, WIDTH);
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
SECTION("Verify with hipMemcpyDtoHAsync api invalid arg call") {
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(nullptr, 0, 0, hData, WIDTH,
|
||||
WIDTH, HEIGHT, hipMemcpyHostToDevice, stream),
|
||||
hipErrorInvalidHandle);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidHandle);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
SECTION("Verify with hipMemcpyDtoHAsync api invalid arg call") {
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(hData, WIDTH, nullptr, 0, 0,
|
||||
WIDTH, HEIGHT, hipMemcpyDeviceToHost, stream),
|
||||
hipErrorInvalidHandle);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidHandle);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
free(hData);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipStreamAttachMemAsync api invalid arg
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipStreamAttachMemAsync") {
|
||||
void* d_memory{nullptr};
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMemPrefetchAsync(reinterpret_cast<void*>(d_memory),
|
||||
0, hipMemAttachHost, 0),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipWaitExternalSemaphoresAsync api invalid arg call
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipWaitExternalSemaphoresAsync") {
|
||||
hipExternalSemaphoreWaitParams wait_params = {};
|
||||
wait_params.params.fence.value = 1;
|
||||
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(nullptr,
|
||||
&wait_params, 1, nullptr),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipSignalExternalSemaphoresAsync api invalid arg call
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipSignalExternalSemaphoresAsync") {
|
||||
hipExternalSemaphoreSignalParams signal_params = {};
|
||||
signal_params.params.fence.value = 1;
|
||||
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(nullptr,
|
||||
&signal_params, 1, nullptr),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipMemPrefetchAsync api invalid arg call
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipMemPrefetchAsync") {
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMemPrefetchAsync(nullptr, 1024, 0),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipMemcpy2DAsync api invalid arg call
|
||||
* Verify hipExtGetLastError status with hipMemset2DAsync api invalid arg call
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpy2DAsync") {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
|
||||
int* A_h{nullptr}, *A_d{nullptr};
|
||||
size_t pitch_A;
|
||||
size_t width{WIDTH * sizeof(int)};
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
// Allocating memory
|
||||
A_h = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
REQUIRE(A_h != nullptr);
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d),
|
||||
&pitch_A, width, WIDTH));
|
||||
REQUIRE(A_d != nullptr);
|
||||
|
||||
// Initialize the data
|
||||
HipTest::setDefaultData<int>(WIDTH*HEIGHT, A_h, nullptr, nullptr);
|
||||
|
||||
SECTION("Verify with hipMemcpy2DAsync api invalid arg call") {
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMemcpy2DAsync(A_h, WIDTH*2, A_d, pitch_A,
|
||||
WIDTH*sizeof(int), WIDTH, hipMemcpyDeviceToHost, stream),
|
||||
hipErrorInvalidPitchValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidPitchValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
SECTION("Verify with hipMemset2DAsync api invalid arg call") {
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMemset2DAsync(A_d, pitch_A, 22,
|
||||
WIDTH*sizeof(int), WIDTH*9, stream),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
|
||||
// DeAllocating the memory
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipMemsetAsync api invalid arg call.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipMemsetAsync") {
|
||||
int *A_d;
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
REQUIRE(A_d != nullptr);
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipMemsetAsync(A_d, 0, Nbytes * 2, stream),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipMemcpyAsync api invalid arg call.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_MemCpyAsync") {
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
|
||||
HipTest::vectorADD<<<1, 1, 0, stream>>>(A_d, B_d, C_d, N);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
|
||||
// testing to check error manually
|
||||
HIP_CHECK_ERROR(hipMemcpyAsync(C_h, C_d, Nbytes+N, hipMemcpyDeviceToHost, 0),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipMemcpyAsync api invalid arg call
|
||||
* Check in other thread this error should not report by hipExtGetLastError()
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
// Inside thread, both hipExtGetLastError() api call should not return error
|
||||
static void thread_wait_func(int sleep_time) {
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
std::this_thread::sleep_for(std::chrono::milliseconds(sleep_time * 1000));
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_MemCpyAsync_thread") {
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
|
||||
HipTest::vectorADD<<<1, 1, 0, stream>>>(A_d, B_d, C_d, N);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
|
||||
std::thread t(thread_wait_func, 2);
|
||||
|
||||
// testing to check error manually
|
||||
HIP_CHECK_ERROR(hipMemcpyAsync(C_h, C_d, Nbytes+N, hipMemcpyDeviceToHost, 0),
|
||||
hipErrorInvalidValue);
|
||||
|
||||
t.join();
|
||||
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipGraphAddMemcpyNode1D api
|
||||
* Create graph with one node as error consciously so it produces an error,
|
||||
* which will be used to verify the behavior of hipExtGetLastError api.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipGraphAddMemcpyNode1D") {
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
size_t NElem{N};
|
||||
|
||||
hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, memcpy_E, kVecAdd;
|
||||
hipGraph_t graph;
|
||||
hipGraphExec_t graphExec;
|
||||
hipStream_t 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(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));
|
||||
|
||||
hipKernelNodeParams kNodeParams{};
|
||||
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(&kVecAdd, graph, nullptr, 0, &kNodeParams));
|
||||
|
||||
// hipGraphAddMemcpyNode1D is called conciously with double the size
|
||||
// so that it produces an error, which will be used to verify the
|
||||
// behavior of hipExtGetLastError api.
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipGraphAddMemcpyNode1D(&memcpy_E, graph, nullptr, 0, C_h,
|
||||
C_d, Nbytes * 2, hipMemcpyDeviceToHost),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
|
||||
// Create dependencies
|
||||
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kVecAdd, 1));
|
||||
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kVecAdd, 1));
|
||||
HIP_CHECK(hipGraphAddDependencies(graph, &kVecAdd, &memcpy_C, 1));
|
||||
|
||||
// Instantiate and launch the graph
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, stream));
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
|
||||
// Verify graph execution result
|
||||
HipTest::checkVectorADD(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(hipStreamDestroy(stream));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipStreamEndCapture api invalid arg call
|
||||
* Create a graph1 with stream with ketnelNode as vector_ADD and
|
||||
* hipStreamEndCapture on graph1 with hipGraphInstantiate to create graphExec
|
||||
* Again hipStreamEndCapture on graph2 which will return hipErrorIllegalState
|
||||
* now verify the behavior of hipExtGetLastError api with this call.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_hipStreamBegin_EndCapture") {
|
||||
int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h;
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
hipGraph_t graph1, graph2;
|
||||
hipGraphExec_t graphExec;
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
|
||||
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HipTest::vectorADD<int><<<1, 1, 0, stream>>>(A_d, B_d, C_d, N);
|
||||
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
HIP_CHECK(hipStreamEndCapture(stream, &graph1));
|
||||
|
||||
// EndCapture is called conciously so that it produces an error,
|
||||
// which will be used to verify the behavior of hipExtGetLastError api.
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipStreamEndCapture(stream, &graph2), hipErrorIllegalState);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorIllegalState);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph1, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, stream));
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
|
||||
// Verify graph execution result
|
||||
HipTest::checkVectorADD(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(graph1));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with hipGraphCreate api invalid arg call.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_error_check_with_hipGraphCreate") {
|
||||
hipGraph_t graph;
|
||||
hipError_t ret;
|
||||
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
ret = hipGraphCreate(&graph, 1);
|
||||
REQUIRE(ret == hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status should update with new api invalid arg call.
|
||||
* Api hipGraphCreate -> return error hipErrorInvalidValue
|
||||
* Api hipDeviceGetGraphMemAttribute -> return error hipErrorInvalidDevice
|
||||
* Now hipExtGetLastError() api shoud return hipErrorInvalidDevice
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_success_before_error_check_again") {
|
||||
int value = 0;
|
||||
hipGraph_t graph;
|
||||
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue);
|
||||
|
||||
HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute(-1,
|
||||
hipGraphMemAttrUsedMemCurrent, &value), hipErrorInvalidDevice);
|
||||
HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidDevice);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipExtGetLastError status with divide_by_zero kernel call
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/errorHandling/hipExtGetLastError.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
static void __global__ devideKernl(int *i, int x, int y) {
|
||||
*i = x/(x-y);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipExtGetLastError_with_Kernel_divide_by_zero") {
|
||||
int *i_d;
|
||||
int i = 9;
|
||||
HIP_CHECK(hipMalloc(&i_d, sizeof(int)));
|
||||
REQUIRE(i_d != nullptr);
|
||||
HIP_CHECK(hipMemcpy(i_d, &i, sizeof(int), hipMemcpyHostToDevice));
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
devideKernl<<<1, 1, 0, stream>>>(i_d, 3, 3);
|
||||
HIP_CHECK(hipExtGetLastError());
|
||||
|
||||
HIP_CHECK(hipFree(i_d));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
/**
|
||||
* End doxygen group hipExtGetLastError.
|
||||
* @}
|
||||
*/
|
||||
@@ -85,13 +85,19 @@ set(TEST_SRC
|
||||
hipGraphAsyncUserObj.cc
|
||||
hipGraphExecBatchMemOpNodeSetParams.cc
|
||||
hipGraphAddBatchMemOpNode.cc
|
||||
hipGraphBatchMemOpNodeSetParams.cc)
|
||||
hipGraphBatchMemOpNodeSetParams.cc
|
||||
hipDrvGraphExecMemcpyNodeSetParams.cc
|
||||
hipDrvGraphAddMemFreeNode.cc)
|
||||
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set(AMD_SRC
|
||||
hipStreamCaptureExtModuleLaunchKernel.cc
|
||||
hipStreamBeginCaptureToGraph.cc
|
||||
hipGetProcAddressGraphApis.cc
|
||||
# Below files are disbled in NVIDIA as PSDB builds are failing due to lower CUDA version.
|
||||
hipGraphExecNodeSetParams.cc
|
||||
hipGraphNodeSetParams.cc
|
||||
hipGraphExecGetFlags.cc
|
||||
)
|
||||
set(TEST_SRC ${TEST_SRC} ${AMD_SRC})
|
||||
endif()
|
||||
|
||||
@@ -0,0 +1,148 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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 <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup hipDrvGraphAddMemFreeNode hipDrvGraphAddMemFreeNode
|
||||
* @{
|
||||
* @ingroup GraphTest
|
||||
* `hipDrvGraphAddMemFreeNode ( hipGraphNode_t * phGraphNode,
|
||||
* hipGraph_t hGraph, const hipGraphNode_t * dependencies,
|
||||
* size_t numDependencies, hipDeviceptr_t dptr)'-
|
||||
* Creates a memory free node and adds it to a graph.
|
||||
*/
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test to verify hipDrvGraphAddMemFreeNode behavior with invalid arguments:
|
||||
* -# Null graph node
|
||||
* -# Null graph
|
||||
* -# Invalid numDependencies for null list of dependencies
|
||||
* -# Invalid numDependencies and valid list for dependencies
|
||||
* -# Null dev_ptr
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/graph/hipDrvGraphAddMemFreeNode.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
TEST_CASE("Unit_hipDrvGraphAddMemFreeNode_Negative_Params") {
|
||||
constexpr size_t N = 1024;
|
||||
hipGraph_t graph;
|
||||
hipGraphNode_t alloc_node, free_node;
|
||||
std::vector<hipGraphNode_t> dependencies;
|
||||
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
|
||||
hipMemAllocNodeParams alloc_param;
|
||||
memset(&alloc_param, 0, sizeof(alloc_param));
|
||||
alloc_param.bytesize = N;
|
||||
alloc_param.poolProps.allocType = hipMemAllocationTypePinned;
|
||||
alloc_param.poolProps.location.id = 0;
|
||||
alloc_param.poolProps.location.type = hipMemLocationTypeDevice;
|
||||
|
||||
HIP_CHECK(hipGraphAddMemAllocNode(&alloc_node, graph, nullptr,
|
||||
0, &alloc_param));
|
||||
REQUIRE(alloc_param.dptr != nullptr);
|
||||
|
||||
SECTION("Passing nullptr to graph node") {
|
||||
HIP_CHECK_ERROR(hipDrvGraphAddMemFreeNode(nullptr, graph,
|
||||
&alloc_node, 1, (hipDeviceptr_t)alloc_param.dptr),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("Passing nullptr to graph") {
|
||||
HIP_CHECK_ERROR(hipDrvGraphAddMemFreeNode(&free_node, nullptr,
|
||||
&alloc_node, 1, (hipDeviceptr_t)alloc_param.dptr),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("Pass invalid numDependencies") {
|
||||
HIP_CHECK_ERROR(hipDrvGraphAddMemFreeNode(&free_node, graph, nullptr,
|
||||
5, (hipDeviceptr_t)alloc_param.dptr),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("Pass invalid numDependencies and valid list for dependencies") {
|
||||
dependencies.push_back(alloc_node);
|
||||
HIP_CHECK_ERROR(hipDrvGraphAddMemFreeNode(&free_node, graph,
|
||||
dependencies.data(), dependencies.size() + 1,
|
||||
(hipDeviceptr_t)alloc_param.dptr),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("Passing nullptr to dev_ptr") {
|
||||
HIP_CHECK_ERROR(hipDrvGraphAddMemFreeNode(&alloc_node, graph,
|
||||
&alloc_node, 1, 0), hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
}
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - It will create memory alloation node and add to the graph then it
|
||||
* will create memory free node to free allocated memory and add to the graph.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/graph/hipDrvGraphAddMemFreeNode.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
TEST_CASE("Unit_hipDrvGraphAddMemFreeNode_Positive") {
|
||||
constexpr size_t N = 1024;
|
||||
hipGraph_t graph;
|
||||
hipGraphExec_t graphExec;
|
||||
hipCtx_t context;
|
||||
hipStream_t streamForGraph;
|
||||
int deviceid = 0;
|
||||
hipGraphNode_t node = nullptr, memFreeNode = nullptr;
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
HIP_CHECK(hipSetDevice(deviceid));
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph));
|
||||
HIP_CHECK(hipCtxCreate(&context, 0, deviceid));
|
||||
|
||||
hipMemAllocNodeParams alloc_param;
|
||||
memset(&alloc_param, 0, sizeof(alloc_param));
|
||||
alloc_param.bytesize = N;
|
||||
alloc_param.poolProps.allocType = hipMemAllocationTypePinned;
|
||||
alloc_param.poolProps.location.id = 0;
|
||||
alloc_param.poolProps.location.type = hipMemLocationTypeDevice;
|
||||
HIP_CHECK(hipGraphAddMemAllocNode(&node, graph, nullptr, 0, &alloc_param));
|
||||
REQUIRE(alloc_param.dptr != nullptr);
|
||||
|
||||
HIP_CHECK(hipDrvGraphAddMemFreeNode(&memFreeNode, graph, &node, 1,
|
||||
(hipDeviceptr_t)alloc_param.dptr));
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, 0));
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph));
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||
HIP_CHECK(hipCtxDestroy(context));
|
||||
}
|
||||
/**
|
||||
* End doxygen group GraphTest.
|
||||
* @}
|
||||
*/
|
||||
@@ -0,0 +1,212 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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 <numeric>
|
||||
|
||||
/**
|
||||
* @addtogroup hipDrvGraphExecMemcpyNodeSetParams hipDrvGraphExecMemcpyNodeSetParams
|
||||
* @{
|
||||
* @ingroup GraphTest
|
||||
* `hipDrvGraphExecMemcpyNodeSetParams((hipGraphExec_t hGraphExec,
|
||||
* hipGraphNode_t hNode, const HIP_MEMCPY3D *copyParams, hipCt_t ctx)` -
|
||||
* Sets the parameters for a memcpy node in the given graphExec
|
||||
*/
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify API behavior with invalid arguments:
|
||||
* -# graphExec is nullptr
|
||||
* -# node is nullptr
|
||||
*
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/graph/hipDrvGraphExecMemcpyNodeSetParams.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
TEST_CASE("Unit_hipDrvGraphExecMemcpyNodeSetParams_Negative") {
|
||||
size_t size = 10;
|
||||
size_t numW = size * sizeof(int);
|
||||
// Host Vectors
|
||||
std::vector<int> A_h(numW);
|
||||
std::vector<int> B_h(numW);
|
||||
hipCtx_t context;
|
||||
// Initialization
|
||||
std::iota(A_h.begin(), A_h.end(), 1);
|
||||
std::fill_n(B_h.begin(), size, 2);
|
||||
int deviceid = 0;
|
||||
hipGraph_t graph;
|
||||
hipStream_t streamForGraph;
|
||||
hipGraphExec_t graphExec;
|
||||
hipGraphNode_t node;
|
||||
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
HIP_CHECK(hipSetDevice(deviceid));
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph));
|
||||
HIP_CHECK(hipCtxCreate(&context, 0, deviceid));
|
||||
|
||||
|
||||
HIP_MEMCPY3D memCpy_params{};
|
||||
|
||||
memset(&memCpy_params, 0x0, sizeof(HIP_MEMCPY3D));
|
||||
memCpy_params.srcXInBytes = 0;
|
||||
memCpy_params.srcY = 0;
|
||||
memCpy_params.srcZ = 0;
|
||||
memCpy_params.dstXInBytes = 0;
|
||||
memCpy_params.dstY = 0;
|
||||
memCpy_params.dstZ = 0;
|
||||
memCpy_params.WidthInBytes = numW;
|
||||
memCpy_params.Height = 1;
|
||||
memCpy_params.Depth = 1;
|
||||
memCpy_params.srcMemoryType = hipMemoryTypeHost;
|
||||
memCpy_params.dstMemoryType = hipMemoryTypeHost;
|
||||
memCpy_params.srcHost = A_h.data();
|
||||
memCpy_params.srcPitch = numW;
|
||||
memCpy_params.srcHeight = 1;
|
||||
memCpy_params.dstHost = B_h.data();
|
||||
memCpy_params.dstPitch = numW;
|
||||
memCpy_params.dstHeight = 1;
|
||||
|
||||
HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr,
|
||||
0, &memCpy_params, context));
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph));
|
||||
SECTION("graphExec is nullptr") {
|
||||
HIP_CHECK_ERROR(hipDrvGraphExecMemcpyNodeSetParams(nullptr, node,
|
||||
&memCpy_params, context), hipErrorInvalidValue);
|
||||
}
|
||||
SECTION("node is nullptr") {
|
||||
HIP_CHECK_ERROR(hipDrvGraphExecMemcpyNodeSetParams(graphExec, nullptr,
|
||||
&memCpy_params, context), hipErrorInvalidValue);
|
||||
}
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||
HIP_CHECK(hipCtxDestroy(context));
|
||||
}
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - It will verify soure data is copied destionation
|
||||
* after adding new memcpy node.
|
||||
* First will create and add mem copy node to graph.
|
||||
* Create another mem copy param node with new source data and
|
||||
* add to the graphExec using hipDrvGraphExecMemcpyNodeSetParams API and
|
||||
* lauch the graph. Compare soure and destination data.
|
||||
*
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/graph/hipDrvGraphExecMemcpyNodeSetParams.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipDrvGraphExecMemcpyNodeSetParams_Positive") {
|
||||
size_t size = 10;
|
||||
size_t numW = size * sizeof(int);
|
||||
// Host Vectors
|
||||
std::vector<int> A_h(numW);
|
||||
std::vector<int> B_h(numW);
|
||||
hipCtx_t context;
|
||||
// Initialization
|
||||
std::iota(A_h.begin(), A_h.end(), 1);
|
||||
std::fill_n(B_h.begin(), size, 0);
|
||||
int deviceid = 0;
|
||||
hipGraph_t graph;
|
||||
hipStream_t streamForGraph;
|
||||
hipGraphExec_t graphExec;
|
||||
hipGraphNode_t node;
|
||||
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
HIP_CHECK(hipSetDevice(deviceid));
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph));
|
||||
HIP_CHECK(hipCtxCreate(&context, 0, deviceid));
|
||||
|
||||
|
||||
HIP_MEMCPY3D memCpy_params{};
|
||||
memset(&memCpy_params, 0x0, sizeof(HIP_MEMCPY3D));
|
||||
memCpy_params.srcXInBytes = 0;
|
||||
memCpy_params.srcY = 0;
|
||||
memCpy_params.srcZ = 0;
|
||||
memCpy_params.dstXInBytes = 0;
|
||||
memCpy_params.dstY = 0;
|
||||
memCpy_params.dstZ = 0;
|
||||
memCpy_params.WidthInBytes = numW;
|
||||
memCpy_params.Height = 1;
|
||||
memCpy_params.Depth = 1;
|
||||
memCpy_params.srcMemoryType = hipMemoryTypeHost;
|
||||
memCpy_params.dstMemoryType = hipMemoryTypeHost;
|
||||
memCpy_params.srcHost = A_h.data();
|
||||
memCpy_params.srcPitch = numW;
|
||||
memCpy_params.srcHeight = 1;
|
||||
memCpy_params.dstHost = B_h.data();
|
||||
memCpy_params.dstPitch = numW;
|
||||
memCpy_params.dstHeight = 1;
|
||||
|
||||
HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr,
|
||||
0, &memCpy_params, context));
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph));
|
||||
REQUIRE(memcmp(A_h.data(), B_h.data(), numW) == 0);
|
||||
// Host Vectors for source
|
||||
std::vector<int> C_h(numW);
|
||||
std::iota(C_h.begin(), C_h.end(), 10);
|
||||
|
||||
HIP_MEMCPY3D memCpy_params2{};
|
||||
memset(&memCpy_params2, 0x0, sizeof(HIP_MEMCPY3D));
|
||||
memCpy_params2.srcXInBytes = 0;
|
||||
memCpy_params2.srcY = 0;
|
||||
memCpy_params2.srcZ = 0;
|
||||
memCpy_params2.dstXInBytes = 0;
|
||||
memCpy_params2.dstY = 0;
|
||||
memCpy_params2.dstZ = 0;
|
||||
memCpy_params2.WidthInBytes = numW;
|
||||
memCpy_params2.Height = 1;
|
||||
memCpy_params2.Depth = 1;
|
||||
memCpy_params2.srcMemoryType = hipMemoryTypeHost;
|
||||
memCpy_params2.dstMemoryType = hipMemoryTypeHost;
|
||||
memCpy_params2.srcHost = C_h.data();
|
||||
memCpy_params2.srcPitch = numW;
|
||||
memCpy_params2.srcHeight = 1;
|
||||
memCpy_params2.dstHost = B_h.data();
|
||||
memCpy_params2.dstPitch = numW;
|
||||
memCpy_params2.dstHeight = 1;
|
||||
|
||||
HIP_CHECK(hipDrvGraphExecMemcpyNodeSetParams(graphExec, node,
|
||||
&memCpy_params2, context));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph));
|
||||
|
||||
REQUIRE(memcmp(C_h.data(), B_h.data(), numW) == 0);
|
||||
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||
HIP_CHECK(hipCtxDestroy(context));
|
||||
}
|
||||
/**
|
||||
* End doxygen group GraphTest.
|
||||
* @}
|
||||
*/
|
||||
|
||||
@@ -20,18 +20,17 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_defgroups.hh>
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_defgroups.hh>
|
||||
#include <memcpy3d_tests_common.hh>
|
||||
#include <numeric>
|
||||
|
||||
// hipDrvGraphAddMemcpyNode API is yet to be implemented in HIP runtime.
|
||||
#if 0
|
||||
/**
|
||||
* @addtogroup hipDrvGraphMemcpyNodeGetParams hipDrvGraphMemcpyNodeGetParams
|
||||
* @{
|
||||
* @ingroup GraphTest
|
||||
* `hipDrvGraphMemcpyNodeGetParams(hipGraphNode_t hNode, HIP_MEMCPY3D* nodeParams)` -
|
||||
* Gets a memcpy node's parameters
|
||||
* `hipDrvGraphMemcpyNodeGetParams(hipGraphNode_t hNode, HIP_MEMCPY3D*
|
||||
* nodeParams)` - Gets a memcpy node's parameters
|
||||
* ________________________
|
||||
* Test cases from other APIs:
|
||||
* - @ref Unit_hipDrvGraphMemcpyNodeSetParams_Positive_Basic
|
||||
@@ -63,37 +62,129 @@ TEST_CASE("Unit_hipDrvGraphMemcpyNodeGetParams_Negative_Parameters") {
|
||||
LinearAllocGuard3D<int> src_alloc(extent);
|
||||
LinearAllocGuard3D<int> dst_alloc(extent);
|
||||
|
||||
auto params =
|
||||
GetDrvMemcpy3DParms(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(),
|
||||
make_hipPos(0, 0, 0), dst_alloc.extent(), hipMemcpyDeviceToDevice);
|
||||
auto params = GetDrvMemcpy3DParms(
|
||||
dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(),
|
||||
make_hipPos(0, 0, 0), dst_alloc.extent(), hipMemcpyDeviceToDevice);
|
||||
|
||||
hipGraph_t graph = nullptr;
|
||||
hipGraphNode_t node = nullptr;
|
||||
|
||||
SECTION("node == nullptr") {
|
||||
HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(nullptr, ¶ms), hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(nullptr, ¶ms),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("pNodeParams == nullptr") {
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, ¶ms, context));
|
||||
HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(node, nullptr), hipErrorInvalidValue);
|
||||
HIP_CHECK(
|
||||
hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, ¶ms, context));
|
||||
HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(node, nullptr),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
}
|
||||
|
||||
#if HT_AMD
|
||||
SECTION("Node is destroyed") {
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, ¶ms, context));
|
||||
HIP_CHECK(
|
||||
hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, ¶ms, context));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(node, ¶ms), hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(node, ¶ms),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipCtxPopCurrent(&context));
|
||||
HIP_CHECK(hipCtxDestroy(context));
|
||||
}
|
||||
#endif // if 0
|
||||
|
||||
/**
|
||||
* End doxygen group GraphTest.
|
||||
* @}
|
||||
*/
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Create graph node with memcopy parameters and add to graph,
|
||||
* get the parameters from the created graph node and compare
|
||||
* the all the values.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/graph/hipDrvGraphMemcpyNodeGetParams.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipDrvGraphMemcpyNodeGetParams_Positive") {
|
||||
size_t size = 10;
|
||||
size_t numW = size * sizeof(int);
|
||||
// Host Vectors
|
||||
std::vector<int> A_h(numW);
|
||||
std::vector<int> B_h(numW);
|
||||
hipCtx_t context;
|
||||
// Initialization
|
||||
std::iota(A_h.begin(), A_h.end(), 1);
|
||||
std::fill_n(B_h.begin(), size, 0);
|
||||
int deviceid = 0;
|
||||
hipGraph_t graph;
|
||||
hipStream_t streamForGraph;
|
||||
hipGraphExec_t graphExec;
|
||||
hipGraphNode_t node;
|
||||
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
HIP_CHECK(hipSetDevice(deviceid));
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph));
|
||||
HIP_CHECK(hipCtxCreate(&context, 0, deviceid));
|
||||
|
||||
HIP_MEMCPY3D memCpy_params{};
|
||||
memset(&memCpy_params, 0x0, sizeof(HIP_MEMCPY3D));
|
||||
memCpy_params.srcXInBytes = 0;
|
||||
memCpy_params.srcY = 0;
|
||||
memCpy_params.srcZ = 0;
|
||||
memCpy_params.dstXInBytes = 0;
|
||||
memCpy_params.dstY = 0;
|
||||
memCpy_params.dstZ = 0;
|
||||
memCpy_params.WidthInBytes = numW;
|
||||
memCpy_params.Height = 1;
|
||||
memCpy_params.Depth = 1;
|
||||
memCpy_params.srcMemoryType = hipMemoryTypeHost;
|
||||
memCpy_params.dstMemoryType = hipMemoryTypeHost;
|
||||
memCpy_params.srcHost = A_h.data();
|
||||
memCpy_params.srcPitch = numW;
|
||||
memCpy_params.srcHeight = 1;
|
||||
memCpy_params.dstHost = B_h.data();
|
||||
memCpy_params.dstPitch = numW;
|
||||
memCpy_params.dstHeight = 1;
|
||||
|
||||
HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &memCpy_params,
|
||||
context));
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph));
|
||||
REQUIRE(memcmp(A_h.data(), B_h.data(), numW) == 0);
|
||||
|
||||
HIP_MEMCPY3D memCpyGetParams{};
|
||||
HIP_CHECK(hipDrvGraphMemcpyNodeGetParams(node, &memCpyGetParams));
|
||||
|
||||
REQUIRE(memCpy_params.srcXInBytes == memCpyGetParams.srcXInBytes);
|
||||
REQUIRE(memCpy_params.srcY == memCpyGetParams.srcY);
|
||||
REQUIRE(memCpy_params.srcZ == memCpyGetParams.srcZ);
|
||||
REQUIRE(memCpy_params.dstXInBytes == memCpyGetParams.dstXInBytes);
|
||||
REQUIRE(memCpy_params.dstY == memCpyGetParams.dstY);
|
||||
REQUIRE(memCpy_params.dstZ == memCpyGetParams.dstZ);
|
||||
REQUIRE(memCpy_params.WidthInBytes == memCpyGetParams.WidthInBytes);
|
||||
REQUIRE(memCpy_params.Height == memCpyGetParams.Height);
|
||||
REQUIRE(memCpy_params.Depth == memCpyGetParams.Depth);
|
||||
REQUIRE(memCpy_params.srcMemoryType == memCpyGetParams.srcMemoryType);
|
||||
REQUIRE(memCpy_params.dstMemoryType == memCpyGetParams.dstMemoryType);
|
||||
REQUIRE(memCpy_params.srcHost == memCpyGetParams.srcHost);
|
||||
REQUIRE(memCpy_params.srcPitch == memCpyGetParams.srcPitch);
|
||||
REQUIRE(memCpy_params.srcHeight == memCpyGetParams.srcHeight);
|
||||
REQUIRE(memCpy_params.dstHost == memCpyGetParams.dstHost);
|
||||
REQUIRE(memCpy_params.dstPitch == memCpyGetParams.dstPitch);
|
||||
REQUIRE(memCpy_params.dstHeight == memCpyGetParams.dstHeight);
|
||||
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||
HIP_CHECK(hipCtxPopCurrent(&context));
|
||||
HIP_CHECK(hipCtxDestroy(context));
|
||||
}
|
||||
/**
|
||||
* End doxygen group GraphTest.
|
||||
* @}
|
||||
*/
|
||||
|
||||
@@ -0,0 +1,149 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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>
|
||||
/**
|
||||
* @addtogroup hipGraphExecGetFlags hipGraphExecGetFlags
|
||||
* @{
|
||||
* @ingroup GraphTest
|
||||
* `hipGraphExecGetFlags(hipGraphExec_t graphExec,
|
||||
* unsigned long long *flags)` -
|
||||
* Return the flags on executable graph
|
||||
*/
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify API behavior with invalid arguments:
|
||||
* -# graphExec is nullptr
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/graph/hipGraphExecGetFlags.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphExecGetFlags_Negative") {
|
||||
hipGraphExec_t graphExec;
|
||||
unsigned long long flags; // NOLINT
|
||||
constexpr size_t Nbytes = 10 * sizeof(int);
|
||||
|
||||
hipGraphNode_t allocNodeA;
|
||||
hipMemAllocNodeParams allocParam;
|
||||
|
||||
hipGraph_t graph;
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
|
||||
memset(&allocParam, 0, sizeof(allocParam));
|
||||
allocParam.bytesize = Nbytes;
|
||||
allocParam.poolProps.allocType = hipMemAllocationTypePinned;
|
||||
allocParam.poolProps.location.id = 0;
|
||||
allocParam.poolProps.location.type = hipMemLocationTypeDevice;
|
||||
HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr,
|
||||
0, &allocParam));
|
||||
REQUIRE(allocParam.dptr != nullptr);
|
||||
|
||||
HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph,
|
||||
hipGraphInstantiateFlagAutoFreeOnLaunch));
|
||||
HIP_CHECK_ERROR(hipGraphExecGetFlags(nullptr, &flags),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - This test will verify the flags what we set while initiating
|
||||
* graph will be matching with flags values getting from
|
||||
* the hipGraphExecGetFlags API call.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/graph/hipGraphExecGetFlags.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphExecGetFlags_positive") {
|
||||
hipGraphExec_t graphExec;
|
||||
unsigned long long flags; // NOLINT
|
||||
hipGraph_t graph;
|
||||
constexpr size_t Nbytes = 10 * sizeof(int);
|
||||
|
||||
hipGraphNode_t allocNodeA;
|
||||
hipMemAllocNodeParams allocParam;
|
||||
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
|
||||
memset(&allocParam, 0, sizeof(allocParam));
|
||||
allocParam.bytesize = Nbytes;
|
||||
allocParam.poolProps.allocType = hipMemAllocationTypePinned;
|
||||
allocParam.poolProps.location.id = 0;
|
||||
allocParam.poolProps.location.type = hipMemLocationTypeDevice;
|
||||
|
||||
HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr,
|
||||
0, &allocParam));
|
||||
REQUIRE(allocParam.dptr != nullptr);
|
||||
|
||||
SECTION("flag is 0") {
|
||||
HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph,
|
||||
0));
|
||||
|
||||
HIP_CHECK(hipGraphExecGetFlags(graphExec, &flags));
|
||||
REQUIRE(flags == 0);
|
||||
}
|
||||
SECTION("flag is hipGraphInstantiateFlagAutoFreeOnLaunch") {
|
||||
HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph,
|
||||
hipGraphInstantiateFlagAutoFreeOnLaunch));
|
||||
|
||||
HIP_CHECK(hipGraphExecGetFlags(graphExec, &flags));
|
||||
REQUIRE(flags == hipGraphInstantiateFlagAutoFreeOnLaunch);
|
||||
}
|
||||
|
||||
// The below feature flags are not implemented
|
||||
// hipGraphInstantiateFlagUpload
|
||||
// hipGraphInstantiateFlagDeviceLaunch
|
||||
// hipGraphInstantiateFlagUseNodePriority
|
||||
#if 0
|
||||
SECTION("flag is hipGraphInstantiateFlagUpload") {
|
||||
HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph,
|
||||
hipGraphInstantiateFlagUpload));
|
||||
|
||||
HIP_CHECK(hipGraphExecGetFlags(graphExec, &flags));
|
||||
REQUIRE(flags == hipGraphInstantiateFlagUpload);
|
||||
}
|
||||
|
||||
SECTION("flag is hipGraphInstantiateFlagDeviceLaunch") {
|
||||
HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph,
|
||||
hipGraphInstantiateFlagDeviceLaunch));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, 0));
|
||||
REQUIRE(flags == hipGraphInstantiateFlagDeviceLaunch);
|
||||
}
|
||||
SECTION("flag is hipGraphInstantiateFlagUseNodePriority") {
|
||||
HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph,
|
||||
hipGraphInstantiateFlagUseNodePriority));
|
||||
HIP_CHECK(hipGraphExecGetFlags(graphExec, &flags));
|
||||
REQUIRE(flags == hipGraphInstantiateFlagUseNodePriority);
|
||||
}
|
||||
#endif
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
}
|
||||
/**
|
||||
* End doxygen group GraphTest.
|
||||
* @}
|
||||
*/
|
||||
|
||||
@@ -19,12 +19,10 @@ 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 <functional>
|
||||
#include <vector>
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
|
||||
#include "graph_memcpy_to_from_symbol_common.hh"
|
||||
|
||||
HIP_GRAPH_MEMCPY_FROM_SYMBOL_NODE_DEFINE_GLOBALS(char)
|
||||
@@ -38,24 +36,27 @@ HIP_GRAPH_MEMCPY_FROM_SYMBOL_NODE_DEFINE_ALTERNATE_GLOBALS(float)
|
||||
HIP_GRAPH_MEMCPY_FROM_SYMBOL_NODE_DEFINE_ALTERNATE_GLOBALS(double)
|
||||
|
||||
template <typename T>
|
||||
void GraphExecMemcpyToSymbolSetParamsShell(const void* symbol, const void* alt_symbol,
|
||||
size_t offset, const std::vector<T> set_values) {
|
||||
const auto f = [alt_symbol, is_arr = set_values.size() > 1](const void* symbol, void* src,
|
||||
size_t count, size_t offset,
|
||||
hipMemcpyKind direction) {
|
||||
void GraphExecMemcpyToSymbolSetParamsShell(const void *symbol,
|
||||
const void *alt_symbol,
|
||||
size_t offset,
|
||||
const std::vector<T> set_values) {
|
||||
const auto f = [alt_symbol, is_arr = set_values.size() > 1](
|
||||
const void *symbol, void *src, size_t count, size_t offset,
|
||||
hipMemcpyKind direction) {
|
||||
hipGraph_t graph = nullptr;
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
|
||||
hipGraphNode_t node = nullptr;
|
||||
HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(
|
||||
&node, graph, nullptr, 0, alt_symbol, reinterpret_cast<T*>(src) + is_arr,
|
||||
count - is_arr * sizeof(T), offset + is_arr * sizeof(T), direction));
|
||||
&node, graph, nullptr, 0, alt_symbol,
|
||||
reinterpret_cast<T *>(src) + is_arr, count - is_arr * sizeof(T),
|
||||
offset + is_arr * sizeof(T), direction));
|
||||
|
||||
hipGraphExec_t graph_exec = nullptr;
|
||||
HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0));
|
||||
|
||||
HIP_CHECK(hipGraphExecMemcpyNodeSetParamsToSymbol(graph_exec, node, symbol, src, count, offset,
|
||||
direction));
|
||||
HIP_CHECK(hipGraphExecMemcpyNodeSetParamsToSymbol(
|
||||
graph_exec, node, symbol, src, count, offset, direction));
|
||||
|
||||
HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread));
|
||||
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
|
||||
@@ -70,26 +71,30 @@ void GraphExecMemcpyToSymbolSetParamsShell(const void* symbol, const void* alt_s
|
||||
}
|
||||
|
||||
/**
|
||||
* @addtogroup hipGraphExecMemcpyNodeSetParamsToSymbol hipGraphExecMemcpyNodeSetParamsToSymbol
|
||||
* @addtogroup hipGraphExecMemcpyNodeSetParamsToSymbol
|
||||
* hipGraphExecMemcpyNodeSetParamsToSymbol
|
||||
* @{
|
||||
* @ingroup GraphTest
|
||||
* `hipGraphExecMemcpyNodeSetParamsToSymbol(hipGraphExec_t hGraphExec, hipGraphNode_t node,
|
||||
* const void *symbol, void *src, size_t count, size_t offset, hipMemcpyKind kind)` -
|
||||
* Sets the parameters for a memcpy node in the given graphExec to copy to a symbol on the device
|
||||
* `hipGraphExecMemcpyNodeSetParamsToSymbol(hipGraphExec_t hGraphExec,
|
||||
* hipGraphNode_t node, const void *symbol, void *src, size_t count, size_t
|
||||
* offset, hipMemcpyKind kind)` - Sets the parameters for a memcpy node in the
|
||||
* given graphExec to copy to a symbol on the device
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify that data is correctly copied to a symbol after node parameters are set following
|
||||
* node addition. A graph is constructed to which a MemcpyToSymbol node is added with valid but
|
||||
* incorrect parameters. After the graph is instantiated the parameters are updated to correct
|
||||
* values and the graph executed. After graph execution, a MemcpyFromSymbol is performed and the
|
||||
* copied values are compared against values known to have been copied to symbol memory previously.
|
||||
* The test is run for scalar, const scalar, array, and const array symbols of types char, int,
|
||||
* float and double. For array symbols, the test is repeated for zero and non-zero offset values.
|
||||
* Verification is performed for destination memory allocated on host and device.
|
||||
* Test source
|
||||
* - Verify that data is correctly copied to a symbol after node parameters
|
||||
* are set following node addition. A graph is constructed to which a
|
||||
* MemcpyToSymbol node is added with valid but incorrect parameters. After the
|
||||
* graph is instantiated the parameters are updated to correct values and the
|
||||
* graph executed. After graph execution, a MemcpyFromSymbol is performed and
|
||||
* the copied values are compared against values known to have been copied to
|
||||
* symbol memory previously. The test is run for scalar, const scalar, array,
|
||||
* and const array symbols of types char, int, float and double. For array
|
||||
* symbols, the test is repeated for zero and non-zero offset values.
|
||||
* Verification is performed for destination memory allocated on host and
|
||||
* device. Test source
|
||||
* ------------------------
|
||||
* - unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc
|
||||
* Test requirements
|
||||
@@ -98,23 +103,23 @@ void GraphExecMemcpyToSymbolSetParamsShell(const void* symbol, const void* alt_s
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic") {
|
||||
SECTION("char") {
|
||||
HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(GraphExecMemcpyToSymbolSetParamsShell, 10,
|
||||
char);
|
||||
HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(
|
||||
GraphExecMemcpyToSymbolSetParamsShell, 10, char);
|
||||
}
|
||||
|
||||
SECTION("int") {
|
||||
HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(GraphExecMemcpyToSymbolSetParamsShell, 10,
|
||||
int);
|
||||
HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(
|
||||
GraphExecMemcpyToSymbolSetParamsShell, 10, int);
|
||||
}
|
||||
|
||||
SECTION("float") {
|
||||
HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(GraphExecMemcpyToSymbolSetParamsShell, 10,
|
||||
float);
|
||||
HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(
|
||||
GraphExecMemcpyToSymbolSetParamsShell, 10, float);
|
||||
}
|
||||
|
||||
SECTION("double") {
|
||||
HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(GraphExecMemcpyToSymbolSetParamsShell, 10,
|
||||
double);
|
||||
HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(
|
||||
GraphExecMemcpyToSymbolSetParamsShell, 10, double);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -132,8 +137,8 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic") {
|
||||
* -# kind is illogical (hipMemcpyDeviceToHost)
|
||||
* -# kind is an invalid enum value
|
||||
* -# Changing memcpy direction
|
||||
* -# Changing src to memory allocated on a different device than the original src
|
||||
* Test source
|
||||
* -# Changing src to memory allocated on a different device than the
|
||||
* original src Test source
|
||||
* ------------------------
|
||||
* - unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc
|
||||
* Test requirements
|
||||
@@ -147,35 +152,37 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Negative_Parameters") {
|
||||
|
||||
LinearAllocGuard<int> var(LinearAllocs::hipMalloc, sizeof(int));
|
||||
hipGraphNode_t node = nullptr;
|
||||
HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&node, graph, nullptr, 0, SYMBOL(int_device_var),
|
||||
var.ptr(), sizeof(*var.ptr()), 0, hipMemcpyDefault));
|
||||
HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(
|
||||
&node, graph, nullptr, 0, SYMBOL(int_device_var), var.ptr(),
|
||||
sizeof(*var.ptr()), 0, hipMemcpyDefault));
|
||||
|
||||
hipGraphExec_t graph_exec = nullptr;
|
||||
HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0));
|
||||
|
||||
SECTION("hGraphExec == nullptr") {
|
||||
HIP_CHECK_ERROR(
|
||||
hipGraphExecMemcpyNodeSetParamsToSymbol(nullptr, node, SYMBOL(int_device_var), var.ptr(),
|
||||
sizeof(*var.ptr()), 0, hipMemcpyDefault),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsToSymbol(
|
||||
nullptr, node, SYMBOL(int_device_var), var.ptr(),
|
||||
sizeof(*var.ptr()), 0, hipMemcpyDefault),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("node == nullptr") {
|
||||
HIP_CHECK_ERROR(
|
||||
hipGraphExecMemcpyNodeSetParamsToSymbol(graph_exec, nullptr, SYMBOL(int_device_var),
|
||||
var.ptr(), sizeof(*var.ptr()), 0, hipMemcpyDefault),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsToSymbol(
|
||||
graph_exec, nullptr, SYMBOL(int_device_var), var.ptr(),
|
||||
sizeof(*var.ptr()), 0, hipMemcpyDefault),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
MemcpyToSymbolCommonNegative(
|
||||
std::bind(hipGraphExecMemcpyNodeSetParamsToSymbol, graph_exec, node, _1, _2, _3, _4, _5),
|
||||
std::bind(hipGraphExecMemcpyNodeSetParamsToSymbol, graph_exec, node, _1,
|
||||
_2, _3, _4, _5),
|
||||
SYMBOL(int_device_var), var.ptr(), sizeof(*var.ptr()));
|
||||
|
||||
SECTION("Changing memcpy direction") {
|
||||
HIP_CHECK_ERROR(
|
||||
hipGraphExecMemcpyNodeSetParamsToSymbol(graph_exec, node, SYMBOL(int_device_var), var.ptr(),
|
||||
sizeof(*var.ptr()), 0, hipMemcpyHostToDevice),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsToSymbol(
|
||||
graph_exec, node, SYMBOL(int_device_var), var.ptr(),
|
||||
sizeof(*var.ptr()), 0, hipMemcpyHostToDevice),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("Changing src allocation device") {
|
||||
@@ -185,9 +192,9 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Negative_Parameters") {
|
||||
HIP_CHECK(hipSetDevice(1));
|
||||
LinearAllocGuard<int> new_var(LinearAllocs::hipMalloc, sizeof(int));
|
||||
HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsToSymbol(
|
||||
graph_exec, node, SYMBOL(int_device_var), new_var.ptr(),
|
||||
sizeof(*new_var.ptr()), 0, static_cast<hipMemcpyKind>(-1)),
|
||||
hipErrorInvalidValue);
|
||||
graph_exec, node, SYMBOL(int_device_var),
|
||||
new_var.ptr(), sizeof(*new_var.ptr()),
|
||||
0, hipMemcpyDefault), hipErrorInvalidValue);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -196,6 +203,6 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Negative_Parameters") {
|
||||
}
|
||||
|
||||
/**
|
||||
* End doxygen group GraphTest.
|
||||
* @}
|
||||
*/
|
||||
* End doxygen group GraphTest.
|
||||
* @}
|
||||
*/
|
||||
|
||||
@@ -0,0 +1,156 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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>
|
||||
|
||||
/**
|
||||
* @addtogroup hipGraphExecNodeSetParams hipGraphExecNodeSetParams
|
||||
* @{
|
||||
* @ingroup GraphTest
|
||||
* `hipGraphExecNodeSetParams(hipGraphExec_t graphExec, hipGraphNode_t node,
|
||||
* hipGraphNodeParams *nodeParams)` -
|
||||
* Updates parameters of a created node on executable graph
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify API behavior with invalid arguments:
|
||||
* -# gGraphExec is nullptr
|
||||
* -# node is nullptr
|
||||
* -# nodeParams is nullptr
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/graph/hipGraphExecNodeSetParams.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphExecNodeSetParams_Negative_Parameters") {
|
||||
hipGraph_t graph;
|
||||
hipGraphExec_t graphExec;
|
||||
hipGraphNode_t node;
|
||||
hipGraphNodeParams node_params = {};
|
||||
char *A_d;
|
||||
size_t Nbytes = 10 * sizeof(char);
|
||||
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
node_params.type = hipGraphNodeTypeMemset;
|
||||
node_params.memset.dst = A_d;
|
||||
node_params.memset.elementSize = sizeof(char);
|
||||
node_params.memset.width = 10;
|
||||
node_params.memset.height = 1;
|
||||
node_params.memset.pitch = 10;
|
||||
node_params.memset.value = 99;
|
||||
|
||||
HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params));
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, 0));
|
||||
|
||||
SECTION("hGraphExec == nullptr") {
|
||||
HIP_CHECK_ERROR(
|
||||
hipGraphExecNodeSetParams(nullptr, node, &node_params),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("node == nullptr") {
|
||||
HIP_CHECK_ERROR(
|
||||
hipGraphExecNodeSetParams(graphExec, nullptr, &node_params),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("node params == nullptr") {
|
||||
HIP_CHECK_ERROR(hipGraphExecNodeSetParams(graphExec, node, nullptr),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
}
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - This will verify the new node param values are successfully
|
||||
* copied to graph node, after launching graphExec
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/graph/hipGraphExecNodeSetParams.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphExecNodeSetParams_Positive") {
|
||||
hipGraph_t graph;
|
||||
hipGraphExec_t graphExec;
|
||||
hipGraphNode_t node;
|
||||
hipGraphNodeParams node_params = {};
|
||||
char *A_d = nullptr, *A_h = nullptr;
|
||||
size_t Nbytes = 10 * sizeof(char);
|
||||
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
node_params.type = hipGraphNodeTypeMemset;
|
||||
node_params.memset.dst = A_d;
|
||||
node_params.memset.elementSize = sizeof(char);
|
||||
node_params.memset.width = 10;
|
||||
node_params.memset.height = 1;
|
||||
node_params.memset.pitch = 10;
|
||||
node_params.memset.value = 99;
|
||||
|
||||
HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params));
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, 0));
|
||||
HIP_CHECK(hipStreamSynchronize(0));
|
||||
A_h = reinterpret_cast<char*>(malloc(Nbytes));
|
||||
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
for (int i = 0; i < 10; i++) {
|
||||
REQUIRE(A_h[i] == 99);
|
||||
}
|
||||
|
||||
hipGraphNodeParams node_params2 = {};
|
||||
node_params2.type = hipGraphNodeTypeMemset;
|
||||
node_params2.memset.dst = A_d;
|
||||
node_params2.memset.elementSize = sizeof(char);
|
||||
node_params2.memset.width = 10;
|
||||
node_params2.memset.height = 1;
|
||||
node_params2.memset.pitch = 10;
|
||||
node_params2.memset.value = 110;
|
||||
|
||||
HIP_CHECK(hipGraphExecNodeSetParams(graphExec, node, &node_params2));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, 0));
|
||||
HIP_CHECK(hipStreamSynchronize(0));
|
||||
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
for (int i = 0; i < 10; i++) {
|
||||
REQUIRE(A_h[i] == 110);
|
||||
}
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
free(A_h);
|
||||
}
|
||||
/**
|
||||
* End doxygen group GraphTest.
|
||||
* @}
|
||||
*/
|
||||
@@ -0,0 +1,163 @@
|
||||
/*
|
||||
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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>
|
||||
|
||||
/**
|
||||
* @addtogroup hipGraphNodeSetParams hipGraphNodeSetParams
|
||||
* @{
|
||||
* @ingroup GraphTest
|
||||
* `hipGraphNodeSetParams(hipGraphNode_t node,
|
||||
* hipGraphNodeParams *nodeParams)` -
|
||||
* Updates parameters of a graph’s node
|
||||
*/
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify API behavior with invalid arguments:
|
||||
* -# node is nullptr
|
||||
* -# nodeParams is nullptr
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/graph/hipGraphNodeSetParams.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphNodeSetParams_Negative_Parameters") {
|
||||
hipGraph_t graph;
|
||||
hipGraphExec_t graphExec;
|
||||
hipGraphNode_t node;
|
||||
hipGraphNodeParams node_params = {};
|
||||
char *A_d;
|
||||
char *A_h;
|
||||
size_t N = 10;
|
||||
size_t Nbytes = N * sizeof(char);
|
||||
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
node_params.type = hipGraphNodeTypeMemset;
|
||||
node_params.memset.dst = A_d;
|
||||
node_params.memset.elementSize = sizeof(char);
|
||||
node_params.memset.width = N;
|
||||
node_params.memset.height = 1;
|
||||
node_params.memset.pitch = N;
|
||||
node_params.memset.value = 99;
|
||||
|
||||
HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params));
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, 0));
|
||||
A_h = reinterpret_cast<char *>(malloc(Nbytes));
|
||||
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
hipGraphNodeParams node_params2 = {};
|
||||
node_params2.type = hipGraphNodeTypeMemset;
|
||||
node_params2.memset.dst = A_d;
|
||||
node_params2.memset.elementSize = sizeof(char);
|
||||
node_params2.memset.width = N;
|
||||
node_params2.memset.height = 1;
|
||||
node_params2.memset.pitch = N;
|
||||
node_params2.memset.value = 110;
|
||||
|
||||
SECTION("node == nullptr") {
|
||||
HIP_CHECK_ERROR(hipGraphNodeSetParams(nullptr, &node_params2),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("nodeParams == nullptr") {
|
||||
HIP_CHECK_ERROR(hipGraphNodeSetParams(node, nullptr),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - This will verify the new node param values are successfully
|
||||
* copied to graph node
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/graph/hipGraphNodeSetParams.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.4
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphNodeSetParams_Positive") {
|
||||
hipGraph_t graph;
|
||||
hipGraphExec_t graphExec;
|
||||
hipGraphNode_t node;
|
||||
hipGraphNodeParams node_params = {};
|
||||
char *A_d, *A_h = nullptr;
|
||||
size_t N = 10;
|
||||
size_t Nbytes = N * sizeof(char);
|
||||
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
node_params.type = hipGraphNodeTypeMemset;
|
||||
node_params.memset.dst = A_d;
|
||||
node_params.memset.elementSize = sizeof(char);
|
||||
node_params.memset.width = N;
|
||||
node_params.memset.height = 1;
|
||||
node_params.memset.pitch = N;
|
||||
node_params.memset.value = 99;
|
||||
|
||||
HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params));
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, 0));
|
||||
HIP_CHECK(hipStreamSynchronize(0));
|
||||
A_h = reinterpret_cast<char*>(malloc(Nbytes));
|
||||
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
for (int i = 0; i < N; i++) {
|
||||
REQUIRE(A_h[i] == 99);
|
||||
}
|
||||
|
||||
hipGraphNodeParams node_params2 = {};
|
||||
node_params2.type = hipGraphNodeTypeMemset;
|
||||
node_params2.memset.dst = A_d;
|
||||
node_params2.memset.elementSize = sizeof(char);
|
||||
node_params2.memset.width = N;
|
||||
node_params2.memset.height = 1;
|
||||
node_params2.memset.pitch = N;
|
||||
node_params2.memset.value = 110;
|
||||
|
||||
HIP_CHECK(hipGraphNodeSetParams(node, &node_params2));
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, 0));
|
||||
HIP_CHECK(hipStreamSynchronize(0));
|
||||
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
for (int i = 0; i < N; i++) {
|
||||
REQUIRE(A_h[i] == 110);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* End doxygen group GraphTest.
|
||||
* @}
|
||||
*/
|
||||
Ссылка в новой задаче
Block a user