AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad

Merge 'develop' into 'amd-staging'

Explorar el Código 
Change-Id: Icaeb728a9f3664da4c79c9f50323042ed7263acc
Este commit está contenido en:
Jenkins
2023-01-10 00:10:09 +00:00
padre 928ee2a571 829dc229b5
commit bac5b07970
Se han modificado 23 ficheros con 5714 adiciones y 262 borrados
Descargar archivo de parche Descargar archivo de diferencias Expandir todos los archivos Contraer todos los archivos
catch/DoxyfileTests
+2540
Ver fichero
La diferencia del archivo ha sido suprimido porque es demasiado grande Cargar Diff
catch/hipTestMain/config/config_amd_linux_common.json
+4 -1
Ver fichero
@@ -14,6 +14,9 @@
"Unit_hipInit_Negative",
"Unit_hipMemset_Negative_OutOfBoundsPtr",
"Unit_hipDeviceReset_Positive_Basic",
"Unit_hipDeviceReset_Positive_Threaded"
"Unit_hipDeviceReset_Positive_Threaded",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode"
]
}
catch/hipTestMain/config/config_amd_windows_common.json
+4 -1
Ver fichero
@@ -93,6 +93,9 @@
"Note: needs to be enabled when streamPerThread issues are fixed",
"Unit_hipStreamSynchronize_NullStreamAndStreamPerThread",
"Note: intermittent Seg fault failure ",
"Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags"
"Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode"
]
}
catch/include/hip_test_common.hh
+38 -4
Ver fichero
@@ -140,6 +140,31 @@ static void initHipCtx(hipCtx_t* pcontext) {
#define HIP_ARRAY hipArray*
#endif
static inline bool IsGfx11() {
#if defined(HT_NVIDIA)
return false;
#elif defined(HT_AMD)
int device = -1;
hipDeviceProp_t props{};
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&props, device));
// Get GCN Arch Name and compare to check if it is gfx11
std::string arch = std::string(props.gcnArchName);
auto pos = arch.find(":");
if (pos != std::string::npos)
arch = arch.substr(0, pos);
if(arch.size() >= 5)
arch = arch.substr(0,5);
return (arch == std::string("gfx11")) ? true : false;
#else
std::cout<<"Have to be either Nvidia or AMD platform, asserting"<<std::endl;
assert(false);
#endif
}
// Utility Functions
namespace HipTest {
@@ -335,6 +360,14 @@ static __global__ void waitKernel(clock_t offset) {
}
}
static __global__ void waitKernel_gfx11(clock_t offset) {
#if HT_AMD
auto start = wall_clock64();
while ((wall_clock64() - start) < offset) {
}
#endif
}
// helper function used to set the device frequency variable
// estimates the number of clock ticks in 1 second
static size_t findTicksPerSecond() {
@@ -350,9 +383,9 @@ static size_t findTicksPerSecond() {
hipEvent_t start, stop;
HIP_CHECK(hipEventCreate(&start));
HIP_CHECK(hipEventCreate(&stop));
auto waitKernel_used = IsGfx11() ? waitKernel_gfx11 : waitKernel;
// Warmup
hipLaunchKernelGGL(waitKernel, dim3(1), dim3(1), 0, 0, clockTicksPerSecond);
hipLaunchKernelGGL(waitKernel_used, dim3(1), dim3(1), 0, 0, clockTicksPerSecond);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
@@ -360,7 +393,7 @@ static size_t findTicksPerSecond() {
// after 10 attempts the result is likely good enough so just accept it
for (int attempts = 10; attempts > 0; --attempts) {
HIP_CHECK(hipEventRecord(start));
hipLaunchKernelGGL(waitKernel, dim3(1), dim3(1), 0, 0, clockTicksPerSecond);
hipLaunchKernelGGL(waitKernel_used, dim3(1), dim3(1), 0, 0, clockTicksPerSecond);
HIP_CHECK(hipEventRecord(stop));
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipEventSynchronize(stop));
@@ -396,7 +429,8 @@ static inline void runKernelForDuration(std::chrono::milliseconds duration,
// precision so that's acceptable.
static size_t ticksPerSecond = findTicksPerSecond();
const auto millis = duration.count();
hipLaunchKernelGGL(waitKernel, dim3(1), dim3(1), 0, stream, ticksPerSecond * millis / 1000);
auto waitKernel_used = IsGfx11() ? waitKernel_gfx11 : waitKernel;
hipLaunchKernelGGL(waitKernel_used, dim3(1), dim3(1), 0, stream, ticksPerSecond * millis / 1000);
}
} // namespace HipTest
catch/include/hip_test_defgroups.hh
+37
Ver fichero
@@ -0,0 +1,37 @@
/*
Copyright (c) 2021 - 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.
*/
// Test groups are named based on the group names from hip_api_runtime.h, with adding "Test" suffix
/**
* @defgroup CallbackTest Callback Activity APIs
* @{
* This section describes tests for the callback/Activity of HIP runtime API.
* @}
*/
/**
* @defgroup GraphTest Graph Management
* @{
* This section describes the graph management types & functions of HIP runtime API.
* @}
*/
catch/include/utils.hh
+5 -1
Ver fichero
@@ -128,7 +128,11 @@ __global__ void Iota(T* const out, size_t pitch, size_t w, size_t h, size_t d) {
inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) {
int ticks_per_ms = 0;
// Clock rate is in kHz => number of clock ticks in a millisecond
HIP_CHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIPCHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeWallClockRate, 0));
} else {
HIPCHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0));
}
Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms);
HIP_CHECK(hipGetLastError());
}
catch/multiproc/hipMemCoherencyTstMProc.cc
+32 -6
Ver fichero
@@ -49,6 +49,19 @@ __global__ void CoherentTst(int *ptr, int PeakClk) {
}
}
__global__ void CoherentTst_gfx11(int *ptr, int PeakClk) {
#if HT_AMD
// Incrementing the value by 1
int64_t GpuFrq = int64_t(PeakClk) * 1000;
int64_t StrtTck = wall_clock64();
atomicAdd(ptr, 1);
// The following while loop checks the value in ptr for around 3-4 seconds
while ((wall_clock64() - StrtTck) <= (3 * GpuFrq)) {
if (atomicCAS(ptr, 3, 4) == 3) break;
}
#endif
}
__global__ void SquareKrnl(int *ptr) {
// ptr value squared here
*ptr = (*ptr) * (*ptr);
@@ -64,14 +77,27 @@ static void TstCoherency(int *Ptr, bool HmmMem) {
HIP_CHECK(hipStreamCreate(&strm));
// storing value 1 in the memory created above
*Ptr = 1;
// Getting gpu frequency
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
if (!HmmMem) {
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void **>(&Dptr), Ptr,
0));
CoherentTst<<<1, 1, 0, strm>>>(Dptr, peak_clk);
if (IsGfx11()) {
HIPCHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeWallClockRate, 0));
} else {
CoherentTst<<<1, 1, 0, strm>>>(Ptr, peak_clk);
HIPCHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
}
if (!HmmMem) {
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void **>(&Dptr), Ptr, 0));
if (IsGfx11()) {
CoherentTst_gfx11<<<1, 1, 0, strm>>>(Dptr, peak_clk);
} else {
CoherentTst<<<1, 1, 0, strm>>>(Dptr, peak_clk);
}
} else {
if (IsGfx11()) {
CoherentTst_gfx11<<<1, 1, 0, strm>>>(Ptr, peak_clk);
} else {
CoherentTst<<<1, 1, 0, strm>>>(Ptr, peak_clk);
}
}
// looping until the value is 2 for 3 seconds
std::chrono::steady_clock::time_point start =
catch/unit/event/Unit_hipEventQuery.cc
+16 -1
Ver fichero
@@ -34,6 +34,20 @@ __global__ void waitKernel(int clockRate, int seconds) {
}
}
__global__ void waitKernel_gfx11(int clockRate, int seconds) {
#if HT_AMD
auto start = wall_clock64();
auto ms = seconds * 1000;
long long waitTill = clockRate * (long long)ms;
while (1) {
auto end = wall_clock64();
if ((end - start) > waitTill) {
return;
}
}
#endif
}
TEST_CASE("Unit_hipEventQuery_DifferentDevice") {
hipEvent_t event1{}, event2{};
HIP_CHECK(hipEventCreate(&event1));
@@ -54,9 +68,10 @@ TEST_CASE("Unit_hipEventQuery_DifferentDevice") {
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipEventRecord(event1, stream));
auto waitKernel_used = IsGfx11() ? waitKernel_gfx11 : waitKernel;
// Start kernel and wait for 3 seconds
// Make sure you increase this time if you add more tests here
waitKernel<<<1, 1, 0, stream>>>(clockRate, 3);
waitKernel_used<<<1, 1, 0, stream>>>(clockRate, 3);
HIP_CHECK(hipEventRecord(event2, stream));
catch/unit/graph/CMakeLists.txt
+2
Ver fichero
@@ -81,6 +81,8 @@ set(TEST_SRC
hipGraphExecChildGraphNodeSetParams.cc
hipStreamGetCaptureInfo_v2.cc
hipUserObjectCreate.cc
hipGraphDebugDotPrint.cc
hipGraphCloneComplx.cc
)
hip_add_exe_to_target(NAME GraphsTest
catch/unit/graph/hipGraphAddKernelNode.cc
+39 -33
Ver fichero
@@ -6,8 +6,10 @@ 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
@@ -22,7 +24,6 @@ THE SOFTWARE.
/* Test verifies hipGraphAddKernelNode API Negative scenarios.
*/
TEST_CASE("Unit_hipGraphAddKernelNode_Negative") {
constexpr int N = 1024;
size_t NElem{N};
@@ -31,7 +32,6 @@ TEST_CASE("Unit_hipGraphAddKernelNode_Negative") {
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
int *A_d, *B_d, *C_d;
hipGraph_t graph;
hipError_t ret;
hipGraphNode_t kNode;
hipKernelNodeParams kNodeParams{};
std::vector<hipGraphNode_t> dependencies;
@@ -41,61 +41,67 @@ TEST_CASE("Unit_hipGraphAddKernelNode_Negative") {
HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N));
HIP_CHECK(hipGraphCreate(&graph, 0));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void*>(&NElem)};
kNodeParams.func = reinterpret_cast<void*>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void **>(kernelArgs);
kNodeParams.extra = nullptr;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
SECTION("Pass pGraphNode as nullptr") {
ret = hipGraphAddKernelNode(nullptr, graph, nullptr, 0, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphAddKernelNode(nullptr, graph, nullptr, 0, &kNodeParams),
hipErrorInvalidValue);
}
SECTION("Pass Graph as nullptr") {
ret = hipGraphAddKernelNode(&kNode, nullptr, nullptr, 0, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, nullptr, nullptr, 0, &kNodeParams),
hipErrorInvalidValue);
}
SECTION("Pass invalid numDependencies") {
ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 11, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, graph, nullptr, 11, &kNodeParams),
hipErrorInvalidValue);
}
SECTION("Pass invalid numDependencies and valid list for dependencies") {
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams));
dependencies.push_back(kNode);
ret = hipGraphAddKernelNode(&kNode, graph,
dependencies.data(), dependencies.size()+1, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, graph, dependencies.data(),
dependencies.size() + 1, &kNodeParams),
hipErrorInvalidValue);
}
SECTION("Pass NodeParams as nullptr") {
ret = hipGraphAddKernelNode(&kNode, graph,
dependencies.data(), dependencies.size(), nullptr);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(
hipGraphAddKernelNode(&kNode, graph, dependencies.data(), dependencies.size(), nullptr),
hipErrorInvalidValue);
}
SECTION("Pass NodeParams func datamember as nullptr") {
#if HT_NVIDIA // on AMD this returns hipErrorInvalidValue
SECTION("Pass NodeParams func data member as nullptr") {
kNodeParams.func = nullptr;
ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams);
REQUIRE(hipSuccess != ret);
HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams),
hipErrorInvalidDeviceFunction);
}
SECTION("Pass kernelParams datamember as nullptr") {
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
#endif
SECTION("Pass kernelParams data member as nullptr") {
kNodeParams.kernelParams = nullptr;
ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams),
hipErrorInvalidValue);
}
#if HT_AMD
// On Cuda setup this test case getting failed
#if HT_AMD // On Cuda setup this test case getting failed
SECTION("Try adding kernel node after destroy the already created graph") {
kNodeParams.kernelParams = reinterpret_cast<void **>(kernelArgs);
HIP_CHECK(hipGraphDestroy(graph));
ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
hipGraph_t destroyed_graph;
HIP_CHECK(hipGraphCreate(&destroyed_graph, 0));
HIP_CHECK(hipGraphDestroy(destroyed_graph));
HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, destroyed_graph, nullptr, 0, &kNodeParams),
hipErrorInvalidValue);
}
#endif
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
HIP_CHECK(hipFree(C_d));
HIP_CHECK(hipGraphDestroy(graph));
}
catch/unit/graph/hipGraphAddMemcpyNode1D.cc
+43 -1
Ver fichero
@@ -24,7 +24,6 @@ Functional -
Memcpy nodes are added and assigned to default device.
2) Allocate memory on default device(Dev 0), Perform memcpy operation for 1D arrays on Peer device(Dev 1) and
verify the results.
Negative -
1) Pass pGraphNode as nullptr and check if api returns error.
2) When graph is un-initialized argument(skipping graph creation), api should return error code.
@@ -198,3 +197,46 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Negative") {
HIP_CHECK(hipFree(A_h));
HIP_CHECK(hipGraphDestroy(graph));
}
TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Negative_Basic") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h;
hipGraphNode_t memcpy_A, memcpy_B, memcpy_C;
hipError_t ret;
hipGraph_t graph;
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
// Pass memcpy direction as hipMemcpyDeviceToHost and
// source pointer as host and destination pointer as device pointer
ret = hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyDeviceToHost);
REQUIRE(hipErrorInvalidValue == ret);
// Pass memcpy direction as hipMemcpyHostToDevice and
// source pointer as device and destination pointer as host pointer
ret = hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
// Pass memcpy direction as hipMemcpyDeviceToDevice and
// pass source pointer as device and destination pointer as host pointer
ret = hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
// Pass memcpy direction as hipMemcpyDeviceToDevice and
// pass source pointer as host and destination pointer as device pointer
ret = hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_d, C_h,
Nbytes, hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphDestroy(graph));
}
catch/unit/graph/hipGraphCloneComplx.cc
+2124
Ver fichero
La diferencia del archivo ha sido suprimido porque es demasiado grande Cargar Diff
catch/unit/graph/hipGraphDebugDotPrint.cc
+246
Ver fichero
@@ -0,0 +1,246 @@
/*
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 <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#define N 1024
#ifdef __linux__
#include <unistd.h>
#include <fstream>
#include <iterator>
#include <algorithm>
#include <string>
static void deleteFile(const char* fName) {
if ( remove(fName) != 0 ) {
INFO("Error in deleting file -" << fName);
} else {
INFO("Successfully deleted file -" << fName);
}
}
static bool checkFileExists(const char* fName) {
return (access(fName, F_OK) != -1);
}
static int countSubstr(std::string input_str, std::string substr) {
int count = 0;
std::string::size_type srch_pos = 0, cur_pos = 0;
while ((cur_pos = input_str.find(substr, srch_pos)) != std::string::npos) {
count++;
srch_pos = (cur_pos + substr.length());
}
return count;
}
static bool validateDotFile(const char* fName,
std::map<std::string, int> *graphData) {
bool isTestPassed = true;
std::ifstream infile(fName);
std::stringstream buffer;
buffer << infile.rdbuf();
std::map<std::string, int>::iterator it;
for (it = (*graphData).begin(); it != (*graphData).end(); it++) {
if (it->second != countSubstr(buffer.str(), it->first)) {
isTestPassed = false;
break;
}
}
return isTestPassed;
}
static void hipGraphDebugDotPrint_Functional(const char* fName,
unsigned int flag = 0) {
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipStream_t stream;
hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, kNodeAdd, memsetNode;
hipKernelNodeParams kNodeParams{};
int *A_d, *B_d, *C_d, *mem_d;
int *A_h, *B_h, *C_h;
hipGraphExec_t graphExec;
size_t NElem{N};
HIP_CHECK(hipMalloc(&mem_d, Nbytes));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kNodeAdd, graph, nullptr, 0, &kNodeParams));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kNodeAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kNodeAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kNodeAdd, &memcpy_C, 1));
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(mem_d);
memsetParams.value = 7;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
std::map<std::string, int> graphData;
graphData["->"] = 3; // number of edges
graphData["MEMCPY"] = 3;
graphData["HtoD"] = 2;
graphData["DtoH"] = 1;
graphData["MEMSET"] = 1;
if ( flag == hipGraphDebugDotFlagsVerbose ) graphData["KERNEL"] = 1;
HIP_CHECK(hipGraphDebugDotPrint(graph, fName, flag));
REQUIRE(true == checkFileExists(fName));
REQUIRE(true == validateDotFile(fName, &graphData));
deleteFile(fName);
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
// Verify graph execution result
HipTest::checkVectorADD<int>(A_h, B_h, C_h, N);
HIP_CHECK(hipFree(mem_d));
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream));
}
/* Functional Test for API - hipGraphDebugDotPrint
Call hipGraphDebugDotPrint and provice path where to write the DOT file.
Verify that DOT file get created or not for each flag passed. */
TEST_CASE("Unit_hipGraphDebugDotPrint_Functional") {
SECTION("Call with hipGraphDebugDotFlagsVerbose flag") {
hipGraphDebugDotPrint_Functional("./graphDotFileFuncVerbose.dot",
hipGraphDebugDotFlagsVerbose);
}
SECTION("Call with hipGraphDebugDotFlagsKernelNodeParams flag") {
hipGraphDebugDotPrint_Functional("./graphDotFileFuncKernelParams.dot",
hipGraphDebugDotFlagsKernelNodeParams);
}
SECTION("Call with hipGraphDebugDotFlagsMemcpyNodeParams flag") {
hipGraphDebugDotPrint_Functional("./graphDotFileFuncMemcpy.dot",
hipGraphDebugDotFlagsMemcpyNodeParams);
}
SECTION("Call with hipGraphDebugDotFlagsMemsetNodeParams flag") {
hipGraphDebugDotPrint_Functional("./graphDotFileFuncMemset.dot",
hipGraphDebugDotFlagsMemsetNodeParams);
}
SECTION("Call with hipGraphDebugDotFlagsHostNodeParams flag") {
hipGraphDebugDotPrint_Functional("./graphDotFileFuncHost.dot",
hipGraphDebugDotFlagsHostNodeParams);
}
SECTION("Call with hipGraphDebugDotFlagsEventNodeParams flag") {
hipGraphDebugDotPrint_Functional("./graphDotFileFuncEvent.dot",
hipGraphDebugDotFlagsEventNodeParams);
}
SECTION("Call with hipGraphDebugDotFlagsExtSemasSignalNodeParams flag") {
hipGraphDebugDotPrint_Functional("./graphDotFileFuncExtSemasSignal.dot",
hipGraphDebugDotFlagsExtSemasSignalNodeParams);
}
SECTION("Call with hipGraphDebugDotFlagsExtSemasWaitNodeParams flag") {
hipGraphDebugDotPrint_Functional("./graphDotFileFuncExtSemasWait.dot",
hipGraphDebugDotFlagsExtSemasWaitNodeParams);
}
SECTION("Call with hipGraphDebugDotFlagsKernelNodeAttributes flag") {
hipGraphDebugDotPrint_Functional("./graphDotFileFuncKernelNodeAttr.dot",
hipGraphDebugDotFlagsKernelNodeAttributes);
}
}
#endif // __linux__
/**
* Negative Test for API - hipGraphDebugDotPrint Argument Check
1) Pass graph as nullptr.
2) Pass graph as uninitialize structure
3) Pass path for dot file to store as nullptr
4) Pass path for dot file to store as empth path
5) Pass flag as hipGraphDebugDotFlags MIN - 1
6) Pass flag as hipGraphDebugDotFlags MAX + 1
7) Pass flag as INT_MAX
*/
#define DOT_FILE_PATH_NEG "./graphDotFileNeg.dot"
TEST_CASE("Unit_hipGraphDebugDotPrint_Argument_Check") {
hipGraph_t graph;
hipError_t ret;
HIP_CHECK(hipGraphCreate(&graph, 0));
SECTION("Pass graph as nullptr") {
ret = hipGraphDebugDotPrint(nullptr, DOT_FILE_PATH_NEG, 0);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass graph as uninitialize structure") {
hipGraph_t graphT{};
ret = hipGraphDebugDotPrint(graphT, DOT_FILE_PATH_NEG, 0);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass path for dot file to store as nullptr") {
ret = hipGraphDebugDotPrint(graph, nullptr, 0);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass path for dot file to store as empth path") {
ret = hipGraphDebugDotPrint(graph, "", 0);
REQUIRE(hipErrorOperatingSystem == ret);
}
SECTION("Pass flag as hipGraphDebugDotFlags MIN - 1") {
ret = hipGraphDebugDotPrint(graph, DOT_FILE_PATH_NEG,
hipGraphDebugDotFlagsVerbose-1);
REQUIRE(hipSuccess == ret);
}
SECTION("Pass flag as hipGraphDebugDotFlags MAX + 1") {
ret = hipGraphDebugDotPrint(graph, DOT_FILE_PATH_NEG,
hipGraphDebugDotFlagsHandles+1);
REQUIRE(hipSuccess == ret);
}
SECTION("Pass flag as INT_MAX") {
ret = hipGraphDebugDotPrint(graph, DOT_FILE_PATH_NEG, INT_MAX);
REQUIRE(hipSuccess == ret);
}
HIP_CHECK(hipGraphDestroy(graph));
}
catch/unit/graph/hipGraphDestroyNode.cc
+160 -1
Ver fichero
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
@@ -29,12 +29,32 @@ Functional ::
1) Create Node and destroy the node
2) Create graph with dependencies and destroy one of the dependency node
before executing the graph.
3) Create a graph with N nodes and (N-1) dependencies between them as shown
below. Start destroying the nodes in iteration from left. In each iteration
verify the number of nodes and dependencies using hipGraphGetNodes and
hipGraphGetEdges.
Node1-->Node2-->Node3->...................->NodeN
4) Create a graph with N nodes and (N-1) dependencies between them as shown
above. Clone the graph. Start destroying the nodes in iteration from left
in the cloned graph. In each iteration verify the number of nodes and
dependencies using hipGraphGetNodes and hipGraphGetEdges. Once all nodes
in the cloned graph are deleted, verify the number of nodes in the original
graph are intact.
5) Create a graph1 with N nodes and (N-1) dependencies between them as shown
above. Create another empty graph0. Add graph1 as child node to graph0.
Delete the child node in graph0. Verify that the nodes in graph1 are still
intact after deleting the child node using hipGraphGetNodes and hipGraphGetEdges.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#define NUM_OF_DUMMY_NODES 8
static __global__ void dummyKernel() {
return;
}
/* This test covers the negative scenarios of
hipGraphDestroyNode API */
@@ -137,3 +157,142 @@ TEST_CASE("Unit_hipGraphDestroyNode_DestroyDependencyNode") {
HIP_CHECK(hipStreamDestroy(streamForGraph));
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Functional Test to test hipGraphDestroyNode using hipGraphGetNodes
* and hipGraphGetEdges APIs.
*/
TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep") {
hipGraph_t graph;
hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES];
hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES];
HIP_CHECK(hipGraphCreate(&graph, 0));
// Create graph with no dependencies
for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) {
void* kernelArgs[] = {nullptr};
kernelNodeParams[i].func = reinterpret_cast<void *>(dummyKernel);
kernelNodeParams[i].gridDim = dim3(1);
kernelNodeParams[i].blockDim = dim3(1);
kernelNodeParams[i].sharedMemBytes = 0;
kernelNodeParams[i].kernelParams = reinterpret_cast<void**>(kernelArgs);
kernelNodeParams[i].extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr,
0, &kernelNodeParams[i]));
}
// Create dependencies between nodes
for (int i = 1; i < NUM_OF_DUMMY_NODES; i++) {
HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode[i-1],
&kernelnode[i], 1));
}
// Start destroying nodes from 0
size_t numOfNodes = 0, numOfDep = 0;
for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) {
// destroy node i
HIP_CHECK(hipGraphDestroyNode(kernelnode[i]));
HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numOfNodes));
REQUIRE(numOfNodes == (NUM_OF_DUMMY_NODES - i - 1));
HIP_CHECK(hipGraphGetEdges(graph, nullptr, nullptr, &numOfDep));
REQUIRE(numOfDep == (NUM_OF_DUMMY_NODES - i - 2));
}
HIP_CHECK(hipGraphDestroyNode(kernelnode[NUM_OF_DUMMY_NODES-1]));
HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numOfNodes));
REQUIRE(numOfNodes == 0);
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Functional Test to test hipGraphDestroyNode using hipGraphGetNodes
* and hipGraphGetEdges APIs on a cloned graph
*/
TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph") {
hipGraph_t graph, clonedgraph;
hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES];
hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES];
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphCreate(&clonedgraph, 0));
// Create graph with no dependencies
for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) {
void* kernelArgs[] = {nullptr};
kernelNodeParams[i].func = reinterpret_cast<void *>(dummyKernel);
kernelNodeParams[i].gridDim = dim3(1);
kernelNodeParams[i].blockDim = dim3(1);
kernelNodeParams[i].sharedMemBytes = 0;
kernelNodeParams[i].kernelParams = reinterpret_cast<void**>(kernelArgs);
kernelNodeParams[i].extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr,
0, &kernelNodeParams[i]));
}
// Create dependencies between nodes
for (int i = 1; i < NUM_OF_DUMMY_NODES; i++) {
HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode[i-1],
&kernelnode[i], 1));
}
HIP_CHECK(hipGraphClone(&clonedgraph, graph));
// Start destroying nodes from 0 and validate number of nodes in
// cloned graph
size_t numOfNodes = 0, numOfDep = 0;
for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) {
hipGraphNode_t node;
// destroy node i
HIP_CHECK(hipGraphNodeFindInClone(&node, kernelnode[i], clonedgraph));
HIP_CHECK(hipGraphDestroyNode(node));
HIP_CHECK(hipGraphGetNodes(clonedgraph, nullptr, &numOfNodes));
REQUIRE(numOfNodes == (NUM_OF_DUMMY_NODES - i - 1));
HIP_CHECK(hipGraphGetEdges(clonedgraph, nullptr, nullptr, &numOfDep));
REQUIRE(numOfDep == (NUM_OF_DUMMY_NODES - i - 2));
}
// Verify the number of nodes in original graph
numOfNodes = 0;
HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numOfNodes));
REQUIRE(numOfNodes == NUM_OF_DUMMY_NODES);
HIP_CHECK(hipGraphDestroy(clonedgraph));
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Functional Test to test hipGraphDestroyNode on child node using
* hipGraphGetNodes and hipGraphGetEdges APIs on a cloned graph.
*/
TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode") {
hipGraph_t graph0, graph1;
hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES], childGraphNode;
hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES];
HIP_CHECK(hipGraphCreate(&graph0, 0));
HIP_CHECK(hipGraphCreate(&graph1, 0));
// Create graph with no dependencies
for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) {
void* kernelArgs[] = {nullptr};
kernelNodeParams[i].func = reinterpret_cast<void *>(dummyKernel);
kernelNodeParams[i].gridDim = dim3(1);
kernelNodeParams[i].blockDim = dim3(1);
kernelNodeParams[i].sharedMemBytes = 0;
kernelNodeParams[i].kernelParams = reinterpret_cast<void**>(kernelArgs);
kernelNodeParams[i].extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph0, nullptr,
0, &kernelNodeParams[i]));
}
// Create dependencies between nodes
for (int i = 1; i < NUM_OF_DUMMY_NODES; i++) {
HIP_CHECK(hipGraphAddDependencies(graph0, &kernelnode[i-1],
&kernelnode[i], 1));
}
// Create child node and add it to graph1
HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph1,
nullptr, 0, graph0));
// delete the child node from graph1
HIP_CHECK(hipGraphDestroyNode(childGraphNode));
// Start destroying nodes from 0
size_t numOfNodes = 0, numOfDep = 0;
for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) {
// destroy node i
HIP_CHECK(hipGraphDestroyNode(kernelnode[i]));
HIP_CHECK(hipGraphGetNodes(graph0, nullptr, &numOfNodes));
REQUIRE(numOfNodes == (NUM_OF_DUMMY_NODES - i - 1));
HIP_CHECK(hipGraphGetEdges(graph0, nullptr, nullptr, &numOfDep));
REQUIRE(numOfDep == (NUM_OF_DUMMY_NODES - i - 2));
}
HIP_CHECK(hipGraphGetNodes(graph1, nullptr, &numOfNodes));
REQUIRE(numOfNodes == 0);
HIP_CHECK(hipGraphDestroy(graph0));
HIP_CHECK(hipGraphDestroy(graph1));
}
catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc
+17 -2
Ver fichero
@@ -67,6 +67,16 @@ static __global__ void sqr_ker_func(int* a, int* b, int clockrate) {
do { cur = clock64()/clockrate - start;}while (cur < wait_t);
}
static __global__ void sqr_ker_func_gfx11(int* a, int* b, int clockrate) {
#if HT_AMD
int tx = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x;
if (tx < LEN) b[tx] = a[tx]*a[tx];
uint64_t wait_t = DELAY_IN_MS,
start = wall_clock64()/clockrate, cur;
do { cur = wall_clock64()/clockrate - start;}while (cur < wait_t);
#endif
}
/**
* Scenario 1: Test to validate setting different events in executable graph.
*/
@@ -106,10 +116,15 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") {
inp_h, memsize, hipMemcpyHostToDevice));
// Get device clock rate
int clkRate = 0;
HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
} else {
HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
}
// kernel1
auto sqr_ker_func_used = IsGfx11() ? sqr_ker_func_gfx11 : sqr_ker_func;
void* kernelArgs[] = {&inp_d, &out_d, reinterpret_cast<void *>(&clkRate)};
kernelNodeParams1.func = reinterpret_cast<void *>(sqr_ker_func);
kernelNodeParams1.func = reinterpret_cast<void *>(sqr_ker_func_used);
kernelNodeParams1.gridDim = dim3(GRID_DIM);
kernelNodeParams1.blockDim = dim3(BLK_DIM);
kernelNodeParams1.sharedMemBytes = 0;
catch/unit/graph/hipGraphExecKernelNodeSetParams.cc
+68 -65
Ver fichero
@@ -6,25 +6,27 @@ in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
Test Case Scenarios :
Negative -
1) Pass hGraphExec as nullptr and verify api returns error code.
2) Pass node as nullptr and verify api returns error code.
3) Pass NodeParams as un-initialized structure object and verify api returns error code.
4) Pass pNodeParams as nullptr and verify api returns error code.
5) Pass NodeParams:func datamember as nullptr and verify api returns error code.
5) Pass NodeParams:func data member as nullptr and verify api returns error code.
Functional -
1) Instantiate a graph with kernel node, obtain executable graph and update
the kernel node params with set and check it is taking effect.
@@ -39,12 +41,11 @@ Functional -
*/
TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Negative") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
hipGraph_t graph;
hipError_t ret;
hipGraphNode_t memcpyNode, kNode{};
hipGraphNode_t kNode{};
hipKernelNodeParams kNodeParams{};
hipStream_t streamForGraph;
int *A_d, *B_d, *C_d;
@@ -55,57 +56,67 @@ TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Negative") {
HIP_CHECK(hipStreamCreate(&streamForGraph));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void*>(&NElem)};
kNodeParams.func = reinterpret_cast<void*>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams));
hipGraphNode_t empty_node;
HIP_CHECK(hipGraphAddEmptyNode(&empty_node, graph, &kNode, 1));
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
SECTION("Pass hipGraphExec as nullptr") {
ret = hipGraphExecKernelNodeSetParams(nullptr, kNode, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(nullptr, kNode, &kNodeParams),
hipErrorInvalidValue);
}
SECTION("Pass Node as nullptr") {
ret = hipGraphExecKernelNodeSetParams(graphExec, nullptr, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, nullptr, &kNodeParams),
hipErrorInvalidValue);
}
#if HT_AMD
/* NodeParams null check is disabled on Nvedia as
/* NodeParams null check is disabled on Nvidia as
* this call gives SIGSEGV error in CUDA setup */
SECTION("Pass NodeParams as nullptr") {
ret = hipGraphExecKernelNodeSetParams(graphExec, kNode, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, kNode, nullptr),
hipErrorInvalidValue);
}
#endif
/* For below 2 scenarios -
In AMD setup this API return - hipErrorInvalidValue and
In CUDA setup this API return - hipErrorInvalidDeviceFunction
As per Cuda spec API can only return "cudaSuccess, cudaErrorInvalidValue".
*/
SECTION("Pass NodeParams as un-initialized structure object") {
hipKernelNodeParams kNodeParams1{};
ret = hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams1);
REQUIRE(hipSuccess != ret);
}
SECTION("Pass NodeParams func datamember as nullptr") {
#if HT_NVIDIA // on AMD this returns hipErrorInvalidValue
SECTION("Pass NodeParams func data member as nullptr") {
kNodeParams.func = nullptr;
ret = hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams);
REQUIRE(hipSuccess != ret);
HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams),
hipErrorInvalidDeviceFunction);
}
#endif
#if HT_NVIDIA // segfaults on AMD
SECTION("Pass kernelParams data member as nullptr") {
kNodeParams.kernelParams = nullptr;
HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams),
hipErrorInvalidValue);
}
#endif
#if HT_NVIDIA // segfaults on AMD
SECTION("node is not a kernel node") {
HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, empty_node, &kNodeParams),
hipErrorInvalidValue);
}
#endif
SECTION("node is not instantiated") {
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams));
HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams),
hipErrorInvalidValue);
}
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
@@ -114,16 +125,15 @@ TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Negative") {
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
/**
* Functional Test for API Exec Kernel Params
*/
TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Functional") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
hipGraph_t graph;
hipGraphNode_t memcpyNode, kNode;
hipKernelNodeParams kNodeParams{}, kNodeParams1{};
@@ -136,43 +146,36 @@ TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Functional") {
HIP_CHECK(hipStreamCreate(&streamForGraph));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, Nbytes,
hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h, Nbytes,
hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void*>(&NElem)};
kNodeParams.func = reinterpret_cast<void*>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(),
dependencies.size(), &kNodeParams));
HIP_CHECK(
hipGraphAddKernelNode(&kNode, graph, dependencies.data(), dependencies.size(), &kNodeParams));
memset(&kNodeParams1, 0, sizeof(kNodeParams1));
kNodeParams1.func = reinterpret_cast<void *>(HipTest::vectorSUB<int>);
kNodeParams1.func = reinterpret_cast<void*>(HipTest::vectorSUB<int>);
kNodeParams1.gridDim = dim3(blocks);
kNodeParams1.blockDim = dim3(threadsPerBlock);
kNodeParams1.sharedMemBytes = 0;
kNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams1.extra = nullptr;
dependencies.clear();
dependencies.push_back(kNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(),
dependencies.size(), C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), dependencies.size(),
C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
REQUIRE(hipSuccess == hipGraphExecKernelNodeSetParams(graphExec, kNode,
&kNodeParams1));
HIP_CHECK(hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams1));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
catch/unit/graph/hipGraphKernelNodeGetParams.cc
+35 -40
Ver fichero
@@ -6,8 +6,10 @@ 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
@@ -18,7 +20,7 @@ THE SOFTWARE.
*/
/**
Testcase Scenarios :
Test Case Scenarios :
Negative -
1) Pass node as nullptr and verify api returns error code.
2) Pass pNodeParams as nullptr and verify api returns error code.
@@ -36,40 +38,42 @@ Functional -
/* Test verifies hipGraphKernelNodeGetParams API Negative scenarios.
*/
TEST_CASE("Unit_hipGraphKernelNodeGetParams_Negative") {
constexpr int N = 1024;
size_t NElem{N};
int *A_d, *B_d, *C_d;
hipError_t ret;
hipGraph_t graph;
hipGraphNode_t kNode;
hipKernelNodeParams kNodeParams{};
HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N));
HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N));
HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N));
HIP_CHECK(hipGraphCreate(&graph, 0));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void*>(&NElem)};
kNodeParams.func = reinterpret_cast<void*>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1);
kNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void **>(kernelArgs);
kNodeParams.extra = nullptr;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams));
SECTION("Pass node as nullptr") {
ret = hipGraphKernelNodeGetParams(nullptr, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphKernelNodeGetParams(nullptr, &kNodeParams), hipErrorInvalidValue);
}
SECTION("Pass kNodeParams as nullptr") {
ret = hipGraphKernelNodeGetParams(kNode, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphKernelNodeGetParams(kNode, nullptr), hipErrorInvalidValue);
}
#if HT_NVIDIA // segfaults on AMD
SECTION("node is not a kernel node") {
hipGraphNode_t empty_node;
HIP_CHECK(hipGraphAddEmptyNode(&empty_node, graph, nullptr, 0));
HIP_CHECK_ERROR(hipGraphKernelNodeGetParams(empty_node, &kNodeParams), hipErrorInvalidValue);
}
#endif
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
HIP_CHECK(hipFree(C_d));
@@ -83,28 +87,20 @@ static bool dim3_compare(dim3 node1, dim3 node2) {
return false;
}
static bool kernelParam_compare(void **p1, void ** p2) {
static bool kernelParam_compare(void** p1, void** p2) {
for (int i = 0; i < 4; i++) {
if (*reinterpret_cast<int *>(p1[i]) != *reinterpret_cast<int *>(p2[i]))
return false;
if (*reinterpret_cast<int*>(p1[i]) != *reinterpret_cast<int*>(p2[i])) return false;
}
return true;
}
static bool node_compare(hipKernelNodeParams *kNode1,
hipKernelNodeParams *kNode2) {
if (!dim3_compare(kNode1->blockDim, kNode2->blockDim))
return false;
if (kNode1->extra != kNode2->extra)
return false;
if (kNode1->func != kNode2->func)
return false;
if (!dim3_compare(kNode1->gridDim, kNode2->gridDim))
return false;
if (!kernelParam_compare(kNode1->kernelParams, kNode2->kernelParams))
return false;
if (kNode1->sharedMemBytes != kNode2->sharedMemBytes)
return false;
static bool node_compare(hipKernelNodeParams* kNode1, hipKernelNodeParams* kNode2) {
if (!dim3_compare(kNode1->blockDim, kNode2->blockDim)) return false;
if (kNode1->extra != kNode2->extra) return false;
if (kNode1->func != kNode2->func) return false;
if (!dim3_compare(kNode1->gridDim, kNode2->gridDim)) return false;
if (!kernelParam_compare(kNode1->kernelParams, kNode2->kernelParams)) return false;
if (kNode1->sharedMemBytes != kNode2->sharedMemBytes) return false;
return true;
}
@@ -121,37 +117,36 @@ TEST_CASE("Unit_hipGraphKernelNodeGetParams_Functional") {
HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N));
HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N));
HIP_CHECK(hipGraphCreate(&graph, 0));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void*>(&NElem)};
kNodeParams.func = reinterpret_cast<void*>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1);
kNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void **>(kernelArgs);
kNodeParams.extra = nullptr;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams));
SECTION("Get Kernel Param and verify.") {
hipKernelNodeParams kNodeGetParams;
HIP_CHECK(hipGraphKernelNodeGetParams(kNode, &kNodeGetParams));
REQUIRE(true == node_compare(&kNodeParams, &kNodeGetParams));
REQUIRE(node_compare(&kNodeParams, &kNodeGetParams));
}
SECTION("Set kernel node params then Get Kernel Param and verify.") {
hipKernelNodeParams kNodeParams1;
kNodeParams1.func =
reinterpret_cast<void *>(HipTest::vectorADDReverse<int>);
kNodeParams1.func = reinterpret_cast<void*>(HipTest::vectorADDReverse<int>);
kNodeParams1.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1);
kNodeParams1.blockDim = dim3(THREADS_PER_BLOCK, 1, 1);
kNodeParams1.sharedMemBytes = 0;
kNodeParams1.kernelParams = reinterpret_cast<void **>(kernelArgs);
kNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams1.extra = nullptr;
HIP_CHECK(hipGraphKernelNodeSetParams(kNode, &kNodeParams1));
hipKernelNodeParams kNodeGetParams1;
HIP_CHECK(hipGraphKernelNodeSetParams(kNode, &kNodeParams1));
HIP_CHECK(hipGraphKernelNodeGetParams(kNode, &kNodeGetParams1));
REQUIRE(true == node_compare(&kNodeParams1, &kNodeGetParams1));
REQUIRE(node_compare(&kNodeParams1, &kNodeGetParams1));
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
HIP_CHECK(hipFree(C_d));
catch/unit/graph/hipGraphKernelNodeSetParams.cc
+73 -70
Ver fichero
@@ -6,19 +6,21 @@ in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
Test Case Scenarios :
Negative -
1) Pass node as nullptr and verify api returns error code.
2) Pass pNodeParams as nullptr and verify api returns error code.
@@ -30,13 +32,12 @@ Functional -
hipGraphKernelNodeSetParams, finally check taking effect after launching graph.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
/* Test verifies hipGraphKernelNodeSetParams API Negative scenarios.
*/
TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") {
constexpr int N = 1024;
size_t NElem{N};
@@ -44,35 +45,53 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") {
constexpr auto threadsPerBlock = 256;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
int *A_d, *B_d, *C_d;
hipError_t ret;
hipGraph_t graph;
hipGraphNode_t kNode;
hipKernelNodeParams kNodeParams{};
HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N));
HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N));
HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N));
HIP_CHECK(hipGraphCreate(&graph, 0));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void*>(&NElem)};
kNodeParams.func = reinterpret_cast<void*>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void **>(kernelArgs);
kNodeParams.extra = nullptr;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams));
SECTION("Pass node as nullptr") {
ret = hipGraphKernelNodeSetParams(nullptr, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(nullptr, &kNodeParams), hipErrorInvalidValue);
}
SECTION("Pass kNodeParams as nullptr") {
ret = hipGraphKernelNodeSetParams(kNode, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(kNode, nullptr), hipErrorInvalidValue);
}
#if HT_NVIDIA // on AMD this returns hipErrorInvalidValue
SECTION("Pass NodeParams func data member as nullptr") {
kNodeParams.func = nullptr;
HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(kNode, &kNodeParams),
hipErrorInvalidDeviceFunction);
}
#endif
#if HT_NVIDIA // segfaults on AMD
SECTION("Pass kernelParams data member as nullptr") {
kNodeParams.kernelParams = nullptr;
HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(kNode, &kNodeParams), hipErrorInvalidValue);
}
#endif
#if HT_NVIDIA // segfaults on AMD
SECTION("node is not a kernel node") {
hipGraphNode_t empty_node;
HIP_CHECK(hipGraphAddEmptyNode(&empty_node, graph, nullptr, 0));
HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(empty_node, &kNodeParams), hipErrorInvalidValue);
}
#endif
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
HIP_CHECK(hipFree(C_d));
@@ -82,12 +101,12 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") {
/**
* Functional Test for API Set Kernel Params
*/
TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
hipGraph_t graph;
hipGraphNode_t memcpyNode, kNode;
hipKernelNodeParams kNodeParams{}, kNodeParams1{};
@@ -100,39 +119,34 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") {
HIP_CHECK(hipStreamCreate(&streamForGraph));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, Nbytes,
hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h, Nbytes,
hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void*>(&NElem)};
kNodeParams.func = reinterpret_cast<void*>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(),
dependencies.size(), &kNodeParams));
HIP_CHECK(
hipGraphAddKernelNode(&kNode, graph, dependencies.data(), dependencies.size(), &kNodeParams));
kNodeParams1.func = reinterpret_cast<void *>(HipTest::vectorSUB<int>);
kNodeParams1.func = reinterpret_cast<void*>(HipTest::vectorSUB<int>);
kNodeParams1.gridDim = dim3(blocks);
kNodeParams1.blockDim = dim3(threadsPerBlock);
kNodeParams1.sharedMemBytes = 0;
kNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams1.extra = nullptr;
HIP_CHECK(hipGraphKernelNodeSetParams(kNode, &kNodeParams1));
dependencies.clear();
dependencies.push_back(kNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(),
dependencies.size(), C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), dependencies.size(),
C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
@@ -147,12 +161,12 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") {
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
static __global__ void ker_vec_add(int *A, int *B) {
static __global__ void ker_vec_add(int* A, int* B) {
int i = threadIdx.x + blockDim.x * blockIdx.x;
A[i] = A[i] + B[i];
}
static __global__ void ker_vec_sub(int *A, int *B) {
static __global__ void ker_vec_sub(int* A, int* B) {
int i = threadIdx.x + blockDim.x * blockIdx.x;
A[i] = A[i] - B[i];
}
@@ -167,7 +181,7 @@ class GraphKernelNodeGetSetParam {
const int blocks = (N / threadsPerBlock);
hipGraphNode_t memcpyH2D_A1, memcpyH2D_A2, memcpyD2H_A3, vec_maths;
hipGraph_t graph;
hipKernelNodeParams kerNodeParams { };
hipKernelNodeParams kerNodeParams{};
int *A1_d, *A2_d, *A1_h, *A2_h, *A3_h;
public:
@@ -179,32 +193,26 @@ class GraphKernelNodeGetSetParam {
HIP_CHECK(hipMalloc(&A2_d, Nbytes));
// Allocate host buffers
A1_h = reinterpret_cast<int*>(malloc(Nbytes));
REQUIRE(A1_h != NULL);
REQUIRE(A1_h != nullptr);
A2_h = reinterpret_cast<int*>(malloc(Nbytes));
REQUIRE(A2_h != NULL);
REQUIRE(A2_h != nullptr);
A3_h = reinterpret_cast<int*>(malloc(Nbytes));
REQUIRE(A3_h != NULL);
REQUIRE(A3_h != nullptr);
// Create all the 3 level graphs
HIP_CHECK(hipGraphCreate(&graph, 0));
void *kernelArgs[] = { &A1_d, &A2_d };
void* kernelArgs[] = {&A1_d, &A2_d};
kerNodeParams.func = reinterpret_cast<void*>(ker_vec_add);
kerNodeParams.gridDim = dim3(blocks);
kerNodeParams.blockDim = dim3(threadsPerBlock);
kerNodeParams.sharedMemBytes = 0;
kerNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kerNodeParams.extra = nullptr;
HIP_CHECK(
hipGraphAddKernelNode(&vec_maths, graph, nullptr, 0, &kerNodeParams));
HIP_CHECK(hipGraphAddKernelNode(&vec_maths, graph, nullptr, 0, &kerNodeParams));
// Add nodes to graph
HIP_CHECK(
hipGraphAddMemcpyNode1D(&memcpyH2D_A1, graph, nullptr, 0, A1_d, A1_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(
hipGraphAddMemcpyNode1D(&memcpyH2D_A2, graph, nullptr, 0, A2_d, A2_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(
hipGraphAddMemcpyNode1D(&memcpyD2H_A3, graph, nullptr, 0, A3_h, A1_d,
Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A1, graph, nullptr, 0, A1_d, A1_h, Nbytes,
hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A2, graph, nullptr, 0, A2_d, A2_h, Nbytes,
hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_A3, graph, nullptr, 0, A3_h, A1_d, Nbytes,
hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A1, &vec_maths, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A2, &vec_maths, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &vec_maths, &memcpyD2H_A3, 1));
@@ -213,20 +221,18 @@ class GraphKernelNodeGetSetParam {
// Fill Random Input Data
void fillRandInpData() {
for (int i = 0; i < N; i++) {
A1_h[i] = (rand() % 256); //NOLINT
A2_h[i] = (rand() % 256); //NOLINT
A1_h[i] = (rand() % 256); // NOLINT
A2_h[i] = (rand() % 256); // NOLINT
}
}
hipGraph_t* getRootGraph() {
return &graph;
}
hipGraph_t* getRootGraph() { return &graph; }
void updateNode() {
size_t numNodes = 0;
HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes));
hipGraphNode_t *nodes = reinterpret_cast<hipGraphNode_t*>(malloc(
numNodes * sizeof(hipGraphNode_t)));
hipGraphNode_t* nodes =
reinterpret_cast<hipGraphNode_t*>(malloc(numNodes * sizeof(hipGraphNode_t)));
HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes));
// Get the Graph node from the embedded graph
size_t nodeIdx = 0;
@@ -246,9 +252,7 @@ class GraphKernelNodeGetSetParam {
}
// Function to validate result
void validateOutData() {
HipTest::checkVectorSUB<int>(A1_h, A2_h, A3_h, N);
}
void validateOutData() { HipTest::checkVectorSUB<int>(A1_h, A2_h, A3_h, N); }
// Destroy resources
~GraphKernelNodeGetSetParam() {
@@ -263,7 +267,7 @@ class GraphKernelNodeGetSetParam {
};
TEST_CASE("Unit_hipGraphKernelNodeGetSetParams_Functional") {
hipGraph_t *graph;
hipGraph_t* graph;
hipStream_t streamForGraph;
hipGraphExec_t graphExec;
GraphKernelNodeGetSetParam GraphKernelNodeGetSetParamObj;
@@ -271,8 +275,7 @@ TEST_CASE("Unit_hipGraphKernelNodeGetSetParams_Functional") {
GraphKernelNodeGetSetParamObj.updateNode();
HIP_CHECK(hipStreamCreate(&streamForGraph));
// Instantiate and launch the childgraph
HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr,
nullptr, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr, nullptr, 0));
GraphKernelNodeGetSetParamObj.fillRandInpData();
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
catch/unit/memory/hipMemPoolApi.cc
+93 -17
Ver fichero
@@ -123,6 +123,21 @@ __global__ void kernel500ms(float* hostRes, int clkRate) {
}
}
__global__ void kernel500ms_gfx11(float* hostRes, int clkRate) {
#if HT_AMD
int tid = threadIdx.x + blockIdx.x * blockDim.x;
hostRes[tid] = tid + 1;
__threadfence_system();
// expecting that the data is getting flushed to host here!
uint64_t start = wall_clock64()/clkRate, cur;
if (clkRate > 1) {
do { cur = wall_clock64()/clkRate-start;}while (cur < wait_ms);
} else {
do { cur = wall_clock64()/start;}while (cur < wait_ms);
}
#endif
}
TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") {
int mem_pool_support = 0;
HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
@@ -147,9 +162,14 @@ TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") {
int blocks = 1024;
int clkRate;
hipMemPoolAttr attr;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream));
@@ -229,9 +249,14 @@ TEST_CASE("Unit_hipMemPoolApi_BasicTrim") {
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
hipMemPoolAttr attr;
attr = hipMemPoolAttrReleaseThreshold;
@@ -312,9 +337,15 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") {
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
}
hipMemPoolAttr attr;
// Not a real free, since kernel isn't done
@@ -329,7 +360,11 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") {
HIP_CHECK(hipStreamSynchronize(stream));
// Second kernel launch with new memory
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipStreamSynchronize(stream));
@@ -369,7 +404,11 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
hipMemPoolAttr attr;
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
} else {
HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
}
float *A, *B, *C;
hipStream_t stream, stream2;
@@ -395,7 +434,11 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
// Run kernel for 500 ms in the first stream
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
@@ -410,7 +453,11 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
REQUIRE(A != B);
// Run kernel with the new memory in the second stream
kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamSynchronize(stream2));
@@ -428,7 +475,13 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
// Run kernel for 500 ms in the first stream
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
@@ -443,7 +496,11 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
REQUIRE(A == B);
// Run kernel with the new memory in the second stream
kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamSynchronize(stream2));
@@ -461,7 +518,12 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
// Run kernel for 500 ms in the first stream
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
@@ -473,7 +535,11 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
REQUIRE(A != B);
// Run kernel with the new memory in the second stream
kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamSynchronize(stream2));
@@ -510,9 +576,15 @@ TEST_CASE("Unit_hipMemPoolApi_Default") {
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
}
hipMemPoolAttr attr;
// Not a real free, since kernel isn't done
@@ -527,7 +599,11 @@ TEST_CASE("Unit_hipMemPoolApi_Default") {
HIP_CHECK(hipStreamSynchronize(stream));
// Second kernel launch with new memory
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream));
catch/unit/memory/hipMemoryAllocateCoherent.cc
+22 -2
Ver fichero
@@ -41,6 +41,21 @@ __global__ void Kernel(float* hostRes, int clkRate) {
}
}
__global__ void Kernel_gfx11(float* hostRes, int clkRate) {
#if HT_AMD
int tid = threadIdx.x + blockIdx.x * blockDim.x;
hostRes[tid] = tid + 1;
__threadfence_system();
// expecting that the data is getting flushed to host here!
uint64_t start = wall_clock64()/clkRate, cur;
if (clkRate > 1) {
do { cur = wall_clock64()/clkRate-start;}while (cur < wait_sec);
} else {
do { cur = wall_clock64()/start;}while (cur < wait_sec);
}
#endif
}
TEST_CASE("Unit_hipHostMalloc_CoherentAccess") {
int blocks = 2;
float* hostRes;
@@ -49,9 +64,14 @@ TEST_CASE("Unit_hipHostMalloc_CoherentAccess") {
hostRes[0] = 0;
hostRes[1] = 0;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
} else {
HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
}
std::cout << clkRate << std::endl;
hipLaunchKernelGGL(HIP_KERNEL_NAME(Kernel), dim3(1), dim3(blocks),
auto Kernel_used = IsGfx11() ? Kernel_gfx11 : Kernel;
hipLaunchKernelGGL(HIP_KERNEL_NAME(Kernel_used), dim3(1), dim3(blocks),
0, 0, hostRes, clkRate);
HIP_CHECK(hipGetLastError());
int eleCounter = 0;
catch/unit/stream/hipStreamACb_MultiThread.cc
+20 -1
Ver fichero
@@ -53,6 +53,24 @@ static __global__ void device_function(float* C_d, float* A_d, size_t Num) {
}
}
static __global__ void device_function_gfx11(float* C_d, float* A_d, size_t Num) {
#if HT_AMD
size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for (size_t i = gputhread; i < Num; i += stride) {
C_d[i] = A_d[i] * A_d[i];
}
// Delay thread 1 only in the GPU
if (gputhread == 1) {
uint64_t wait_t = 3200000000, start = wall_clock64(), cur;
do {
cur = wall_clock64() - start;
} while (cur < wait_t);
}
#endif
}
static void HIPRT_CB Thread1_Callback(hipStream_t stream, hipError_t status,
void* userData) {
@@ -128,7 +146,8 @@ TEST_CASE("Unit_hipStreamAddCallback_MultipleThreads") {
constexpr unsigned threadsPerBlock = 256;
constexpr unsigned blocks = (N + 255)/threadsPerBlock;
hipLaunchKernelGGL((device_function), dim3(blocks),
auto device_function_used = IsGfx11() ? device_function_gfx11 : device_function;
hipLaunchKernelGGL((device_function_used), dim3(blocks),
dim3(threadsPerBlock), 0,
mystream, C_d, A_d, N);
HIP_CHECK(hipGetLastError());
catch/unit/stream/hipStreamWaitEvent.cc
+19 -4
Ver fichero
@@ -94,6 +94,20 @@ __global__ void waitKernel(int clockRate, int seconds) {
}
}
__global__ void waitKernel_gfx11(int clockRate, int seconds) {
#if HT_AMD
auto start = wall_clock64();
auto ms = seconds * 1000;
long long waitTill = clockRate * (long long)ms;
while (1) {
auto end = wall_clock64();
if ((end - start) > waitTill) {
return;
}
}
#endif
}
TEST_CASE("Unit_hipStreamWaitEvent_Default") {
hipStream_t stream{nullptr};
hipEvent_t waitEvent{nullptr};
@@ -111,7 +125,8 @@ TEST_CASE("Unit_hipStreamWaitEvent_Default") {
HIP_CHECK(hipGetDeviceProperties(&prop, deviceId));
auto clockRate = prop.clockRate;
waitKernel<<<1, 1, 0, stream>>>(clockRate, 2); // Wait for 2 seconds
auto waitKernel_used = IsGfx11() ? waitKernel_gfx11 : waitKernel;
waitKernel_used<<<1, 1, 0, stream>>>(clockRate, 2); // Wait for 2 seconds
HIP_CHECK(hipEventRecord(waitEvent, stream));
@@ -145,8 +160,8 @@ TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") {
hipDeviceProp_t prop{};
HIP_CHECK(hipGetDeviceProperties(&prop, deviceId));
auto clockRate = prop.clockRate;
waitKernel<<<1, 1, 0, blockedStreamA>>>(clockRate,
auto waitKernel_used = IsGfx11() ? waitKernel_gfx11 : waitKernel;
waitKernel_used<<<1, 1, 0, blockedStreamA>>>(clockRate,
3); // wait for 3 seconds
HIP_CHECK(hipEventRecord(waitEvent, blockedStreamA));
@@ -155,7 +170,7 @@ TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") {
HIP_CHECK(hipStreamWaitEvent(streamBlockedOnStreamA, waitEvent, 0));
waitKernel<<<1, 1, 0, streamBlockedOnStreamA>>>(clockRate, 2); // Wait for 2 seconds
waitKernel_used<<<1, 1, 0, streamBlockedOnStreamA>>>(clockRate, 2); // Wait for 2 seconds
HIP_CHECK(hipStreamSynchronize(unblockingStream));
catch/unit/streamperthread/hipStreamPerThrdTsts.cc
+77 -12
Ver fichero
@@ -95,6 +95,39 @@ __global__ void StreamPerThrd1(int *A, int Pk_Clk) {
*A = 1;
}
__global__ void StreamPerThrd_gfx11(int *Ad, int *Ad1, size_t n, int Pk_Clk,
int Wait, int WaitEvnt = 0) {
#if HT_AMD
size_t index = blockIdx.x * blockDim.x + threadIdx.x;
if (index < n) {
Ad[index] = Ad[index] + 10;
}
if (Wait) {
int64_t GpuFrq = (Pk_Clk * 1000);
int64_t StrtTck = wall_clock64();
if (index == 0) {
// The following while loop checks the value in ptr for around 4 seconds
while ((wall_clock64() - StrtTck) <= (6 * GpuFrq)) {
}
if (WaitEvnt == 1) {
*Ad1 = 1;
}
}
}
#endif
}
__global__ void StreamPerThrd1_gfx11(int *A, int Pk_Clk) {
#if HT_AMD
int64_t GpuFrq = (Pk_Clk * 1000);
int64_t StrtTck = wall_clock64();
// The following while loop checks the value in ptr for around 3-4 seconds
while ((wall_clock64() - StrtTck) <= (3 * GpuFrq)) {
}
*A = 1;
#endif
}
__global__ void MiniKernel(int *A) {
if (*A == 0) {
*A = 2; // Fail condition
@@ -189,12 +222,18 @@ static void EventSync() {
HIP_CHECK(hipEventCreate(&start));
HIP_CHECK(hipEventCreate(&end));
HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice));
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
dim3 dimBlock(blockSize, 1, 1);
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
HIP_CHECK(hipEventRecord(start, hipStreamPerThread));
StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL, NumElms,
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeWallClockRate, 0));
StreamPerThrd_gfx11<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL, NumElms,
peak_clk, 0);
} else {
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL, NumElms,
peak_clk, 0);
}
HIP_CHECK(hipEventRecord(end, hipStreamPerThread));
HIP_CHECK(hipEventSynchronize(end));
HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost));
@@ -226,12 +265,18 @@ TEST_CASE("Unit_hipStreamPerThreadTst_StrmQuery") {
Ah[i] = CONST_NUM;
}
HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice));
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
dim3 dimBlock(blockSize, 1, 1);
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
SECTION("Test working of hipStreamQuery") {
StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL,
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeWallClockRate, 0));
StreamPerThrd_gfx11<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL,
NumElms, peak_clk, 1);
} else {
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL,
NumElms, peak_clk, 1);
}
err = hipStreamQuery(hipStreamPerThread);
if (err != hipErrorNotReady) {
WARN("hipStreamQuery on hipStreamPerThread didnt return expected error!");
@@ -245,7 +290,11 @@ TEST_CASE("Unit_hipStreamPerThreadTst_StrmQuery") {
HIP_CHECK(hipHostMalloc(&Hptr, sizeof(int)));
*Hptr = 0;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d), Hptr, 0));
StreamPerThrd1<<<1, 1, 0, hipStreamPerThread>>>(A_d, peak_clk);
if (IsGfx11()) {
StreamPerThrd1_gfx11<<<1, 1, 0, hipStreamPerThread>>>(A_d, peak_clk);
} else {
StreamPerThrd1<<<1, 1, 0, hipStreamPerThread>>>(A_d, peak_clk);
}
HIP_CHECK(hipStreamAddCallback(hipStreamPerThread, CallBackFunctn, A_d, 0));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
HIP_CHECK(hipHostFree(Hptr));
@@ -277,11 +326,17 @@ TEST_CASE("Unit_hipStreamPerThread_MangdMem") {
hipStreamPerThread));
}
int peak_clk;
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
dim3 dimBlock(blockSize, 1, 1);
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Hmm, NULL,
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeWallClockRate, 0));
StreamPerThrd_gfx11<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Hmm, NULL,
NumElms, peak_clk, 0);
} else {
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Hmm, NULL,
NumElms, peak_clk, 0);
}
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
// Validating the result
int MisMatch = 0;
@@ -313,11 +368,17 @@ TEST_CASE("Unit_hipStreamPerThread_ChildProc") {
Ah[i] = CONST_NUM;
}
HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice));
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
dim3 dimBlock(blockSize, 1, 1);
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL,
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeWallClockRate, 0));
StreamPerThrd_gfx11<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL,
NumElms, peak_clk, 0);
} else{
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL,
NumElms, peak_clk, 0);
}
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost));
int MisMatch = 0;
@@ -380,13 +441,17 @@ TEST_CASE("Unit_hipStreamPerThread_StrmWaitEvt") {
HIP_CHECK(hipMalloc(&Ad1, sizeof(int)));
HIP_CHECK(hipMemset(Ad1, 0, sizeof(int)));
int peak_clk;
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
dim3 dimBlock(blockSize, 1, 1);
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
hipEvent_t e1;
HIPCHECK(hipEventCreate(&e1));
StreamPerThrd<<<dimGrid, dimBlock, 0, Strm>>>(Ad, Ad1, NumElms,
peak_clk, 1, 1);
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeWallClockRate, 0));
StreamPerThrd_gfx11<<<dimGrid, dimBlock, 0, Strm>>>(Ad, Ad1, NumElms, peak_clk, 1, 1);
} else {
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
StreamPerThrd<<<dimGrid, dimBlock, 0, Strm>>>(Ad, Ad1, NumElms, peak_clk, 1, 1);
}
HIP_CHECK(hipEventRecord(e1, Strm));
HIP_CHECK(hipStreamWaitEvent(hipStreamPerThread, e1, 0 /*flags*/));
MiniKernel<<<1, 1, 0, hipStreamPerThread>>>(Ad1);