EXSWHTEC-169 - Implement additional tests for Kernel Graph Node APIs (#7)
- Tidy up hipGraphAddKernelNode tests - Tidy up hipGraphKernelNodeGetParams tests - Tidy up hipGraphKernelNodeSetParams tests - Tidy up hipGraphExecKernelNodeSetParams tests. - Disable failing test sections on AMD.
Этот коммит содержится в:
Mirza Halilčević
коммит произвёл
GitHub
GitHub
родитель
4dc52105c0
Коммит
e710eeea8d
@@ -6,19 +6,21 @@ 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
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY 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
|
||||
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 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
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
Test Case Scenarios :
|
||||
Negative -
|
||||
1) Pass node as nullptr and verify api returns error code.
|
||||
2) Pass pNodeParams as nullptr and verify api returns error code.
|
||||
@@ -30,13 +32,12 @@ Functional -
|
||||
hipGraphKernelNodeSetParams, finally check taking effect after launching graph.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
|
||||
/* Test verifies hipGraphKernelNodeSetParams API Negative scenarios.
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") {
|
||||
constexpr int N = 1024;
|
||||
size_t NElem{N};
|
||||
@@ -44,35 +45,53 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") {
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
int *A_d, *B_d, *C_d;
|
||||
hipError_t ret;
|
||||
hipGraph_t graph;
|
||||
hipGraphNode_t kNode;
|
||||
hipKernelNodeParams kNodeParams{};
|
||||
|
||||
HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N));
|
||||
HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N));
|
||||
HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N));
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
|
||||
|
||||
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
|
||||
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;
|
||||
|
||||
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
|
||||
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams));
|
||||
|
||||
SECTION("Pass node as nullptr") {
|
||||
ret = hipGraphKernelNodeSetParams(nullptr, &kNodeParams);
|
||||
REQUIRE(hipErrorInvalidValue == ret);
|
||||
HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(nullptr, &kNodeParams), hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("Pass kNodeParams as nullptr") {
|
||||
ret = hipGraphKernelNodeSetParams(kNode, nullptr);
|
||||
REQUIRE(hipErrorInvalidValue == ret);
|
||||
HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(kNode, nullptr), hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
#if HT_NVIDIA // on AMD this returns hipErrorInvalidValue
|
||||
SECTION("Pass NodeParams func data member as nullptr") {
|
||||
kNodeParams.func = nullptr;
|
||||
HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(kNode, &kNodeParams),
|
||||
hipErrorInvalidDeviceFunction);
|
||||
}
|
||||
#endif
|
||||
|
||||
#if HT_NVIDIA // segfaults on AMD
|
||||
SECTION("Pass kernelParams data member as nullptr") {
|
||||
kNodeParams.kernelParams = nullptr;
|
||||
HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(kNode, &kNodeParams), hipErrorInvalidValue);
|
||||
}
|
||||
#endif
|
||||
|
||||
#if HT_NVIDIA // segfaults on AMD
|
||||
SECTION("node is not a kernel node") {
|
||||
hipGraphNode_t empty_node;
|
||||
HIP_CHECK(hipGraphAddEmptyNode(&empty_node, graph, nullptr, 0));
|
||||
HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(empty_node, &kNodeParams), hipErrorInvalidValue);
|
||||
}
|
||||
#endif
|
||||
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
HIP_CHECK(hipFree(B_d));
|
||||
HIP_CHECK(hipFree(C_d));
|
||||
@@ -82,12 +101,12 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") {
|
||||
/**
|
||||
* Functional Test for API Set Kernel Params
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") {
|
||||
constexpr size_t N = 1024;
|
||||
constexpr size_t Nbytes = N * sizeof(int);
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
hipGraph_t graph;
|
||||
hipGraphNode_t memcpyNode, kNode;
|
||||
hipKernelNodeParams kNodeParams{}, kNodeParams1{};
|
||||
@@ -100,39 +119,34 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") {
|
||||
|
||||
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));
|
||||
|
||||
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));
|
||||
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>);
|
||||
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));
|
||||
HIP_CHECK(
|
||||
hipGraphAddKernelNode(&kNode, graph, dependencies.data(), dependencies.size(), &kNodeParams));
|
||||
|
||||
kNodeParams1.func = reinterpret_cast<void *>(HipTest::vectorSUB<int>);
|
||||
kNodeParams1.func = reinterpret_cast<void*>(HipTest::vectorSUB<int>);
|
||||
kNodeParams1.gridDim = dim3(blocks);
|
||||
kNodeParams1.blockDim = dim3(threadsPerBlock);
|
||||
kNodeParams1.sharedMemBytes = 0;
|
||||
kNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs);
|
||||
kNodeParams1.extra = nullptr;
|
||||
HIP_CHECK(hipGraphKernelNodeSetParams(kNode, &kNodeParams1));
|
||||
|
||||
dependencies.clear();
|
||||
dependencies.push_back(kNode);
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(),
|
||||
dependencies.size(), C_h, C_d,
|
||||
Nbytes, hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), dependencies.size(),
|
||||
C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
// Instantiate and launch the graph
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
|
||||
@@ -147,12 +161,12 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") {
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||
}
|
||||
|
||||
static __global__ void ker_vec_add(int *A, int *B) {
|
||||
static __global__ void ker_vec_add(int* A, int* B) {
|
||||
int i = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
A[i] = A[i] + B[i];
|
||||
}
|
||||
|
||||
static __global__ void ker_vec_sub(int *A, int *B) {
|
||||
static __global__ void ker_vec_sub(int* A, int* B) {
|
||||
int i = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
A[i] = A[i] - B[i];
|
||||
}
|
||||
@@ -167,7 +181,7 @@ class GraphKernelNodeGetSetParam {
|
||||
const int blocks = (N / threadsPerBlock);
|
||||
hipGraphNode_t memcpyH2D_A1, memcpyH2D_A2, memcpyD2H_A3, vec_maths;
|
||||
hipGraph_t graph;
|
||||
hipKernelNodeParams kerNodeParams { };
|
||||
hipKernelNodeParams kerNodeParams{};
|
||||
int *A1_d, *A2_d, *A1_h, *A2_h, *A3_h;
|
||||
|
||||
public:
|
||||
@@ -179,32 +193,26 @@ class GraphKernelNodeGetSetParam {
|
||||
HIP_CHECK(hipMalloc(&A2_d, Nbytes));
|
||||
// Allocate host buffers
|
||||
A1_h = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
REQUIRE(A1_h != NULL);
|
||||
REQUIRE(A1_h != nullptr);
|
||||
A2_h = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
REQUIRE(A2_h != NULL);
|
||||
REQUIRE(A2_h != nullptr);
|
||||
A3_h = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
REQUIRE(A3_h != NULL);
|
||||
REQUIRE(A3_h != nullptr);
|
||||
// Create all the 3 level graphs
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
void *kernelArgs[] = { &A1_d, &A2_d };
|
||||
void* kernelArgs[] = {&A1_d, &A2_d};
|
||||
kerNodeParams.func = reinterpret_cast<void*>(ker_vec_add);
|
||||
kerNodeParams.gridDim = dim3(blocks);
|
||||
kerNodeParams.blockDim = dim3(threadsPerBlock);
|
||||
kerNodeParams.sharedMemBytes = 0;
|
||||
kerNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
|
||||
kerNodeParams.extra = nullptr;
|
||||
HIP_CHECK(
|
||||
hipGraphAddKernelNode(&vec_maths, graph, nullptr, 0, &kerNodeParams));
|
||||
HIP_CHECK(hipGraphAddKernelNode(&vec_maths, graph, nullptr, 0, &kerNodeParams));
|
||||
// Add nodes to graph
|
||||
HIP_CHECK(
|
||||
hipGraphAddMemcpyNode1D(&memcpyH2D_A1, graph, nullptr, 0, A1_d, A1_h,
|
||||
Nbytes, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(
|
||||
hipGraphAddMemcpyNode1D(&memcpyH2D_A2, graph, nullptr, 0, A2_d, A2_h,
|
||||
Nbytes, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(
|
||||
hipGraphAddMemcpyNode1D(&memcpyD2H_A3, graph, nullptr, 0, A3_h, A1_d,
|
||||
Nbytes, hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A1, graph, nullptr, 0, A1_d, A1_h, Nbytes,
|
||||
hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A2, graph, nullptr, 0, A2_d, A2_h, Nbytes,
|
||||
hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_A3, graph, nullptr, 0, A3_h, A1_d, Nbytes,
|
||||
hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A1, &vec_maths, 1));
|
||||
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A2, &vec_maths, 1));
|
||||
HIP_CHECK(hipGraphAddDependencies(graph, &vec_maths, &memcpyD2H_A3, 1));
|
||||
@@ -213,20 +221,18 @@ class GraphKernelNodeGetSetParam {
|
||||
// Fill Random Input Data
|
||||
void fillRandInpData() {
|
||||
for (int i = 0; i < N; i++) {
|
||||
A1_h[i] = (rand() % 256); //NOLINT
|
||||
A2_h[i] = (rand() % 256); //NOLINT
|
||||
A1_h[i] = (rand() % 256); // NOLINT
|
||||
A2_h[i] = (rand() % 256); // NOLINT
|
||||
}
|
||||
}
|
||||
|
||||
hipGraph_t* getRootGraph() {
|
||||
return &graph;
|
||||
}
|
||||
hipGraph_t* getRootGraph() { return &graph; }
|
||||
|
||||
void updateNode() {
|
||||
size_t numNodes = 0;
|
||||
HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes));
|
||||
hipGraphNode_t *nodes = reinterpret_cast<hipGraphNode_t*>(malloc(
|
||||
numNodes * sizeof(hipGraphNode_t)));
|
||||
hipGraphNode_t* nodes =
|
||||
reinterpret_cast<hipGraphNode_t*>(malloc(numNodes * sizeof(hipGraphNode_t)));
|
||||
HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes));
|
||||
// Get the Graph node from the embedded graph
|
||||
size_t nodeIdx = 0;
|
||||
@@ -246,9 +252,7 @@ class GraphKernelNodeGetSetParam {
|
||||
}
|
||||
|
||||
// Function to validate result
|
||||
void validateOutData() {
|
||||
HipTest::checkVectorSUB<int>(A1_h, A2_h, A3_h, N);
|
||||
}
|
||||
void validateOutData() { HipTest::checkVectorSUB<int>(A1_h, A2_h, A3_h, N); }
|
||||
|
||||
// Destroy resources
|
||||
~GraphKernelNodeGetSetParam() {
|
||||
@@ -263,7 +267,7 @@ class GraphKernelNodeGetSetParam {
|
||||
};
|
||||
|
||||
TEST_CASE("Unit_hipGraphKernelNodeGetSetParams_Functional") {
|
||||
hipGraph_t *graph;
|
||||
hipGraph_t* graph;
|
||||
hipStream_t streamForGraph;
|
||||
hipGraphExec_t graphExec;
|
||||
GraphKernelNodeGetSetParam GraphKernelNodeGetSetParamObj;
|
||||
@@ -271,8 +275,7 @@ TEST_CASE("Unit_hipGraphKernelNodeGetSetParams_Functional") {
|
||||
GraphKernelNodeGetSetParamObj.updateNode();
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph));
|
||||
// Instantiate and launch the childgraph
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr,
|
||||
nullptr, 0));
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr, nullptr, 0));
|
||||
GraphKernelNodeGetSetParamObj.fillRandInpData();
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph));
|
||||
|
||||
Ссылка в новой задаче
Block a user