From 08aa662166323aa0cc9e205ad71f917f30a02b0e Mon Sep 17 00:00:00 2001 From: Anusha GodavarthySurya Date: Tue, 13 Jul 2021 22:53:37 -0700 Subject: [PATCH] SWDEV-240806 - Update graph test to use hipGraphGetNodes, hipGraphGetRootNodes and hipGraphAddDependencies Change-Id: I33249d2b625d8ff03c9fe4a71a541cac5fef9d24 --- tests/src/runtimeApi/graph/hipGraph.cpp | 145 +++++++++++++----- .../graph/hipSimpleGraphWithKernel.cpp | 142 +++++++++++++++++ 2 files changed, 248 insertions(+), 39 deletions(-) create mode 100644 tests/src/runtimeApi/graph/hipSimpleGraphWithKernel.cpp diff --git a/tests/src/runtimeApi/graph/hipGraph.cpp b/tests/src/runtimeApi/graph/hipGraph.cpp index c38e4f2f75..427e583181 100644 --- a/tests/src/runtimeApi/graph/hipGraph.cpp +++ b/tests/src/runtimeApi/graph/hipGraph.cpp @@ -21,40 +21,34 @@ #include /* HIT_START * BUILD: %t %s ../../test_common.cpp - * TEST: %t EXCLUDE_HIP_PLATFORM nvidia + * TEST: %t * HIT_END */ #define THREADS_PER_BLOCK 512 -#define GRAPH_LAUNCH_ITERATIONS 3 +#define GRAPH_LAUNCH_ITERATIONS 1000 __global__ void reduce(float* d_in, double* d_out, size_t inputSize, size_t outputSize) { - // sdata is allocated in the kernel call: 3rd arg to <<>> int myId = threadIdx.x + blockDim.x * blockIdx.x; int tid = threadIdx.x; - // do reduction in global mem for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { if (tid < s) { d_in[myId] += d_in[myId + s]; } - __syncthreads(); // make sure all adds at one stage are done! + __syncthreads(); } - // only thread 0 writes result for this block back to global mem if (tid == 0) { int blkx = blockIdx.x; d_out[blockIdx.x] = d_in[myId]; } } __global__ void reduceFinal(double* d_in, double* d_out, size_t inputSize) { - // sdata is allocated in the kernel call: 3rd arg to <<>> int myId = threadIdx.x + blockDim.x * blockIdx.x; int tid = threadIdx.x; - // do reduction in global mem for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { if (tid < s) { d_in[myId] += d_in[myId + s]; } - __syncthreads(); // make sure all adds at one stage are done! + __syncthreads(); } - // only thread 0 writes result for this block back to global mem if (tid == 0) { *d_out = d_in[myId]; } @@ -62,6 +56,55 @@ __global__ void reduceFinal(double* d_in, double* d_out, size_t inputSize) { void init_input(float* a, size_t size) { for (size_t i = 0; i < size; i++) a[i] = (rand() & 0xFF) / (float)RAND_MAX; } + +bool hipWithoutGraphs(float* inputVec_h, float* inputVec_d, double* outputVec_d, double* result_d, + size_t inputSize, size_t numOfBlocks) { + hipStream_t stream1, stream2, stream3; + hipEvent_t forkStreamEvent, memsetEvent1, memsetEvent2; + double result_h = 0.0; + HIPCHECK(hipStreamCreate(&stream1)); + HIPCHECK(hipStreamCreate(&stream2)); + HIPCHECK(hipStreamCreate(&stream3)); + HIPCHECK(hipEventCreate(&forkStreamEvent)); + HIPCHECK(hipEventCreate(&memsetEvent1)); + HIPCHECK(hipEventCreate(&memsetEvent2)); + auto start = std::chrono::high_resolution_clock::now(); + for (int i = 0; i < GRAPH_LAUNCH_ITERATIONS; i++) { + HIPCHECK(hipMemcpyAsync(inputVec_d, inputVec_h, sizeof(float) * inputSize, hipMemcpyDefault, + stream1)); + HIPCHECK(hipMemsetAsync(outputVec_d, 0, sizeof(double) * numOfBlocks, stream2)); + HIPCHECK(hipEventRecord(memsetEvent1, stream2)); + HIPCHECK(hipMemsetAsync(result_d, 0, sizeof(double), stream3)); + HIPCHECK(hipEventRecord(memsetEvent2, stream3)); + HIPCHECK(hipStreamWaitEvent(stream1, memsetEvent1, 0)); + hipLaunchKernelGGL(reduce, dim3(inputSize / THREADS_PER_BLOCK, 1, 1), + dim3(THREADS_PER_BLOCK, 1, 1), 0, stream1, inputVec_d, outputVec_d, + inputSize, numOfBlocks); + HIPCHECK(hipStreamWaitEvent(stream1, memsetEvent2, 0)); + hipLaunchKernelGGL(reduceFinal, dim3(1, 1, 1), dim3(THREADS_PER_BLOCK, 1, 1), 0, stream1, + outputVec_d, result_d, numOfBlocks); + HIPCHECK(hipMemcpyAsync(&result_h, result_d, sizeof(double), hipMemcpyDefault, stream1)); + HIPCHECK(hipStreamSynchronize(stream1)); + } + auto stop = std::chrono::high_resolution_clock::now(); + auto result = std::chrono::duration(stop - start); + std::cout << "Time taken for hipWithoutGraphs : " + << std::chrono::duration_cast(result).count() + << " millisecs " << std::endl; + HIPCHECK(hipStreamDestroy(stream1)); + HIPCHECK(hipStreamDestroy(stream2)); + HIPCHECK(hipStreamDestroy(stream3)); + double result_h_cpu = 0.0; + for (int i = 0; i < inputSize; i++) { + result_h_cpu += inputVec_h[i]; + } + if (result_h_cpu != result_h) { + printf("Final reduced sum = %lf %lf\n", result_h_cpu, result_h); + return false; + } + return true; +} + bool hipGraphsUsingStreamCapture(float* inputVec_h, float* inputVec_d, double* outputVec_d, double* result_d, size_t inputSize, size_t numOfBlocks) { hipStream_t stream1, stream2, stream3, streamForGraph; @@ -75,6 +118,7 @@ bool hipGraphsUsingStreamCapture(float* inputVec_h, float* inputVec_d, double* o HIPCHECK(hipEventCreate(&forkStreamEvent)); HIPCHECK(hipEventCreate(&memsetEvent1)); HIPCHECK(hipEventCreate(&memsetEvent2)); + auto start = std::chrono::high_resolution_clock::now(); HIPCHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); HIPCHECK(hipEventRecord(forkStreamEvent, stream1)); HIPCHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0)); @@ -94,16 +138,34 @@ bool hipGraphsUsingStreamCapture(float* inputVec_h, float* inputVec_d, double* o outputVec_d, result_d, numOfBlocks); HIPCHECK(hipMemcpyAsync(&result_h, result_d, sizeof(double), hipMemcpyDefault, stream1)); HIPCHECK(hipStreamEndCapture(stream1, &graph)); + hipGraphNode_t* nodes = NULL; + size_t numNodes = 0; + HIPCHECK(hipGraphGetNodes(graph, nodes, &numNodes)); + printf("\nNum of nodes in the graph created using stream capture API = %zu\n", numNodes); + HIPCHECK(hipGraphGetRootNodes(graph, nodes, &numNodes)); + printf("Num of root nodes in the graph created using stream capture API = %zu\n", numNodes); hipGraphExec_t graphExec; + HIPCHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + auto start1 = std::chrono::high_resolution_clock::now(); for (int i = 0; i < GRAPH_LAUNCH_ITERATIONS; i++) { HIPCHECK(hipGraphLaunch(graphExec, streamForGraph)); } HIPCHECK(hipStreamSynchronize(streamForGraph)); + auto stop = std::chrono::high_resolution_clock::now(); + auto resultWithInit = std::chrono::duration(stop - start); + auto resultWithoutInit = std::chrono::duration(stop - start1); + std::cout << "Time taken for hipGraphsUsingStreamCapture with Init: " + << std::chrono::duration_cast(resultWithInit).count() + << " milliseconds without Init:" + << std::chrono::duration_cast(resultWithoutInit).count() + << " milliseconds " << std::endl; + HIPCHECK(hipGraphExecDestroy(graphExec)); HIPCHECK(hipGraphDestroy(graph)); HIPCHECK(hipStreamDestroy(stream1)); HIPCHECK(hipStreamDestroy(stream2)); + HIPCHECK(hipStreamDestroy(stream3)); HIPCHECK(hipStreamDestroy(streamForGraph)); double result_h_cpu = 0.0; for (int i = 0; i < inputSize; i++) { @@ -123,25 +185,18 @@ bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d, hipGraphNode_t memcpyNode, kernelNode, memsetNode; double result_h = 0.0; HIPCHECK(hipStreamCreate(&streamForGraph)); + auto start = std::chrono::high_resolution_clock::now(); hipKernelNodeParams kernelNodeParams = {0}; - hipMemcpy3DParms memcpyParams = {0}; hipMemsetParams memsetParams = {0}; - memcpyParams.srcArray = NULL; - memcpyParams.srcPos = make_hipPos(0, 0, 0); - memcpyParams.srcPtr = make_hipPitchedPtr(inputVec_h, sizeof(float) * inputSize, inputSize, 1); - memcpyParams.dstArray = NULL; - memcpyParams.dstPos = make_hipPos(0, 0, 0); - memcpyParams.dstPtr = make_hipPitchedPtr(inputVec_d, sizeof(float) * inputSize, inputSize, 1); - memcpyParams.extent = make_hipExtent(sizeof(float) * inputSize, 1, 1); - memcpyParams.kind = hipMemcpyHostToDevice; memsetParams.dst = (void*)outputVec_d; memsetParams.value = 0; memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(float); // elementSize can be max 4 bytes + memsetParams.elementSize = sizeof(float); memsetParams.width = numOfBlocks * 2; memsetParams.height = 1; HIPCHECK(hipGraphCreate(&graph, 0)); - HIPCHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &memcpyParams)); + HIPCHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, inputVec_d, inputVec_h, + sizeof(float) * inputSize, hipMemcpyHostToDevice)); HIPCHECK(hipGraphAddMemsetNode(&memsetNode, graph, NULL, 0, &memsetParams)); nodeDependencies.push_back(memsetNode); nodeDependencies.push_back(memcpyNode); @@ -176,26 +231,33 @@ bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d, nodeDependencies.size(), &kernelNodeParams)); nodeDependencies.clear(); nodeDependencies.push_back(kernelNode); - memset(&memcpyParams, 0, sizeof(memcpyParams)); - memcpyParams.srcArray = NULL; - memcpyParams.srcPos = make_hipPos(0, 0, 0); - memcpyParams.srcPtr = make_hipPitchedPtr(result_d, sizeof(double), 1, 1); - memcpyParams.dstArray = NULL; - memcpyParams.dstPos = make_hipPos(0, 0, 0); - memcpyParams.dstPtr = make_hipPitchedPtr(&result_h, sizeof(double), 1, 1); - memcpyParams.extent = make_hipExtent(sizeof(double), 1, 1); - memcpyParams.kind = hipMemcpyDeviceToHost; - HIPCHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), - nodeDependencies.size(), &memcpyParams)); + HIPCHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nodeDependencies.data(), + nodeDependencies.size(), &result_h, result_d, sizeof(double), + hipMemcpyDeviceToHost)); nodeDependencies.clear(); nodeDependencies.push_back(memcpyNode); hipGraphNode_t hostNode; hipGraphExec_t graphExec; + hipGraphNode_t* nodes = NULL; + size_t numNodes = 0; + HIPCHECK(hipGraphGetNodes(graph, nodes, &numNodes)); + printf("\nNum of nodes in the graph created using hipGraphsManual API = %zu\n", numNodes); + HIPCHECK(hipGraphGetRootNodes(graph, nodes, &numNodes)); + printf("Num of root nodes in the graph created using hipGraphsManual API = %zu\n", numNodes); HIPCHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + auto start1 = std::chrono::high_resolution_clock::now(); for (int i = 0; i < GRAPH_LAUNCH_ITERATIONS; i++) { HIPCHECK(hipGraphLaunch(graphExec, streamForGraph)); } HIPCHECK(hipStreamSynchronize(streamForGraph)); + auto stop = std::chrono::high_resolution_clock::now(); + auto resultWithInit = std::chrono::duration(stop - start); + auto resultWithoutInit = std::chrono::duration(stop - start1); + std::cout << "Time taken for hipGraphsManual with Init: " + << std::chrono::duration_cast(resultWithInit).count() + << " milliseconds without Init:" + << std::chrono::duration_cast(resultWithoutInit).count() + << " milliseconds " << std::endl; HIPCHECK(hipGraphExecDestroy(graphExec)); HIPCHECK(hipGraphDestroy(graph)); @@ -211,9 +273,9 @@ bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d, return true; } int main(int argc, char** argv) { - size_t size = 1 << 12; // number of elements to reduce + size_t size = 1 << 12; size_t maxBlocks = 512; - hipSetDevice(0); // + hipSetDevice(0); printf("%zu elements\n", size); printf("threads per block = %d\n", THREADS_PER_BLOCK); printf("Graph Launch iterations = %d\n", GRAPH_LAUNCH_ITERATIONS); @@ -224,17 +286,22 @@ int main(int argc, char** argv) { HIPCHECK(hipMalloc(&outputVec_d, sizeof(double) * maxBlocks)); HIPCHECK(hipMalloc(&result_d, sizeof(double))); init_input(inputVec_h, size); - bool status1 = hipGraphsManual(inputVec_h, inputVec_d, outputVec_d, result_d, size, maxBlocks); - bool status2 = + bool status1 = hipWithoutGraphs(inputVec_h, inputVec_d, outputVec_d, result_d, size, maxBlocks); + bool status2 = hipGraphsManual(inputVec_h, inputVec_d, outputVec_d, result_d, size, maxBlocks); + bool status3 = hipGraphsUsingStreamCapture(inputVec_h, inputVec_d, outputVec_d, result_d, size, maxBlocks); + HIPCHECK(hipFree(inputVec_d)); HIPCHECK(hipFree(outputVec_d)); HIPCHECK(hipFree(result_d)); if (!status1) { - failed("Failed during hip Graph Manual\n"); + failed("Failed during hip without graph\n"); } if (!status2) { - failed("Failed during hip Graphs during stream capture\n"); + failed("Failed during hip graph manual\n"); + } + if (!status3) { + failed("Failed during hipGraph with capture\n"); } passed(); -} +} \ No newline at end of file diff --git a/tests/src/runtimeApi/graph/hipSimpleGraphWithKernel.cpp b/tests/src/runtimeApi/graph/hipSimpleGraphWithKernel.cpp new file mode 100644 index 0000000000..7baaab3f76 --- /dev/null +++ b/tests/src/runtimeApi/graph/hipSimpleGraphWithKernel.cpp @@ -0,0 +1,142 @@ +#include +#include +#include "hip/hip_runtime.h" +#include +#include +#include +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * TEST: %t + * HIT_END + */ +#define N 1024 * 1024 +#define NSTEP 1000 +#define NKERNEL 25 +#define CONSTANT 5.34 + +__global__ void simpleKernel(float* out_d, float* in_d) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < N) out_d[idx] = CONSTANT * in_d[idx]; +} + +bool hipTestWithGraph() { + int deviceId; + HIPCHECK(hipGetDevice(&deviceId)); + hipDeviceProp_t props; + HIPCHECK(hipGetDeviceProperties(&props, deviceId)); + + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + float *in_h, *out_h; + in_h = new float[N]; + out_h = new float[N]; + for (int i = 0; i < N; i++) { + in_h[i] = i; + } + + float *in_d, *out_d; + HIPCHECK(hipMalloc(&in_d, N * sizeof(float))); + HIPCHECK(hipMalloc(&out_d, N * sizeof(float))); + HIPCHECK(hipMemcpy(in_d, in_h, N * sizeof(float), hipMemcpyHostToDevice)); + + auto start = std::chrono::high_resolution_clock::now(); + // start CPU wallclock timer + bool graphCreated = false; + hipGraph_t graph; + hipGraphExec_t instance; + + hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal); + for (int ikrnl = 0; ikrnl < NKERNEL; ikrnl++) { + simpleKernel<<>>(out_d, in_d); + } + hipStreamEndCapture(stream, &graph); + hipGraphInstantiate(&instance, graph, NULL, NULL, 0); + + auto start1 = std::chrono::high_resolution_clock::now(); + for (int istep = 0; istep < NSTEP; istep++) { + hipGraphLaunch(instance, stream); + hipStreamSynchronize(stream); + } + auto stop = std::chrono::high_resolution_clock::now(); + auto resultWithInit = std::chrono::duration(stop - start); + auto resultWithoutInit = std::chrono::duration(stop - start1); + std::cout << "Time taken for graph with Init: " + << std::chrono::duration_cast(resultWithInit).count() + << " milliseconds without Init:" + << std::chrono::duration_cast(resultWithoutInit).count() + << " milliseconds " << std::endl; + + HIPCHECK(hipMemcpy(out_h, out_d, N * sizeof(float), hipMemcpyDeviceToHost)); + for (int i = 0; i < N; i++) { + if (float(in_h[i] * CONSTANT) != out_h[i]) { + return false; + } + } + delete[] in_h; + delete[] out_h; + HIPCHECK(hipFree(in_d)); + HIPCHECK(hipFree(out_d)); + return true; +} + +bool hipTestWithoutGraph() { + int deviceId; + HIPCHECK(hipGetDevice(&deviceId)); + hipDeviceProp_t props; + HIPCHECK(hipGetDeviceProperties(&props, deviceId)); + printf("info: running on device #%d %s\n", deviceId, props.name); + + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + float *in_h, *out_h; + in_h = new float[N]; + out_h = new float[N]; + for (int i = 0; i < N; i++) { + in_h[i] = i; + } + + float *in_d, *out_d; + HIPCHECK(hipMalloc(&in_d, N * sizeof(float))); + HIPCHECK(hipMalloc(&out_d, N * sizeof(float))); + HIPCHECK(hipMemcpy(in_d, in_h, N * sizeof(float), hipMemcpyHostToDevice)); + + // start CPU wallclock timer + auto start = std::chrono::high_resolution_clock::now(); + for (int istep = 0; istep < NSTEP; istep++) { + for (int ikrnl = 0; ikrnl < NKERNEL; ikrnl++) { + simpleKernel<<>>(out_d, in_d); + } + HIPCHECK(hipStreamSynchronize(stream)); + } + auto stop = std::chrono::high_resolution_clock::now(); + auto result = std::chrono::duration(stop - start); + std::cout << "Time taken for test without graph: " + << std::chrono::duration_cast(result).count() + << " millisecs " << std::endl; + HIPCHECK(hipMemcpy(out_h, out_d, N * sizeof(float), hipMemcpyDeviceToHost)); + for (int i = 0; i < N; i++) { + if (float(in_h[i] * CONSTANT) != out_h[i]) { + return false; + } + } + delete[] in_h; + delete[] out_h; + HIPCHECK(hipFree(in_d)); + HIPCHECK(hipFree(out_d)); + return true; +} + +int main(int argc, char* argv[]) { + bool status1, status2; + status1 = hipTestWithoutGraph(); + status2 = hipTestWithGraph(); + if (!status1) { + failed("Failed during test with hip graph\n"); + } + if (!status2) { + failed("Failed during test without graph\n"); + } + passed(); +}