AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

SWDEV-240806 - Update graph test to use hipGraphGetNodes, hipGraphGetRootNodes and hipGraphAddDependencies

Telusuri Sumber 
Change-Id: I33249d2b625d8ff03c9fe4a71a541cac5fef9d24
This commit is contained in:
Anusha GodavarthySurya
2021-07-13 22:53:37 -07:00
committed by Anusha Godavarthy Surya
orang tua a0b301089c
melakukan 08aa662166
2 mengubah file dengan 248 tambahan dan 39 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
tests/src/runtimeApi/graph/hipGraph.cpp
+106 -39
Melihat File
@@ -21,40 +21,34 @@
#include <vector>
/* 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 <<<b, t, shmem>>>
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 <<<b, t, shmem>>>
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<double, std::milli>(stop - start);
std::cout << "Time taken for hipWithoutGraphs : "
<< std::chrono::duration_cast<std::chrono::milliseconds>(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<double, std::milli>(stop - start);
auto resultWithoutInit = std::chrono::duration<double, std::milli>(stop - start1);
std::cout << "Time taken for hipGraphsUsingStreamCapture with Init: "
<< std::chrono::duration_cast<std::chrono::milliseconds>(resultWithInit).count()
<< " milliseconds without Init:"
<< std::chrono::duration_cast<std::chrono::milliseconds>(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<double, std::milli>(stop - start);
auto resultWithoutInit = std::chrono::duration<double, std::milli>(stop - start1);
std::cout << "Time taken for hipGraphsManual with Init: "
<< std::chrono::duration_cast<std::chrono::milliseconds>(resultWithInit).count()
<< " milliseconds without Init:"
<< std::chrono::duration_cast<std::chrono::milliseconds>(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();
}
}
tests/src/runtimeApi/graph/hipSimpleGraphWithKernel.cpp
+142
Melihat File
@@ -0,0 +1,142 @@
#include <stdio.h>
#include <iostream>
#include "hip/hip_runtime.h"
#include <test_common.h>
#include <chrono>
#include <unistd.h>
/* 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<<<dim3(N / 512, 1, 1), dim3(512, 1, 1), 0, stream>>>(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<double, std::milli>(stop - start);
auto resultWithoutInit = std::chrono::duration<double, std::milli>(stop - start1);
std::cout << "Time taken for graph with Init: "
<< std::chrono::duration_cast<std::chrono::milliseconds>(resultWithInit).count()
<< " milliseconds without Init:"
<< std::chrono::duration_cast<std::chrono::milliseconds>(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<<<dim3(N / 512, 1, 1), dim3(512, 1, 1), 0, stream>>>(out_d, in_d);
}
HIPCHECK(hipStreamSynchronize(stream));
}
auto stop = std::chrono::high_resolution_clock::now();
auto result = std::chrono::duration<double, std::milli>(stop - start);
std::cout << "Time taken for test without graph: "
<< std::chrono::duration_cast<std::chrono::milliseconds>(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();
}