From 12fb9350ad07a62b1906b67b08fcc9bb25c98c0f Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 2 Mar 2023 07:01:07 +0530 Subject: [PATCH 1/2] SWDEV-360327 - [catch2][dtest] Test cases added for hipLaunchHostFunc API (#170) Change-Id: I0d92ed0d8e97a2f7b34cc89dd6dcffa2236ee308 --- .../config/config_amd_windows_MI2xx.json | 5 +- .../config/config_amd_windows_common.json | 10 +- catch/unit/stream/CMakeLists.txt | 1 + catch/unit/stream/hipLaunchHostFunc.cc | 556 ++++++++++++++++++ 4 files changed, 570 insertions(+), 2 deletions(-) create mode 100644 catch/unit/stream/hipLaunchHostFunc.cc diff --git a/catch/hipTestMain/config/config_amd_windows_MI2xx.json b/catch/hipTestMain/config/config_amd_windows_MI2xx.json index 7898310ea5..3121abb125 100644 --- a/catch/hipTestMain/config/config_amd_windows_MI2xx.json +++ b/catch/hipTestMain/config/config_amd_windows_MI2xx.json @@ -93,6 +93,9 @@ "Unit_hipInit_Negative", "Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime", "Unit_hipStreamBeginCapture_captureComplexGraph", - "Unit_hipMemGetAddressRange_Negative" + "Unit_hipMemGetAddressRange_Negative", + "Unit_hipStreamValue_Wait64_Blocking_NoMask_Nor", + "Unit_hipLaunchHostFunc_Graph", + "Unit_hipLaunchHostFunc_KernelHost" ] } diff --git a/catch/hipTestMain/config/config_amd_windows_common.json b/catch/hipTestMain/config/config_amd_windows_common.json index 0e87e64b1f..73265230ea 100644 --- a/catch/hipTestMain/config/config_amd_windows_common.json +++ b/catch/hipTestMain/config/config_amd_windows_common.json @@ -56,6 +56,7 @@ "Unit_hipStreamWaitEvent_DifferentStreams", "Unit_hipStreamQuery_WithFinishedWork", "Unit_hipDeviceGetCacheConfig_Positive_Basic", + "Unit_hipDeviceGetCacheConfig_Positive_Basic", "Unit_hipDeviceGetCacheConfig_Positive_Threaded", "Unit_hipStreamValue_Wait32_Blocking_Mask_Gte", "Unit_hipStreamValue_Wait32_Blocking_Mask_Eq_1", @@ -96,10 +97,17 @@ "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", "Unit_hipMemAdvise_AccessedBy_All_Devices", "Unit_hipMemAdvise_No_Flag_Interference", + "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode", "Unit_hipMemGetAddressRange_Negative", - "Unit_hipMemGetAddressRange_Positive" + "Unit_hipMemGetAddressRange_Positive", + "Unit_hipGraphAddEventRecordNode_MultipleRun", + "Unit_hipDeviceGetPCIBusId_Negative_PartialFill", + "Unit_hipStreamValue_Wait64_Blocking_NoMask_Nor", + "Unit_hipStreamQuery_WithFinishedWork", + "Unit_hipLaunchHostFunc_Graph", + "Unit_hipLaunchHostFunc_KernelHost" ] } diff --git a/catch/unit/stream/CMakeLists.txt b/catch/unit/stream/CMakeLists.txt index b2228290fe..a82fa769c9 100644 --- a/catch/unit/stream/CMakeLists.txt +++ b/catch/unit/stream/CMakeLists.txt @@ -15,6 +15,7 @@ set(TEST_SRC hipStreamQuery.cc hipDeviceGetStreamPriorityRange.cc hipStreamACb_StrmSyncTiming.cc + hipLaunchHostFunc.cc ) if(HIP_PLATFORM MATCHES "amd") diff --git a/catch/unit/stream/hipLaunchHostFunc.cc b/catch/unit/stream/hipLaunchHostFunc.cc new file mode 100644 index 0000000000..a3f49387ba --- /dev/null +++ b/catch/unit/stream/hipLaunchHostFunc.cc @@ -0,0 +1,556 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include +#define GRIDSIZE 512 +#define BLOCKSIZE 256 +#define NUM_OF_STREAM 3 +#define THREADS_PER_BLOCK 512 +#define GRAPH_LAUNCH_ITERATIONS 1000 + +static __global__ void reduce(float* d_in, double* d_out) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + int tid = threadIdx.x; + for (int s = blockDim.x / 2; s > 0; s >>= 1) { + if (tid < s) { + d_in[myId] += d_in[myId + s]; + } + __syncthreads(); + } + if (tid == 0) { + d_out[blockIdx.x] = d_in[myId]; + } +} +static __global__ void reduceFinal(double* d_in, double* d_out) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + int tid = threadIdx.x; + for (int s = blockDim.x / 2; s > 0; s >>= 1) { + if (tid < s) { + d_in[myId] += d_in[myId + s]; + } + __syncthreads(); + } + if (tid == 0) { + *d_out = d_in[myId]; + } +} + +static void init_input(float* a, size_t size) { + unsigned int seed = time(nullptr); + for (size_t i = 0; i < size; i++) { + a[i] = (HipTest::RAND_R(&seed) & 0xFF) / static_cast(RAND_MAX); + } +} + +static bool gPassed = true; +static void *gusrptr; +static void *ptr0xff = reinterpret_cast(0xffffffff); +static size_t NSize = GRIDSIZE * BLOCKSIZE; +static size_t Nbytes = NSize * sizeof(float); + +typedef struct userDataStruct { + float *A_h; + float *C_h; + float *A_d; + float *C_d; + bool isPassed; + bool isOpCompleted; +} usrDataS; + +// Common callback function. +static void Fn_validateSq(void* userData) { + REQUIRE(userData != nullptr); + usrDataS *ptrUsrData = reinterpret_cast(userData); + for (size_t i = 0; i < NSize; i++) { + if (ptrUsrData->C_h[i] != + (ptrUsrData->A_h[i] * ptrUsrData->A_h[i])) { + ptrUsrData->isPassed = false; + return; + } + } + ptrUsrData->isPassed = true; +} + +// Test scenario 1 +// simple scenario that validates passing userData to host function. +static void Fn_ChkUserdataPtr(void* userData) { + gPassed = true; + if (gusrptr != userData) { + gPassed = false; + } +} + +TEST_CASE("Unit_hipLaunchHostFunc_basic") { + hipStream_t mystream; + HIP_CHECK(hipStreamCreate(&mystream)); + gusrptr = ptr0xff; + gPassed = true; + HIP_CHECK(hipLaunchHostFunc(mystream, Fn_ChkUserdataPtr, gusrptr)); + HIP_CHECK(hipStreamSynchronize(mystream)); + HIP_CHECK(hipStreamDestroy(mystream)); + REQUIRE(gPassed); +} + +// Negative test scenario for hipLaunchHostFunc +TEST_CASE("Unit_hipLaunchHostFunc_Negative") { + hipStream_t mystream; + HIP_CHECK(hipStreamCreate(&mystream)); + + SECTION("Pass nullptr as function") { + REQUIRE(hipLaunchHostFunc(mystream, nullptr, 0) == hipErrorInvalidValue); + } + HIP_CHECK(hipStreamDestroy(mystream)); +} + +// Local Function +static void launchOperationOnStrm(usrDataS *usrDataptr, hipStream_t stream) { + usrDataptr->isPassed = false; + HIP_CHECK(hipMallocAsync(reinterpret_cast(&(usrDataptr->A_d)), + Nbytes, stream)); + HIP_CHECK(hipMallocAsync(reinterpret_cast(&(usrDataptr->C_d)), + Nbytes, stream)); + HIP_CHECK(hipMemcpyAsync(usrDataptr->A_d, usrDataptr->A_h, Nbytes, + hipMemcpyHostToDevice, stream)); + hipLaunchKernelGGL((HipTest::vector_square), dim3(GRIDSIZE), + dim3(BLOCKSIZE), 0, stream, usrDataptr->A_d, + usrDataptr->C_d, NSize); + HIP_CHECK(hipMemcpyAsync(usrDataptr->C_h, usrDataptr->C_d, Nbytes, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipLaunchHostFunc(stream, Fn_validateSq, + reinterpret_cast(usrDataptr))); + HIP_CHECK(hipFreeAsync(reinterpret_cast(usrDataptr->A_d), + stream)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(usrDataptr->C_d), + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(usrDataptr->isPassed); +} + +// Test scenario 2 +// scenario that validates the host launch function on 3 different streams, +// created stream, default/null stream and hipStreamPerThread. +TEST_CASE("Unit_hipLaunchHostFunc_streams") { + hipStream_t stream[NUM_OF_STREAM]; + HIP_CHECK(hipStreamCreate(&stream[0])); + stream[1] = 0; // Null stream + stream[2] = hipStreamPerThread; + usrDataS *usrDataptr = reinterpret_cast( + malloc(sizeof(usrDataS))); + REQUIRE(usrDataptr != nullptr); + usrDataptr->A_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr->A_h != nullptr); + usrDataptr->C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr->C_h != nullptr); + for (size_t i = 0; i < NSize; i++) { + usrDataptr->A_h[i] = 21.0f; + } + for (int idx = 0; idx < NUM_OF_STREAM; idx++) { + launchOperationOnStrm(usrDataptr, stream[idx]); + } + HIP_CHECK(hipStreamDestroy(stream[0])); + free(usrDataptr->A_h); + free(usrDataptr->C_h); + free(usrDataptr); +} + +// Test scenario 3 +// test case to validate hipLaunchHostFunc for multi stream scenario. +// create 2 different streams and call hipLaunchHostFunc, stream synchronize. +static void Fn_validateMul_stream(void* userData) { + REQUIRE(userData != nullptr); + usrDataS *ptrUsrData = reinterpret_cast(userData); + for (size_t i = 0; i < NSize; i++) { + if (ptrUsrData->C_h[i] != + (ptrUsrData->A_h[i] * ptrUsrData->A_h[i])) { + ptrUsrData->isPassed = false; + return; + } + } + ptrUsrData->isPassed = true; +} + +TEST_CASE("Unit_hipLaunchHostFunc_multistreams") { + hipStream_t mystream1, mystream2; + HIP_CHECK(hipStreamCreateWithFlags(&mystream1, hipStreamNonBlocking)); + HIP_CHECK(hipStreamCreateWithFlags(&mystream2, hipStreamNonBlocking)); + usrDataS *usrDataptr1 = reinterpret_cast( + malloc(sizeof(usrDataS))); + REQUIRE(usrDataptr1 != nullptr); + usrDataS *usrDataptr2 = reinterpret_cast( + malloc(sizeof(usrDataS))); + REQUIRE(usrDataptr2 != nullptr); + usrDataptr1->A_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr1->A_h != nullptr); + usrDataptr1->C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr1->C_h != nullptr); + // input data + for (size_t i = 0; i < NSize; i++) { + usrDataptr1->A_h[i] = 11.0f; + } + usrDataptr1->isPassed = false; + usrDataptr2->isPassed = false; + HIP_CHECK(hipMallocAsync(reinterpret_cast(&(usrDataptr1->A_d)), + Nbytes, mystream1)); + HIP_CHECK(hipMallocAsync(reinterpret_cast(&(usrDataptr1->C_d)), + Nbytes, mystream1)); + HIP_CHECK(hipMemcpyAsync(usrDataptr1->A_d, usrDataptr1->A_h, Nbytes, + hipMemcpyHostToDevice, mystream1)); + const unsigned blocks = GRIDSIZE; + const unsigned threadsPerBlock = BLOCKSIZE; + hipLaunchKernelGGL((HipTest::vector_square), dim3(blocks), + dim3(threadsPerBlock), 0, mystream1, usrDataptr1->A_d, + usrDataptr1->C_d, NSize); + HIP_CHECK(hipMemcpyAsync(usrDataptr1->C_h, usrDataptr1->C_d, Nbytes, + hipMemcpyDeviceToHost, mystream1)); + HIP_CHECK(hipLaunchHostFunc(mystream1, Fn_validateMul_stream, + reinterpret_cast(usrDataptr1))); + // launch kernel function for mystream2 + usrDataptr2->A_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr2->A_h != nullptr); + usrDataptr2->C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr2->C_h != nullptr); + // input data + for (size_t i = 0; i < NSize; i++) { + usrDataptr2->A_h[i] = 9.0f; + } + HIP_CHECK(hipMallocAsync(reinterpret_cast(&(usrDataptr2->A_d)), + Nbytes, mystream2)); + HIP_CHECK(hipMallocAsync(reinterpret_cast(&(usrDataptr2->C_d)), + Nbytes, mystream2)); + HIP_CHECK(hipMemcpyAsync(usrDataptr2->A_d, usrDataptr2->A_h, Nbytes, + hipMemcpyHostToDevice, mystream2)); + hipLaunchKernelGGL((HipTest::vector_square), dim3(blocks), + dim3(threadsPerBlock), 0, mystream2, usrDataptr2->A_d, + usrDataptr2->C_d, NSize); + HIP_CHECK(hipMemcpyAsync(usrDataptr2->C_h, usrDataptr2->C_d, Nbytes, + hipMemcpyDeviceToHost, mystream2)); + HIP_CHECK(hipLaunchHostFunc(mystream2, Fn_validateMul_stream, + reinterpret_cast(usrDataptr2))); + HIP_CHECK(hipFreeAsync(reinterpret_cast(usrDataptr1->A_d), + mystream1)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(usrDataptr1->C_d), + mystream1)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(usrDataptr2->A_d), + mystream2)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(usrDataptr2->C_d), + mystream2)); + HIP_CHECK(hipStreamSynchronize(mystream1)); + HIP_CHECK(hipStreamSynchronize(mystream2)); + HIP_CHECK(hipStreamDestroy(mystream1)); + HIP_CHECK(hipStreamDestroy(mystream2)); + REQUIRE(usrDataptr1->isPassed); + REQUIRE(usrDataptr2->isPassed); + free(usrDataptr1->A_h); + free(usrDataptr1->C_h); + free(usrDataptr2->A_h); + free(usrDataptr2->C_h); + free(usrDataptr2); + free(usrDataptr1); +} + +// Test scenario 4 +// test case to validate hipLaunchHostFunc for the kernel, +// validate hipLaunchHostFunc after kernel launch. +static void Fn_Completion_state(void* userData) { + REQUIRE(userData != nullptr); + usrDataS *ptrUsrData = reinterpret_cast(userData); + ptrUsrData->isOpCompleted = true; +} + +TEST_CASE("Unit_hipLaunchHostFunc_KernelHost") { + hipStream_t stream1, stream2, stream3; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&stream3)); + usrDataS *usrDataptr = reinterpret_cast( + malloc(sizeof(usrDataS))); + REQUIRE(usrDataptr != nullptr); + usrDataptr->A_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr->A_h != nullptr); + usrDataptr->C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr->C_h != nullptr); + // input data + for (size_t i = 0; i < NSize; i++) { + usrDataptr->A_h[i] = 7.0f; + } + usrDataptr->isOpCompleted = false; + HIP_CHECK(hipMallocAsync(reinterpret_cast(&(usrDataptr->A_d)), + Nbytes, stream1)); + HIP_CHECK(hipMallocAsync(reinterpret_cast(&(usrDataptr->C_d)), + Nbytes, stream1)); + HIP_CHECK(hipMemcpyAsync(usrDataptr->A_d, usrDataptr->A_h, Nbytes, + hipMemcpyHostToDevice, stream1)); + HIP_CHECK(hipLaunchHostFunc(stream1, Fn_Completion_state, + reinterpret_cast(usrDataptr))); + while (!usrDataptr->isOpCompleted) { + std::this_thread::sleep_for(std::chrono::microseconds(100000)); + } // Sleep for 100 ms*/ + usrDataptr->isOpCompleted = false; + const unsigned blocks = GRIDSIZE; + const unsigned threadsPerBlock = BLOCKSIZE; + hipLaunchKernelGGL((HipTest::vector_square), dim3(blocks), + dim3(threadsPerBlock), 0, stream2, usrDataptr->A_d, + usrDataptr->C_d, NSize); + HIP_CHECK(hipLaunchHostFunc(stream2, Fn_Completion_state, + reinterpret_cast(usrDataptr))); + while (!usrDataptr->isOpCompleted) { + std::this_thread::sleep_for(std::chrono::microseconds(100000)); + } // Sleep for 100 ms*/ + usrDataptr->isOpCompleted = false; + HIP_CHECK(hipMemcpyAsync(usrDataptr->C_h, usrDataptr->C_d, Nbytes, + hipMemcpyDeviceToHost, stream3)); + HIP_CHECK(hipLaunchHostFunc(stream2, Fn_Completion_state, + reinterpret_cast(usrDataptr))); + HIP_CHECK(hipFreeAsync(reinterpret_cast(usrDataptr->A_d), + stream3)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(usrDataptr->C_d), + stream3)); + while (!usrDataptr->isOpCompleted) { + std::this_thread::sleep_for(std::chrono::microseconds(100000)); + } // Sleep for 100 ms*/ + for (size_t i = 0; i < NSize; i++) { + if (usrDataptr->C_h[i] != + (usrDataptr->A_h[i] * usrDataptr->A_h[i])) { + REQUIRE(false); + } + } + HIP_CHECK(hipStreamSynchronize(stream3)); + HIP_CHECK(hipStreamDestroy(stream3)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); + free(usrDataptr->A_h); + free(usrDataptr->C_h); + free(usrDataptr); +} + +// Test scenario 5 +// scenario that validates the host launch function on multi device +// environment. +TEST_CASE("Unit_hipLaunchHostFunc_multidevice") { + int num_devices; + HIP_CHECK(hipGetDeviceCount(&num_devices)); + if (num_devices < 2) { + SUCCEED("Skipping the testcases as numDevices < 2"); + return; + } + usrDataS *usrDataptr = reinterpret_cast( + malloc(sizeof(usrDataS))); + REQUIRE(usrDataptr != nullptr); + usrDataptr->A_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr->A_h != nullptr); + usrDataptr->C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr->C_h != nullptr); + for (size_t i = 0; i < NSize; i++) { + usrDataptr->A_h[i] = 21.0f; + } + for (int dev = 0; dev < num_devices; dev++) { + HIP_CHECK(hipSetDevice(dev)); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + launchOperationOnStrm(usrDataptr, stream); + HIP_CHECK(hipStreamDestroy(stream)); + } + free(usrDataptr->A_h); + free(usrDataptr->C_h); + free(usrDataptr); +} + +// Test scenario 6 +// scenario that validates the host launch function on created +// stream with same priority. +TEST_CASE("Unit_hipLaunchHostFunc_Samepriority") { + int priority = 0; + unsigned int flags = 0; + usrDataS *usrDataptr = reinterpret_cast( + malloc(sizeof(usrDataS))); + REQUIRE(usrDataptr != nullptr); + usrDataptr->A_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr->A_h != nullptr); + usrDataptr->C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr->C_h != nullptr); + for (size_t i = 0; i < NSize; i++) { + usrDataptr->A_h[i] = 21.0f; + } + for (int idx = 0; idx < NUM_OF_STREAM; idx++) { + hipStream_t stream[NUM_OF_STREAM]; + HIP_CHECK(hipStreamCreateWithPriority(&stream[idx], flags, priority)); + launchOperationOnStrm(usrDataptr, stream[idx]); + HIP_CHECK(hipStreamDestroy(stream[idx])); + } + free(usrDataptr->A_h); + free(usrDataptr->C_h); + free(usrDataptr); +} + +// Test scenario 7 +// scenario that validates the host launch function on +// created stream with different priority. +TEST_CASE("Unit_hipLaunchHostFunc_Diffpriority") { + int priority; + int priority_low{}; + int priority_high{}; + unsigned int flags = 0; + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + int numOfPriorities = priority_low - priority_high; + const float arr_size = numOfPriorities + 1; + hipStream_t *stream = reinterpret_cast( + malloc(arr_size*sizeof(hipStream_t))); + stream[0] = 0; + int count = 1; + // Create a stream for each of the priority levels + for (priority = priority_high; priority < priority_low; priority++) { + HIP_CHECK(hipStreamCreateWithPriority(&stream[count++], flags, priority)); + } + usrDataS *usrDataptr = reinterpret_cast( + malloc(sizeof(usrDataS))); + REQUIRE(usrDataptr != nullptr); + usrDataptr->A_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr->A_h != nullptr); + usrDataptr->C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(usrDataptr->C_h != nullptr); + for (size_t i = 0; i < NSize; i++) { + usrDataptr->A_h[i] = 11.0f; + } + for (int idx = 0; idx < arr_size; idx++) { + launchOperationOnStrm(usrDataptr, stream[idx]); + } + count = 1; + for (priority = priority_high; priority < priority_low; priority++) { + HIP_CHECK(hipStreamDestroy(stream[count++])); + } + free(usrDataptr->A_h); + free(usrDataptr->C_h); + free(usrDataptr); +} + +// Test scenario 8 +// create a graph by using hipGraphsUsingStreamCapture and call host function. + +typedef struct callBackData { + const char* fn_name; + double* data; +} callBackData_t; +double result_gpu = 0.0; +void myHostNodeCallback(void* data) { + static int iter = 0; + iter++; + // Check status of GPU after stream operations are done + callBackData_t* tmp = reinterpret_cast(data); + // checkCudaErrors(tmp->status); + double* result = reinterpret_cast(tmp->data); + const char* function = reinterpret_cast(tmp->fn_name); + if (iter == GRAPH_LAUNCH_ITERATIONS) + printf("[%s] Host callback final reduced sum = %lf\n", function, *result); + result_gpu = *result; + *result = 0.0; // reset the result +} + +TEST_CASE("Unit_hipLaunchHostFunc_Graph") { + size_t size = 1 << 12; + size_t maxBlocks = 512; + float *inputVec_d = NULL, *inputVec_h = NULL; + double *outputVec_d = NULL, *result_d; + inputVec_h = reinterpret_cast(malloc(sizeof(float) * size)); + HIP_CHECK(hipMalloc(&inputVec_d, sizeof(float) * size)); + HIP_CHECK(hipMalloc(&outputVec_d, sizeof(double) * maxBlocks)); + HIP_CHECK(hipMalloc(&result_d, sizeof(double))); + init_input(inputVec_h, size); + hipStream_t stream1, stream2, stream3, streamForGraph; + hipEvent_t forkStreamEvent, memsetEvent1, memsetEvent2; + hipGraph_t graph; + double result_h = 0.0; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&stream3)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipEventCreate(&forkStreamEvent)); + HIP_CHECK(hipEventCreate(&memsetEvent1)); + HIP_CHECK(hipEventCreate(&memsetEvent2)); + auto start = std::chrono::high_resolution_clock::now(); + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(forkStreamEvent, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0)); + HIP_CHECK(hipStreamWaitEvent(stream3, forkStreamEvent, 0)); + HIP_CHECK( + hipMemcpyAsync(inputVec_d, inputVec_h, sizeof(float) * size, + hipMemcpyDefault, stream1)); + HIP_CHECK(hipMemsetAsync(outputVec_d, 0, sizeof(double) * maxBlocks, + stream2)); + HIP_CHECK(hipEventRecord(memsetEvent1, stream2)); + HIP_CHECK(hipMemsetAsync(result_d, 0, sizeof(double), stream3)); + HIP_CHECK(hipEventRecord(memsetEvent2, stream3)); + HIP_CHECK(hipStreamWaitEvent(stream1, memsetEvent1, 0)); + hipLaunchKernelGGL(reduce, dim3(size / THREADS_PER_BLOCK, 1, 1), + dim3(THREADS_PER_BLOCK, 1, 1), 0, stream1, + inputVec_d, outputVec_d); + HIP_CHECK(hipStreamWaitEvent(stream1, memsetEvent2, 0)); + hipLaunchKernelGGL(reduceFinal, dim3(1, 1, 1), dim3(THREADS_PER_BLOCK, 1, 1), + 0, stream1, outputVec_d, result_d); + HIP_CHECK(hipMemcpyAsync(&result_h, result_d, sizeof(double), + hipMemcpyDefault, stream1)); + + callBackData_t hostFnData; + hostFnData.data = &result_h; + hostFnData.fn_name = "hipGraphsUsingStreamCapture"; + hipHostFn_t fn = myHostNodeCallback; + HIP_CHECK(hipLaunchHostFunc(stream1, fn, &hostFnData)); + + HIP_CHECK(hipStreamEndCapture(stream1, &graph)); + hipGraphNode_t* nodes = NULL; + size_t numNodes = 0; + HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes)); + printf("\nNum of nodes in the graph created using stream" + "capture API = %zu\n", numNodes); + HIP_CHECK(hipGraphGetRootNodes(graph, nodes, &numNodes)); + printf("root nodes in the graph created using stream capture API = %zu\n", + numNodes); + hipGraphExec_t graphExec; + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + auto start1 = std::chrono::high_resolution_clock::now(); + for (int i = 0; i < GRAPH_LAUNCH_ITERATIONS; i++) { + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + } + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + auto stop = std::chrono::high_resolution_clock::now(); + auto WithInit = std::chrono::duration(stop - start); + auto WithoutInit = std::chrono::duration(stop - start1); + std::cout << "Time taken for hipGraphsUsingStreamCapture with Init: " + << std::chrono::duration_cast(WithInit).count() + << " milliseconds without Init:" + << std::chrono::duration_cast(WithoutInit).count() + << " milliseconds " << std::endl; + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream3)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + double result_h_cpu = 0.0; + for (size_t i = 0; i < size; i++) { + result_h_cpu += inputVec_h[i]; + } + REQUIRE(result_h_cpu == result_gpu); + HIP_CHECK(hipFree(inputVec_d)); + HIP_CHECK(hipFree(outputVec_d)); + HIP_CHECK(hipFree(result_d)); +} From dc93a4e6cd43d049d424f2bba68d5ebec2b091a6 Mon Sep 17 00:00:00 2001 From: milos-mozetic <118800401+milos-mozetic@users.noreply.github.com> Date: Thu, 2 Mar 2023 07:59:30 +0100 Subject: [PATCH 2/2] EXSWHTEC-223 - Implement HIP API coverage tool (#98) - Upload the initial version of HIP API coverage tool - Add argument to fetch path to the HIP/include directory Within the new repository structure, HIP tests are separated from the HIP headers. Added parameter to the executable used to pass the path to the HIP/include directory. - Add tracking of test cases via doxygen comments - Add test cases to the HTML report - Adjust table size for HIP API test cases - Remove style class from the empty row in test cases table - Create TestCaseOccurrence to display more information about found Test Cases - Change the name of table columns to differentiate information - Remove duplicated path to the hiprtc.h file - Add newline at the end of file - Add newline at the end of makefile and coverage.css --- utils/coverage/Makefile | 38 ++ .../coverageReportHTML/resources/amber.png | Bin 0 -> 141 bytes .../coverageReportHTML/resources/coverage.css | 227 ++++++++++++ .../coverageReportHTML/resources/emerald.png | Bin 0 -> 141 bytes .../coverageReportHTML/resources/glass.png | Bin 0 -> 167 bytes .../coverageReportHTML/resources/ruby.png | Bin 0 -> 141 bytes .../coverageReportHTML/resources/snow.png | Bin 0 -> 141 bytes .../coverageReportHTML/resources/updown.png | Bin 0 -> 117 bytes utils/coverage/hipAPI.cpp | 222 ++++++++++++ utils/coverage/hipAPI.h | 75 ++++ utils/coverage/hipAPICoverageUtils.cpp | 284 +++++++++++++++ utils/coverage/hipAPICoverageUtils.h | 32 ++ utils/coverage/hipAPIGroup.cpp | 335 ++++++++++++++++++ utils/coverage/hipAPIGroup.h | 53 +++ utils/coverage/mainCoverage.cpp | 74 ++++ utils/coverage/reportGenerators.cpp | 307 ++++++++++++++++ utils/coverage/reportGenerators.h | 41 +++ 17 files changed, 1688 insertions(+) create mode 100644 utils/coverage/Makefile create mode 100644 utils/coverage/coverageReportHTML/resources/amber.png create mode 100644 utils/coverage/coverageReportHTML/resources/coverage.css create mode 100644 utils/coverage/coverageReportHTML/resources/emerald.png create mode 100644 utils/coverage/coverageReportHTML/resources/glass.png create mode 100644 utils/coverage/coverageReportHTML/resources/ruby.png create mode 100644 utils/coverage/coverageReportHTML/resources/snow.png create mode 100644 utils/coverage/coverageReportHTML/resources/updown.png create mode 100644 utils/coverage/hipAPI.cpp create mode 100644 utils/coverage/hipAPI.h create mode 100644 utils/coverage/hipAPICoverageUtils.cpp create mode 100644 utils/coverage/hipAPICoverageUtils.h create mode 100644 utils/coverage/hipAPIGroup.cpp create mode 100644 utils/coverage/hipAPIGroup.h create mode 100644 utils/coverage/mainCoverage.cpp create mode 100644 utils/coverage/reportGenerators.cpp create mode 100644 utils/coverage/reportGenerators.h diff --git a/utils/coverage/Makefile b/utils/coverage/Makefile new file mode 100644 index 0000000000..ee590bb186 --- /dev/null +++ b/utils/coverage/Makefile @@ -0,0 +1,38 @@ +# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +CC=g++ +CPPFLAGS=-std=c++17 +SRC=mainCoverage.cpp hipAPI.cpp hipAPIGroup.cpp reportGenerators.cpp hipAPICoverageUtils.cpp +OBJ=generateHipAPICoverage + +default_target: all +.PHONY : default_target + +all: ${SRC} + ${CC} ${CPPFLAGS} $^ -o ${OBJ} + +clean: + rm ${OBJ} + rm *.xml + rm -r ./coverageReportHTML/testModules + rm -r ./coverageReportHTML/testAPIs + rm ./coverageReportHTML/*.html +.PHONY : clean diff --git a/utils/coverage/coverageReportHTML/resources/amber.png b/utils/coverage/coverageReportHTML/resources/amber.png new file mode 100644 index 0000000000000000000000000000000000000000..2cab170d8359081983a4e343848dfe06bc490f12 GIT binary patch literal 141 zcmeAS@N?(olHy`uVBq!ia0vp^j3CU&3?x-=hn)ga>?NMQuI!iC1^G2tW}LqE04T&+ z;1OBOz`!j8!i<;h*8KqrvZOouIx;Y9?C1WI$O`1M1^9%x{(levWG?NMQuI!iC1^Jb!lvI6;R0X`wF(yt=9xVZRt1vCRixIA4P dLn>}1Cji+@42)0J?}79&c)I$ztaD0e0sy@GAL0N2 literal 0 HcmV?d00001 diff --git a/utils/coverage/coverageReportHTML/resources/glass.png b/utils/coverage/coverageReportHTML/resources/glass.png new file mode 100644 index 0000000000000000000000000000000000000000..e1abc00680a3093c49fdb775ae6bdb6764c95af2 GIT binary patch literal 167 zcmeAS@N?(olHy`uVBq!ia0vp^j3CU&3?x-=hn)gaEa{HEjtmSN`?>!lvI6;R0X`wF z|Ns97GD8ntt^-nxB|(0{3=Yq3q=7g|-tI089jvk*Kn`btM`SSr1Gf+eGhVt|_XjA* zUgGKN%6^Gmn4d%Ph(nkFP>9RZ#WAE}PI3Z}&BVayv3^M*kj3EX>gTe~DWM4f=_Dpv literal 0 HcmV?d00001 diff --git a/utils/coverage/coverageReportHTML/resources/ruby.png b/utils/coverage/coverageReportHTML/resources/ruby.png new file mode 100644 index 0000000000000000000000000000000000000000..991b6d4ec9e78be165e3ef757eed1aada287364d GIT binary patch literal 141 zcmeAS@N?(olHy`uVBq!ia0vp^j3CU&3?x-=hn)ga>?NMQuI!iC1^FceV#7`HfI^%F z9+AZi4BSE>%y{W;-5;PJOS+@4BLl<6e(pbstUx|nfKQ0)e^Y%R^MdiLxj>4`)5S5Q b;#P73kj=!v_*DHKNFRfztDnm{r-UW|iOwIS literal 0 HcmV?d00001 diff --git a/utils/coverage/coverageReportHTML/resources/snow.png b/utils/coverage/coverageReportHTML/resources/snow.png new file mode 100644 index 0000000000000000000000000000000000000000..2cdae107fceec6e7f02ac7acb4a34a82a540caa5 GIT binary patch literal 141 zcmeAS@N?(olHy`uVBq!ia0vp^j3CU&3?x-=hn)ga>?NMQuI!iC1^MM!lvI6;R0X`wF|Ns97GD8ntt^-nBo-U3d c6}OTTfNUlP#;5A{K>8RwUHx3vIVCg!071?oo&W#< literal 0 HcmV?d00001 diff --git a/utils/coverage/coverageReportHTML/resources/updown.png b/utils/coverage/coverageReportHTML/resources/updown.png new file mode 100644 index 0000000000000000000000000000000000000000..aa56a238b3e6c435265250f9266cd1b8caba0f20 GIT binary patch literal 117 zcmeAS@N?(olHy`uVBq!ia0vp^AT}Qd8;}%R+`Ae`*?77*hG?8mPH5^{)z4*}Q$iB}huR`+ literal 0 HcmV?d00001 diff --git a/utils/coverage/hipAPI.cpp b/utils/coverage/hipAPI.cpp new file mode 100644 index 0000000000..39ef99e4a2 --- /dev/null +++ b/utils/coverage/hipAPI.cpp @@ -0,0 +1,222 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hipAPI.h" + +FileOccurrence::FileOccurrence(std::string file_name, int line_number): + file_name{file_name}, line_number{line_number} {} + +TestCaseOccurrence::TestCaseOccurrence(std::string test_case_name, std::string file_name, int line_number): + FileOccurrence{file_name, line_number}, test_case_name{test_case_name} {} + +bool operator==(const HipAPI& l_hip_api, const HipAPI& r_hip_api) { + return l_hip_api.api_name == r_hip_api.api_name; +} + +bool operator<(const HipAPI& l_hip_api, const HipAPI& r_hip_api) { + return l_hip_api.api_name < r_hip_api.api_name; +} + +HipAPI::HipAPI(std::string api_name, bool deprecated_flag, std::string api_group_name): + api_name{api_name}, deprecated{deprecated_flag}, api_group_name{api_group_name} {} + +std::string HipAPI::getName() const { + return api_name; +} + +std::string HipAPI::getGroupName() const { + return api_group_name; +} + +int HipAPI::getNumberOfCalls() const { + return file_occurrences.size(); +} + +int HipAPI::getNumberOfTestCases() const { + return test_cases.size(); +} + +void HipAPI::addFileOccurrence(FileOccurrence file_occurrence) { + file_occurrences.push_back(file_occurrence); +} + +void HipAPI::addTestCase(TestCaseOccurrence test_case) { + test_cases.push_back(test_case); +} + +bool HipAPI::isDeprecated() const +{ + return deprecated; +} + +std::string HipAPI::getBasicStatsXML() const +{ + std::stringstream xml_node; + xml_node << "\t\t\n"; + + if (!deprecated) { + xml_node << "\t\t\t" << api_name << "\n"; + } else { + xml_node << "\t\t\t" << "[DEPRECATED] " << api_name << "\n"; + } + + if (!file_occurrences.empty()) { + xml_node << "\t\t\t" << file_occurrences.size() << "\n"; + xml_node << "\t\t\t\n"; + for (auto const& file_occurrence: file_occurrences) { + xml_node << "\t\t\t\t" << file_occurrence.file_name << ":" << file_occurrence.line_number << "\n"; + } + xml_node << "\t\t\t\n"; + } + + xml_node << "\t\t\n"; + return xml_node.str(); +} + +std::string HipAPI::createHTMLReport() const { + std::stringstream html_report; + std::string one_tab{"\n\t"}; + std::string two_tabs{"\n\t\t"}; + std::string three_tabs{"\n\t\t\t"}; + std::string four_tabs{"\n\t\t\t\t"}; + std::string five_tabs{"\n\t\t\t\t\t"}; + std::string six_tabs{"\n\t\t\t\t\t\t"}; + + html_report << ""; + html_report << "" << one_tab << ""; + html_report << one_tab << "" << api_name << " Coverage report" << one_tab << "" << one_tab<< ""; + html_report << one_tab << "" << one_tab << ""; + html_report << two_tabs << ""; + html_report << two_tabs << "\n"; + html_report << two_tabs << "" << three_tabs << ""; + html_report << two_tabs << ""; + + html_report << two_tabs << "\n"; + html_report << one_tab << "
" << api_name << " Coverage report
" << four_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + html_report << five_tabs << ""; + html_report << four_tabs << "
"; + html_report << six_tabs << "
Group:" << api_group_name << "
Calls within test source files:" << file_occurrences.size() << "
Number of test cases:" << test_cases.size() << "
"; + html_report << three_tabs << "
"; + + html_report << one_tab << "
"; + // Add info about test cases + html_report << one_tab << ""; + if (!test_cases.empty()) { + html_report << two_tabs << ""; + html_report << three_tabs << ""; + html_report << three_tabs << ""; + html_report << three_tabs << ""; + html_report << two_tabs << ""; + + html_report << two_tabs << ""; + html_report << three_tabs << ""; + html_report << three_tabs << ""; + html_report << three_tabs << ""; + html_report << two_tabs << ""; + + for (auto const& test_case: test_cases) { + html_report << two_tabs << ""; + html_report << three_tabs << ""; + html_report << three_tabs << ""; + html_report << three_tabs << ""; + html_report << two_tabs << ""; + } + } else { + html_report << two_tabs << ""; + html_report << three_tabs << ""; + html_report << two_tabs << ""; + } + html_report << one_tab << "



Test case IDTest case occurrence in fileLine number
" << test_case.test_case_name << "" << test_case.file_name << "" << test_case.line_number << "

There are no test cases detected within doxygen comments.
"; + + html_report << one_tab << "
"; + html_report << one_tab << ""; + html_report << two_tabs << ""; + html_report << one_tab << "
"; + + // Add info about API occurrences in the test files. + html_report << one_tab << ""; + if (!file_occurrences.empty()) { + html_report << two_tabs << ""; + html_report << three_tabs << ""; + html_report << three_tabs << ""; + html_report << two_tabs << ""; + + html_report << two_tabs << ""; + html_report << three_tabs << ""; + html_report << three_tabs << ""; + html_report << two_tabs << ""; + + for (auto const& file_occurrence: file_occurrences) { + html_report << two_tabs << ""; + html_report << three_tabs << ""; + html_report << three_tabs << ""; + html_report << two_tabs << ""; + } + } else { + html_report << two_tabs << ""; + html_report << three_tabs << ""; + html_report << two_tabs << ""; + } + html_report << one_tab << "


API occurrence in fileLine number
" << file_occurrence.file_name << "" << file_occurrence.line_number << "

There are no occurrences within test source files.
"; + html_report << one_tab << "
"; + + html_report << one_tab << "
"; + html_report << one_tab << ""; + html_report << two_tabs << ""; + + time_t now{time(nullptr)}; + std::string date{asctime(gmtime(&now))}; + html_report << two_tabs << ""; + html_report << one_tab << "
Generated: " << date; + html_report << two_tabs << " UTC
"; + html_report << one_tab << "
"; + html_report << "\n\n"; + + return html_report.str(); +} diff --git a/utils/coverage/hipAPI.h b/utils/coverage/hipAPI.h new file mode 100644 index 0000000000..b14d3b7c6e --- /dev/null +++ b/utils/coverage/hipAPI.h @@ -0,0 +1,75 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include + +/* +Helper class used to store information in what file has HIP API been detected, +and on what line of code in that file. +*/ +class FileOccurrence { + public: + std::string file_name; + int line_number; + FileOccurrence(std::string file_name, int line_number); +}; + +/* +Helper class used to store information in what file has the API Test Case been detected, +and on what line of code in that file. +*/ +class TestCaseOccurrence : public FileOccurrence { + public: + std::string test_case_name; + TestCaseOccurrence(std::string test_case_name, std::string file_name, int line_number); +}; + +/* +Class used to store infromation about each HIP API. All information +is related to the API name, number of calls from test .cc files, +and its status of deprecation. +*/ +class HipAPI { + friend bool operator==(const HipAPI& l_hip_api, const HipAPI& r_hip_api); + friend bool operator<(const HipAPI& l_hip_api, const HipAPI& r_hip_api); + + public: + HipAPI(std::string api_name, bool deprecated_flag, std::string api_group_name); + std::string getName() const; + std::string getGroupName() const; + int getNumberOfCalls() const; + int getNumberOfTestCases() const; + void addFileOccurrence(FileOccurrence file_occurence); + void addTestCase(TestCaseOccurrence test_case); + bool isDeprecated() const; + std::string getBasicStatsXML() const; + std::string createHTMLReport() const; + private: + std::string api_name; + int number_of_calls; + bool deprecated; + std::string api_group_name; + std::vector file_occurrences; + std::vector test_cases; +}; diff --git a/utils/coverage/hipAPICoverageUtils.cpp b/utils/coverage/hipAPICoverageUtils.cpp new file mode 100644 index 0000000000..022815768c --- /dev/null +++ b/utils/coverage/hipAPICoverageUtils.cpp @@ -0,0 +1,284 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hipAPICoverageUtils.h" + +/* +Used to find all HIP API occurrences within the passed test .cc files. +Occurrence is detected when the HIP API is called in several scenarios. + - API Call with assertion, e.g. REQUIRE(hipAPI()); + - API call with assignment, e.g. result = hipAPI(); + - API call with parameter, e.g. ASSERT_EQUAL(hipSuccess, hipGetDevice(&deviceId)); +*/ +void findAPICallInFile(HipAPI& hip_api, std::string test_module_file) { + std::fstream test_module_file_handler; + test_module_file_handler.open(test_module_file); + + int line_number{0}; + std::string line; + /* + Closed brackets might be in another line, if the API + call is multiline. + */ + std::string api_call_with_assert{"(" + hip_api.getName() + "("}; + std::string api_call_with_assignment{"= " + hip_api.getName() + "("}; + std::string api_call_with_parameter{", " + hip_api.getName() + "("}; + + while (std::getline(test_module_file_handler, line)) { + ++line_number; + if ((line.find(api_call_with_assert) != std::string::npos) || + (line.find(api_call_with_assignment) != std::string::npos) || + (line.find(api_call_with_parameter) != std::string::npos)) { + hip_api.addFileOccurrence(FileOccurrence(test_module_file, line_number)); + } + } + + test_module_file_handler.close(); +} + +/* +Used to find all HIP API test cases within the passed test .cc files. +Matching test case is detected when the HIP API in defined within doxygen comment. +*/ +void findAPITestCaseInFile(HipAPI& hip_api, std::string test_module_file) { + std::fstream test_module_file_handler; + test_module_file_handler.open(test_module_file); + + int line_number{0}; + std::string line; + + std::string add_group_definition{"@addtogroup"}; + std::string ref_test_case{"@ref"}; + std::string test_case_definition{"TEST_CASE("}; + std::string current_api_name{"None"}; + std::string test_case{"None"}; + + while (std::getline(test_module_file_handler, line)) { + ++line_number; + if (line.find(add_group_definition) != std::string::npos) { + current_api_name = line.substr(line.find(add_group_definition) + 1); + current_api_name = current_api_name.substr(current_api_name.rfind(" ") + 1); + } + + if (hip_api.getName() != current_api_name) { + continue; + } + + if (line.find(ref_test_case) != std::string::npos) { + test_case = line.substr(line.rfind(" ") + 1); + hip_api.addTestCase(TestCaseOccurrence{test_case, test_module_file, line_number}); + continue; + } + + if (line.find(test_case_definition) != std::string::npos) { + test_case = line.substr(line.find("\"") + 1); + test_case = test_case.substr(0, test_case.find("\"")); + hip_api.addTestCase(TestCaseOccurrence{test_case, test_module_file, line_number}); + } + } + + test_module_file_handler.close(); +} + +/* +Used to iterate through all passed test .cc files and search for passed +HIP API instance. This instance shall be used to update occurrences. +*/ +void searchForAPI(HipAPI& hip_api, std::vector& test_module_files) { + std::cout << "Searching for " << hip_api.getName() << " in test module files." << std::endl; + for (auto const& test_module_file: test_module_files) { + findAPICallInFile(hip_api, test_module_file); + findAPITestCaseInFile(hip_api, test_module_file); + } +} + +/* +Used to extract all HIP APIs from the passed header file. +*/ +std::vector extractHipAPIs(std::string& hip_api_header_file, + std::vector& api_group_names, bool start_groups) { + std::fstream hip_header_file_handler; + hip_header_file_handler.open(hip_api_header_file); + + std::string line; + std::vector hip_apis; + + /* + If the HIP API is deprecated, it will be marked with + the following deprecation line in code. + */ + std::string deprecated_line{"DEPRECATED("}; + /* + Each HIP API has prefix hip in the name. Groups are marked with @defgroup, and the + main group that shall be considered is HIP API. Before that group is defined, lines + of code shall not be considered. + */ + std::string hip_api_prefix{"hip"}; + std::string group_definition{"@defgroup"}; + std::string add_group_definition{"@addtogroup"}; + std::string start_of_api_groups{"HIP API"}; + std::string end_of_api_groups{"doxygen end HIP API"}; + std::string end_group_definition{"@}"}; + + bool deprecated_flag{false}; + bool api_group_names_start{start_groups}; + std::vector api_group_names_tracker{api_group_names}; + int line_number{0}; + + while (std::getline(hip_header_file_handler, line)) { + ++line_number; + + // Declarations of the HIP APIs start after the HIP API group has been defined. + if ((line.find(group_definition) != std::string::npos || line.find(add_group_definition) != std::string::npos) + && line.find(start_of_api_groups) != std::string::npos) { + api_group_names_start = true; + continue; + } + + // If the API Groups have not started yet, go to the next file line. + if (!api_group_names_start) { + continue; + } + + // If the end of HIP API group has been detected, break the loop. + if (line.find(end_of_api_groups) != std::string::npos) { + break; + } + + /* + If the API is deprecated, raise a flag and go to the + next line where the API is declared. + */ + if (line.find(deprecated_line) != std::string::npos) { + std::getline(hip_header_file_handler, line); + ++line_number; + deprecated_flag = true; + } else { + deprecated_flag = false; + } + + if (line.find(group_definition) != std::string::npos) { + std::string group_name = line.substr(line.find(group_definition) + group_definition.length() + 1); + group_name = group_name.substr(group_name.find(' ') + 1); + api_group_names.push_back(group_name); + api_group_names_tracker.push_back(group_name); + } + else if (line.find(add_group_definition) != std::string::npos) + { + std::string group_name = line.substr(line.find(add_group_definition) + add_group_definition.length() + 1); + group_name = group_name.substr(group_name.find(' ') + 1); + api_group_names.push_back(group_name); + api_group_names_tracker.push_back(group_name); + } + + /* + Possible case is that there are nested groups. While api_group_names + vector contains all detected groups, api_group_names_tracker is responsible + to track the last defined group, because of the nested cases. + */ + if (line.find(end_group_definition) != std::string::npos) { + api_group_names_tracker.pop_back(); + } + + /* + The line which contains HIP API declaration looks like: + hipError_t hipMalloc(void** ptr, size_t size); + The name of HIP API is found by following steps: + - Take the substring from the start to the first open bracket. + - Extract the name from that substring by finding the last "hip". + Avoid comments. + */ + if (line.find(hip_api_prefix) != std::string::npos && + line.find("(") != std::string::npos && + line.find(" ") != 0 && + line.find(" *") != 0) { + std::string api_name_no_brackets{line.substr(0, line.find("("))}; + /* + If there is no hip substring, then there is no valid API in that line. + */ + if (api_name_no_brackets.rfind(hip_api_prefix) == std::string::npos) { + continue; + } + + /* + Extract the substring that starts from the last hip substring to the + end of that string (until the open bracket from the original string). + Remove all spaces if they exist in the parsed string, e.g., + hipError_t hipDeviceSetLimit ( enum hipLimit_t limit, size_t value );. + */ + std::string api_name{api_name_no_brackets.substr(api_name_no_brackets.rfind(hip_api_prefix))}; + api_name.erase(std::remove(api_name.begin(), api_name.end(), ' '), api_name.end()); + + if (!api_group_names_tracker.empty()) { + /* + If the API is not present in the global list of HIP APIs, add it. + */ + HipAPI hip_api{api_name, deprecated_flag, api_group_names_tracker.back()}; + if(std::find(hip_apis.begin(), hip_apis.end(), hip_api) == hip_apis.end()) + { + hip_apis.push_back(hip_api); + } + } else { + std::cout << "[SKIP_FROM_COV] Group not detected for \"" << api_name << "\" in file \"" + << hip_api_header_file << "\", line " << line_number << std::endl; + } + } + } + + hip_header_file_handler.close(); + return hip_apis; +} + +/* +Used to extract test .cc files from the passed tests root directory. +Goes through all subdirectories and searches for .cc and .hh files for +HIP API invocations. +Implements BFS algorithm to avoid recursion. +*/ +std::vector extractTestModuleFiles(std::string& tests_root_directory) +{ + std::vector directory_queue; + directory_queue.push_back(tests_root_directory); + std::vector test_module_files; + + while (!directory_queue.empty()) { + std::string processed_entry = directory_queue.back(); + directory_queue.pop_back(); + for (const auto& entry: std::filesystem::directory_iterator(processed_entry)) { + if (std::filesystem::is_directory(entry.path())) { + directory_queue.push_back(entry.path()); + } else { + if (entry.path().string().find(".cc") != std::string::npos || + entry.path().string().find(".hh") != std::string::npos) { + test_module_files.push_back(entry.path()); + } + } + } + } + + return test_module_files; +} + +std::string findAbsolutePathOfFile(std::string file_path) +{ + return std::filesystem::canonical(std::filesystem::absolute(file_path)); +} diff --git a/utils/coverage/hipAPICoverageUtils.h b/utils/coverage/hipAPICoverageUtils.h new file mode 100644 index 0000000000..68454fd863 --- /dev/null +++ b/utils/coverage/hipAPICoverageUtils.h @@ -0,0 +1,32 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include "hipAPIGroup.h" + +void findAPICallInFile(HipAPI& hip_api, std::string test_module_file); +void findAPITestCaseInFile(HipAPI& hip_api, std::string test_module_file); +void searchForAPI(HipAPI& hip_api, std::vector& test_module_files); +std::vector extractHipAPIs(std::string& hip_api_header_file, std::vector& api_group_names, bool start_groups); +std::vector extractTestModuleFiles(std::string& tests_root_directory); +std::string findAbsolutePathOfFile(std::string file_path); diff --git a/utils/coverage/hipAPIGroup.cpp b/utils/coverage/hipAPIGroup.cpp new file mode 100644 index 0000000000..b24e2f1baa --- /dev/null +++ b/utils/coverage/hipAPIGroup.cpp @@ -0,0 +1,335 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hipAPIGroup.h" + +bool operator==(const HipAPIGroup& l_hip_api_group, const HipAPIGroup& r_hip_api_group) { + return l_hip_api_group.group_name == r_hip_api_group.group_name; +} + +HipAPIGroup::HipAPIGroup(std::string group_name, std::vector& hip_apis): + group_name{group_name}, number_of_api_calls{0}, percentage_of_called_apis{0.f}, + total_number_of_apis{0}, number_of_test_cases{0} +{ + for (auto const& hip_api: hip_apis) { + if (hip_api.getGroupName() != group_name) { + continue; + } + + if (hip_api.isDeprecated()) { + deprecated_apis.push_back(hip_api); + } else { + if (hip_api.getNumberOfCalls()) { + called_apis.push_back(hip_api); + } else { + not_called_apis.push_back(hip_api); + } + } + + number_of_api_calls += hip_api.getNumberOfCalls(); + number_of_test_cases += hip_api.getNumberOfTestCases(); + } + + total_number_of_apis = called_apis.size() + not_called_apis.size() + deprecated_apis.size(); + if (not_called_apis.empty()) { + percentage_of_called_apis = 100.f; + } else { + percentage_of_called_apis = 100.f * called_apis.size() / (total_number_of_apis - deprecated_apis.size()); + } +} + +std::string HipAPIGroup::getName() const { + return group_name; +} + +int HipAPIGroup::getTotalNumberOfAPIs() const { + return total_number_of_apis; +} + +int HipAPIGroup::getTotalNumberOfCalls() const { + return number_of_api_calls; +} + +int HipAPIGroup::getTotalNumberOfTestCases() const { + return number_of_test_cases; +} + +int HipAPIGroup::getNumberOfCalledAPIs() const { + return called_apis.size(); +} + +int HipAPIGroup::getNumberOfNotCalledAPIs() const { + return not_called_apis.size(); +} + +int HipAPIGroup::getNumberOfDeprecatedAPIs() const { + return deprecated_apis.size(); +} + +float HipAPIGroup::getPercentageOfCalledAPIs() const { + return percentage_of_called_apis; +} + +std::string HipAPIGroup::getBasicStatsXML() const { + std::stringstream xml_node; + + std::string tag_name; + std::transform(group_name.begin(), group_name.end(), std::back_inserter(tag_name), ::toupper); + std::replace(tag_name.begin(), tag_name.end(), ' ', '-'); + + xml_node << "\n<" << tag_name << ">"; + + xml_node << "\n\t"; + xml_node << "\n\t\t" << total_number_of_apis << ""; + xml_node << "\n\t\t" << number_of_api_calls << ""; + xml_node << "\n\t\t" << called_apis.size() << ""; + xml_node << "\n\t\t" << not_called_apis.size() << ""; + xml_node << "\n\t\t" << deprecated_apis.size() << ""; + xml_node << "\n\t\t" << percentage_of_called_apis << "%"; + xml_node << "\n\t"; + + if (!called_apis.empty()) { + xml_node << "\n\t\n"; + for (auto const& hip_api: called_apis) { + xml_node << hip_api.getBasicStatsXML(); + } + xml_node << "\t"; + } + + if (!not_called_apis.empty()) { + xml_node << "\n\t\n"; + for (auto const& hip_api: not_called_apis) { + xml_node << hip_api.getBasicStatsXML(); + } + xml_node << "\t"; + } + + if (!deprecated_apis.empty()) { + xml_node << "\n\t\n"; + for (auto const& hip_api: deprecated_apis) { + xml_node << hip_api.getBasicStatsXML(); + } + xml_node << "\t"; + } + + xml_node << "\n"; + return xml_node.str(); +} + +std::string HipAPIGroup::getBasicStatsHTML() const +{ + std::stringstream html_object; + std::string two_tabs{"\n\t\t"}; + std::string three_tabs{"\n\t\t\t"}; + std::string four_tabs{"\n\t\t\t\t"}; + std::string five_tabs{"\n\t\t\t\t\t"}; + + // Determine font class from coverage.css and image for color bar. + std::string font_class; + std::string color_bar; + + if (percentage_of_called_apis < 40.f) { + font_class = "coverNumLo"; + color_bar = "resources/ruby.png"; + } else if (percentage_of_called_apis < 80.f) { + font_class = "coverNumMed"; + color_bar = "resources/amber.png"; + } else { + font_class = "coverNumHi"; + color_bar = "resources/emerald.png"; + } + + html_object << two_tabs << ""; + html_object << three_tabs << "" << group_name << ""; + html_object << three_tabs << "" << total_number_of_apis << ""; + html_object << three_tabs << "" << number_of_api_calls << ""; + html_object << three_tabs << "" << number_of_test_cases << ""; + html_object << three_tabs << "" << called_apis.size() << ""; + html_object << three_tabs << "" << not_called_apis.size() << ""; + html_object << three_tabs << "" << deprecated_apis.size() << ""; + html_object << three_tabs << "" << std::fixed << std::setprecision(2) << percentage_of_called_apis << "%"; + + html_object << three_tabs << ""; + html_object << four_tabs << ""; + html_object << five_tabs << "
\"""; + html_object << "\""
"; + html_object << four_tabs << ""; + html_object << three_tabs << ""; + + return html_object.str(); +} + +std::string HipAPIGroup::createHTMLReport() const +{ + std::stringstream html_report; + std::string one_tab{"\n\t"}; + std::string two_tabs{"\n\t\t"}; + std::string three_tabs{"\n\t\t\t"}; + std::string four_tabs{"\n\t\t\t\t"}; + std::string five_tabs{"\n\t\t\t\t\t"}; + std::string six_tabs{"\n\t\t\t\t\t\t"}; + + html_report << ""; + html_report << "" << one_tab << ""; + html_report << one_tab << "" << group_name << " Coverage report"; + html_report << one_tab << "" << one_tab<< ""; + html_report << one_tab << "" << one_tab << ""; + html_report << two_tabs << ""; + html_report << two_tabs << "\n"; + html_report << two_tabs << "" << three_tabs << ""; + html_report << two_tabs << ""; + + html_report << two_tabs << "\n"; + html_report << one_tab << "
" << group_name << " Coverage report
" << four_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + // Determine font class from coverage.css file based on coverage percentage. + std::string font_class; + if (percentage_of_called_apis < 40.f) { + font_class = "headerCovTableEntryLo"; + } else if (percentage_of_called_apis < 80.f) { + font_class = "headerCovTableEntryMed"; + } else { + font_class = "headerCovTableEntryHi"; + } + + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + + html_report << five_tabs << ""; + html_report << four_tabs << "
Value
Total number of detected HIP APIs:" << total_number_of_apis << "
HIP API calls within test source files:" << number_of_api_calls << "
Total number of test cases:" << number_of_test_cases << "
HIP APIs that are called at least once:" << called_apis.size() << "
HIP APIs that are not called at all:" << not_called_apis.size() << "
HIP APIs that are marked as deprecated:" << deprecated_apis.size() << "
Test coverage by implemented tests for the HIP APIs:" << std::fixed << std::setprecision(2) << percentage_of_called_apis << "%
"; + html_report << three_tabs << "
"; + + // Add info about Test module APIs. + html_report << one_tab << "
"; + html_report << one_tab << ""; + html_report << two_tabs << ""; + html_report << three_tabs << ""; + html_report << three_tabs << ""; + html_report << two_tabs << ""; + html_report << two_tabs << ""; + + html_report << three_tabs << ""; + + html_report << three_tabs << ""; + + html_report << three_tabs << ""; + + html_report << two_tabs << ""; + + html_report << one_tab << "

"; + html_report << three_tabs << "
"; + html_report << four_tabs << ""; + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + for (auto const& hip_api: called_apis) { + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + } + + html_report << four_tabs << "
Called APIs
" << hip_api.getName() << "
"; + html_report << three_tabs << "		"; + html_report << four_tabs << ""; + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + for (auto const& hip_api: not_called_apis) { + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + } + html_report << four_tabs << "
Not called APIs
" << hip_api.getName() << "
"; + html_report << three_tabs << "		"; + html_report << four_tabs << ""; + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + for (auto const& hip_api: deprecated_apis) { + html_report << five_tabs << ""; + html_report << six_tabs << ""; + html_report << five_tabs << ""; + } + html_report << four_tabs << "
Deprecated APIs
" << hip_api.getName() << "
"; + html_report << three_tabs << "
"; + html_report << one_tab << "
"; + html_report << one_tab << "
"; + + html_report << one_tab << ""; + html_report << two_tabs << ""; + + time_t now{time(nullptr)}; + std::string date{asctime(gmtime(&now))}; + html_report << two_tabs << ""; + html_report << one_tab << "
Generated: " << date; + html_report << two_tabs << " UTC
"; + html_report << one_tab << "
"; + html_report << "\n\n"; + + return html_report.str(); +} diff --git a/utils/coverage/hipAPIGroup.h b/utils/coverage/hipAPIGroup.h new file mode 100644 index 0000000000..c9b07b8656 --- /dev/null +++ b/utils/coverage/hipAPIGroup.h @@ -0,0 +1,53 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hipAPI.h" +#include +#include + +class HipAPIGroup{ + friend bool operator==(const HipAPIGroup& l_hip_api_group, const HipAPIGroup& r_hip_api_group); + + public: + HipAPIGroup(std::string group_name, std::vector& hip_apis); + std::string getName() const; + int getTotalNumberOfAPIs() const; + int getTotalNumberOfCalls() const; + int getTotalNumberOfTestCases() const; + int getNumberOfCalledAPIs() const; + int getNumberOfNotCalledAPIs() const; + int getNumberOfDeprecatedAPIs() const; + float getPercentageOfCalledAPIs() const; + std::string getBasicStatsXML() const; + std::string getBasicStatsHTML() const; + std::string createHTMLReport() const; + private: + std::string group_name; + int total_number_of_apis; + int number_of_api_calls; + float percentage_of_called_apis; + int number_of_test_cases; + std::string parent_group_name; + std::vector called_apis; + std::vector not_called_apis; + std::vector deprecated_apis; +}; diff --git a/utils/coverage/mainCoverage.cpp b/utils/coverage/mainCoverage.cpp new file mode 100644 index 0000000000..65de5890a8 --- /dev/null +++ b/utils/coverage/mainCoverage.cpp @@ -0,0 +1,74 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "reportGenerators.h" + +int main(int argc, char** argv) +{ + if (argc != 2) { + std::cout << "Please provide the path to the cloned HIP/include/ directory as an argument! Only one argument supported." << std::endl; + std::cout << "\tExample: ./generateHipAPICoverage /workspace/user1/HIP/include/" << std::endl; + return -1; + } + std::string hip_include_path = argv[1]; + /* + Relative paths to all needed files, as it is expected that the application + is called from the HIP/tests/catch/coverage directory. + */ + std::string hip_api_header_file{findAbsolutePathOfFile(hip_include_path + "/hip/hip_runtime_api.h")}; + std::string hip_rtc_header_file{findAbsolutePathOfFile(hip_include_path + "/hip/hiprtc.h")}; + std::string tests_root_directory{findAbsolutePathOfFile("../../catch")}; + + std::vector api_group_names; + // Extract all HIP API declarations from the HIP API header file. + std::vector hip_apis{extractHipAPIs(hip_api_header_file, api_group_names, false)}; + std::cout << "Number of detected HIP APIs from " << hip_api_header_file << ": " << hip_apis.size() << std::endl; + + std::vector hip_rtc_apis{extractHipAPIs(hip_rtc_header_file, api_group_names, true)}; + std::cout << "Number of detected HIP APIs from " << hip_rtc_header_file << ": " << hip_rtc_apis.size() << std::endl; + hip_apis.insert(hip_apis.end(), hip_rtc_apis.begin(), hip_rtc_apis.end()); + + // Extract all test module .cc files that shall be used for API searching. + std::cout << "Searching for HIP API calls in source files within " << tests_root_directory << "." << std::endl; + std::vector test_module_files{extractTestModuleFiles(tests_root_directory)}; + + // Search for each HIP API in the extracted test .cc files. + for(HipAPI& hip_api: hip_apis) { + searchForAPI(hip_api, test_module_files); + } + + std::vector hip_api_groups; + for (auto const& api_group_name: api_group_names) { + HipAPIGroup hip_api_group{api_group_name, hip_apis}; + // Avoid having duplicated groups. + if (std::find(hip_api_groups.begin(), hip_api_groups.end(), hip_api_group) == hip_api_groups.end()) { + hip_api_groups.push_back(hip_api_group); + } + } + + std::cout << "Generating XML report files." << std::endl; + generateXMLReportFiles(hip_apis, hip_api_groups); + std::cout << "Generating HTML report files." << std::endl; + generateHTMLReportFiles(hip_apis, hip_api_groups, tests_root_directory, hip_api_header_file, hip_rtc_header_file); + + return 0; +} diff --git a/utils/coverage/reportGenerators.cpp b/utils/coverage/reportGenerators.cpp new file mode 100644 index 0000000000..7b74282929 --- /dev/null +++ b/utils/coverage/reportGenerators.cpp @@ -0,0 +1,307 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "reportGenerators.h" + +BasicAPIStats::BasicAPIStats(std::vector& hip_api_groups): + number_of_called_apis{0}, number_of_not_called_apis{0}, + number_of_deprecated_apis{0}, total_number_of_api_calls{0}, + total_number_of_test_cases{0} +{ + for (auto const& hip_api_group: hip_api_groups) { + number_of_called_apis += hip_api_group.getNumberOfCalledAPIs(); + number_of_not_called_apis += hip_api_group.getNumberOfNotCalledAPIs(); + number_of_deprecated_apis += hip_api_group.getNumberOfDeprecatedAPIs(); + total_number_of_api_calls += hip_api_group.getTotalNumberOfCalls(); + total_number_of_test_cases += hip_api_group.getTotalNumberOfTestCases(); + } + total_number_of_apis = number_of_called_apis + number_of_not_called_apis + number_of_deprecated_apis; + tests_coverage_percentage = 100.f * number_of_called_apis / (number_of_called_apis + number_of_not_called_apis); +} + +float BasicAPIStats::getLowCoverageLimit() const { + return 40.f; +} + +float BasicAPIStats::getMediumCoverageLimit() const { + return 80.f; +} + +void generateXMLReportFiles(std::vector& hip_apis, std::vector& hip_api_groups) { + BasicAPIStats basic_stats{hip_api_groups}; + + std::cout << "Total number of HIP API calls: " << basic_stats.total_number_of_api_calls << std::endl; + std::cout << "Number of the HIP APIs that are called at least once: " << basic_stats.number_of_called_apis << std::endl; + std::cout << "Number of the HIP APIs that are not called at all: " << basic_stats.number_of_not_called_apis << std::endl; + std::cout << "Number of the HIP APIs that are marked as deprecated: " << basic_stats.number_of_deprecated_apis << std::endl; + std::cout << "Test coverage by implemented tests, for the HIP APIs that are not marked as deprecated: "; + std::cout << basic_stats.tests_coverage_percentage << "%" << std::endl; + + /* + Generate XML file that contains relevant information about test coverage. + The XML file is created using raw handling of XML files, as there is no need + for the additional 3rd party library that implements XML file CRUD operations. + */ + std::fstream coverage_report; + std::string report_file_name{"CoverageReport.xml"}; + coverage_report.open(report_file_name, std::ios::out); + + time_t now{time(nullptr)}; + std::string date{asctime(gmtime(&now))}; + coverage_report << "" << date << "\n"; + + coverage_report << "\n"; + + coverage_report << "\t\n\t\tTotal number of detected HIP APIs."; + coverage_report << "\n\t\t" << hip_apis.size() << "\n\t\n"; + + coverage_report << "\t\n\t\tTotal number of HIP API calls within test source files."; + coverage_report << "\n\t\t" << basic_stats.total_number_of_api_calls << "\n\t\n"; + + coverage_report << "\t\n\t\tNumber of the HIP APIs that are called at least once."; + coverage_report << "\n\t\t" << basic_stats.number_of_called_apis << "\n\t\n"; + + coverage_report << "\t\n\t\tNumber of the HIP APIs that are not called at all."; + coverage_report << "\n\t\t" << basic_stats.number_of_not_called_apis << "\n\t\n"; + + coverage_report << "\t\n\t\tNumber of the HIP APIs that are marked as deprecated."; + coverage_report << "\n\t\t" << basic_stats.number_of_deprecated_apis << "\n\t\n"; + + coverage_report << "\t\n\t\tTest coverage by implemented tests for the HIP APIs that are not marked as deprecated."; + coverage_report << "\n\t\t" << basic_stats.tests_coverage_percentage << "%\n\t"; + + coverage_report << "\n"; + + for (auto const& hip_api_group: hip_api_groups) { + coverage_report << hip_api_group.getBasicStatsHTML(); + } + + coverage_report.close(); + std::cout << "Generated XML report file " << findAbsolutePathOfFile(report_file_name) << std::endl; +} + +void generateHTMLReportFiles(std::vector& hip_apis, std::vector& hip_api_groups, + std::string tests_root_directory, std::string hipApiHeaderFile, std::string hip_rtc_header_file) { + BasicAPIStats basic_stats{hip_api_groups}; + + std::fstream coverage_report; + // Main HTML report file. + std::string report_file_name{"./coverageReportHTML/CoverageReport.html"}; + // Directories used to store generated HTML files. + std::string test_modules_directory{"./coverageReportHTML/testModules"}; + std::string test_apis_directory{"./coverageReportHTML/testAPIs"}; + std::filesystem::create_directories(test_modules_directory); + std::filesystem::create_directories(test_apis_directory); + + coverage_report.open(report_file_name, std::ios::out); + + // Helper strings with tabs and newlines for better HTML formatting. + std::string one_tab{"\n\t"}; + std::string two_tabs{"\n\t\t"}; + std::string three_tabs{"\n\t\t\t"}; + std::string four_tabs{"\n\t\t\t\t"}; + std::string five_tabs{"\n\t\t\t\t\t"}; + std::string six_tabs{"\n\t\t\t\t\t\t"}; + + /* + Create HTML file which contains report from coverage. There is no need for + 3rd party HTML libraries as the HTML report file is pretty simple and only + consists of tables and appropriate data. + It is better to open CoverageReport.html file in browser and view page + source, as it is much more clear. + */ + coverage_report << ""; + coverage_report << "" << one_tab << ""; + coverage_report << one_tab << "HIP API Coverage report"; + coverage_report << one_tab << "" << one_tab << ""; + coverage_report << one_tab << "" << one_tab << ""; + coverage_report << two_tabs << ""; + coverage_report << two_tabs << "\n"; + coverage_report << two_tabs << "" << three_tabs << ""; + coverage_report << two_tabs << ""; + + coverage_report << two_tabs << "\n"; + coverage_report << one_tab << "
HIP API Coverage report
" << four_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << five_tabs << ""; + + coverage_report << five_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << five_tabs << ""; + + coverage_report << five_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << five_tabs << ""; + + coverage_report << five_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << five_tabs << ""; + + coverage_report << five_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << five_tabs << ""; + + coverage_report << five_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << five_tabs << ""; + + coverage_report << five_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << five_tabs << ""; + + // Based on the tests coverage percentage, pick a color for displaying it. + std::string font_class; + if (basic_stats.tests_coverage_percentage < basic_stats.getLowCoverageLimit()) { + font_class = "headerCovTableEntryLo"; + } + else if (basic_stats.tests_coverage_percentage < basic_stats.getMediumCoverageLimit()) { + font_class = "headerCovTableEntryMed"; + } + else { + font_class = "headerCovTableEntryHi"; + } + + coverage_report << five_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << six_tabs << ""; + coverage_report << five_tabs << ""; + + coverage_report << five_tabs << ""; + coverage_report << four_tabs << "
Catch2 tests location:" << tests_root_directory << "Value
Source files included:" << hipApiHeaderFile << "Total number of detected HIP APIs:" << basic_stats.total_number_of_apis << "
" << hip_rtc_header_file << "HIP API calls within test source files:" << basic_stats.total_number_of_api_calls << "
Total number of test cases:" << basic_stats.total_number_of_test_cases << "
HIP APIs that are called at least once:" << basic_stats.number_of_called_apis << "
HIP APIs that are not called at all:" << basic_stats.number_of_not_called_apis << "
HIP APIs that are marked as deprecated:" << basic_stats.number_of_deprecated_apis << "
Test coverage by implemented tests for the HIP APIs:" << + std::fixed << std::setprecision(2) << basic_stats.tests_coverage_percentage << "%
"; + coverage_report << three_tabs << "
"; + + // Add info about HIP API Groups. + coverage_report << one_tab << "
"; + coverage_report << one_tab << ""; + coverage_report << two_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << two_tabs << ""; + + coverage_report << two_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << two_tabs << ""; + + /* + Get basic stats for each API Group in HTML format and append it to the main HTML. + Create an HTML page for each API Group for more detailed information, as they are + used as hyperlinks from the main HTML page. + */ + for (auto const& hip_api_group: hip_api_groups) { + coverage_report << hip_api_group.getBasicStatsHTML(); + + std::fstream coverage_module_report; + std::string report_module_file_name{test_modules_directory + "/" + hip_api_group.getName() + ".html"}; + coverage_module_report.open(report_module_file_name, std::ios::out); + coverage_module_report << hip_api_group.createHTMLReport(); + coverage_module_report.close(); + } + + coverage_report << two_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << three_tabs << ""; + coverage_report << two_tabs << ""; + + coverage_report << one_tab << "

ModuleHIP APIsHIP API CallsTest casesCalled APIsNot called APIsDeprecated APIsCoverage
" << basic_stats.total_number_of_apis << "" << basic_stats.total_number_of_api_calls << "" << basic_stats.total_number_of_test_cases << "" << basic_stats.number_of_called_apis << "" << basic_stats.number_of_not_called_apis << "" << basic_stats.number_of_deprecated_apis << "" << + std::fixed << std::setprecision(2) << basic_stats.tests_coverage_percentage << "%
"; + coverage_report << one_tab << "
"; + coverage_report << one_tab << "
"; + + coverage_report << one_tab << ""; + coverage_report << two_tabs << ""; + + time_t now{time(nullptr)}; + std::string date{asctime(gmtime(&now))}; + coverage_report << two_tabs << ""; + coverage_report << one_tab << "
Generated: " << date; + coverage_report << two_tabs << " UTC
"; + coverage_report << one_tab << "
"; + coverage_report << "\n\n"; + + coverage_report.close(); + + // Create HTML report for each API, as they are used as hyperlinks from Groups HTML. + for (auto const& hip_api: hip_apis) { + std::fstream coverage_api_report; + std::string report_api_file_name{test_apis_directory + "/" + hip_api.getName() + ".html"}; + coverage_api_report.open(report_api_file_name, std::ios::out); + coverage_api_report << hip_api.createHTMLReport(); + coverage_api_report.close(); + } + + std::cout << "Generated HTML report file " << findAbsolutePathOfFile(report_file_name) << std::endl; +} diff --git a/utils/coverage/reportGenerators.h b/utils/coverage/reportGenerators.h new file mode 100644 index 0000000000..60bb69a2a9 --- /dev/null +++ b/utils/coverage/reportGenerators.h @@ -0,0 +1,41 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hipAPICoverageUtils.h" + +class BasicAPIStats { + public: + int number_of_called_apis; + int number_of_not_called_apis; + int number_of_deprecated_apis; + int total_number_of_api_calls; + int total_number_of_test_cases; + int total_number_of_apis; + float tests_coverage_percentage; + BasicAPIStats(std::vector& hip_api_groups); + float getLowCoverageLimit() const; + float getMediumCoverageLimit() const; +}; + +void generateXMLReportFiles(std::vector& hip_apis, std::vector& hip_api_groups); +void generateHTMLReportFiles(std::vector& hip_apis, std::vector& hip_api_groups, + std::string tests_root_directory, std::string hipApiHeaderFile, std::string hip_rtc_header_file);