diff --git a/bin/hipcc.bat b/bin/hipcc.bat index a19ce9d80f..7a3c25ba8d 100644 --- a/bin/hipcc.bat +++ b/bin/hipcc.bat @@ -1,2 +1,2 @@ -@IF DEFINED HIP_PATH (set HIPCC="%HIP_PATH%/bin/hipcc") ELSE (set HIPCC="%~dp0/hipcc") +set HIPCC="%~dp0/hipcc" @perl %HIPCC% %* diff --git a/bin/hipconfig.bat b/bin/hipconfig.bat index 9ba2d4205c..47d9500c71 100644 --- a/bin/hipconfig.bat +++ b/bin/hipconfig.bat @@ -1,2 +1,2 @@ -@IF DEFINED HIP_PATH (set HIPCONFIG="%HIP_PATH%/bin/hipconfig") ELSE (set HIPCONFIG="%~dp0/hipconfig") +set HIPCONFIG="%~dp0/hipconfig" @perl %HIPCONFIG% %* diff --git a/cmake/FindHIP.cmake b/cmake/FindHIP.cmake index 75806e1b5f..d56b20f26e 100644 --- a/cmake/FindHIP.cmake +++ b/cmake/FindHIP.cmake @@ -50,8 +50,8 @@ mark_as_advanced(HIP_HOST_COMPILATION_CPP) get_filename_component(_IMPORT_PREFIX "${CMAKE_CURRENT_LIST_DIR}/../" REALPATH) -# HIP is supported on Linux only -if(UNIX AND NOT APPLE AND NOT CYGWIN) +# HIP is currently not supported for apple +if(NOT APPLE) # Search for HIP installation if(NOT HIP_ROOT_DIR) # Search in user specified path first @@ -94,7 +94,6 @@ if(UNIX AND NOT APPLE AND NOT CYGWIN) # Now search in default paths find_program(HIP_HIPCC_EXECUTABLE hipcc) endif() - mark_as_advanced(HIP_HIPCC_EXECUTABLE) # Find HIPCONFIG executable find_program( @@ -113,7 +112,12 @@ if(UNIX AND NOT APPLE AND NOT CYGWIN) # Now search in default paths find_program(HIP_HIPCONFIG_EXECUTABLE hipconfig) endif() + if(NOT UNIX) + set(HIP_HIPCONFIG_EXECUTABLE "${HIP_HIPCONFIG_EXECUTABLE}.bat") + set(HIP_HIPCC_EXECUTABLE "${HIP_HIPCC_EXECUTABLE}.bat") + endif() mark_as_advanced(HIP_HIPCONFIG_EXECUTABLE) + mark_as_advanced(HIP_HIPCC_EXECUTABLE) # Find HIPCC_CMAKE_LINKER_HELPER executable find_program( diff --git a/include/hip/hip_bfloat16.h b/include/hip/hip_bfloat16.h index 642090b4f5..eb99206632 100644 --- a/include/hip/hip_bfloat16.h +++ b/include/hip/hip_bfloat16.h @@ -56,7 +56,7 @@ struct hip_bfloat16 enum truncate_t { - truncate + truncate_0 }; __host__ __device__ hip_bfloat16() = default; diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 7d4d58ec24..f2a24e6dbe 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -445,6 +445,8 @@ typedef enum hipDeviceAttribute_t { ///< hipStreamWaitValue64(), '0' otherwise. hipDeviceAttributeImageSupport, ///< '1' if Device supports image, '0' otherwise. + hipDeviceAttributeMultiprocessorBoostCount, ///< All available boost compute units for the device + hipDeviceAttributeAmdSpecificEnd = 19999, hipDeviceAttributeVendorSpecificBegin = 20000, // Extended attributes for vendors @@ -4302,6 +4304,11 @@ typedef enum hipStreamUpdateCaptureDependenciesFlags { hipStreamSetCaptureDependencies, ///< Replace the dependency set with the new nodes } hipStreamUpdateCaptureDependenciesFlags; +typedef enum hipGraphInstantiateFlags { + hipGraphInstantiateFlagAutoFreeOnLaunch = + 1, ///< Automatically free memory allocated in a graph before relaunching. +} hipGraphInstantiateFlags; + /** * @brief Begins graph capture on a stream. * @@ -4399,6 +4406,31 @@ hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t size_t numDependencies, unsigned int flags __dparm(0)); +/** + * @brief Enqueues a host function call in a stream. + * + * @param [in] stream - stream to enqueue work to. + * @param [in] fn - function to call once operations enqueued preceeding are complete. + * @param [in] userData - User-specified data to be passed to the function. + * @returns #hipSuccess, #hipErrorInvalidResourceHandle, #hipErrorInvalidValue, + * #hipErrorNotSupported + * @warning : This API is marked as beta, meaning, while this is feature complete, + * it is still open to changes and may have outstanding issues. + */ +hipError_t hipLaunchHostFunc(hipStream_t stream, hipHostFn_t fn, void* userData); + +/** + * @brief Swaps the stream capture mode of a thread. + * + * @param [in] mode - Pointer to mode value to swap with the current mode + * @returns #hipSuccess, #hipErrorInvalidValue + * + * @warning : This API is marked as beta, meaning, while this is feature complete, + * it is still open to changes and may have outstanding issues. + * + */ +hipError_t hipThreadExchangeStreamCaptureMode(hipStreamCaptureMode* mode); + /** * @brief Creates a graph * diff --git a/samples/0_Intro/module_api/launchKernelHcc.cpp b/samples/0_Intro/module_api/launchKernelHcc.cpp index 39d3bd31d2..464f6d8851 100644 --- a/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/samples/0_Intro/module_api/launchKernelHcc.cpp @@ -69,9 +69,6 @@ int main() { HIP_CHECK(hipModuleLoad(&Module, fileName)); HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name)); - uint32_t len = LEN; - uint32_t one = 1; - struct { void* _Ad; void* _Bd; diff --git a/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp b/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp index 581f742834..e094f70d07 100644 --- a/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp +++ b/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp @@ -401,7 +401,6 @@ void ResultDatabase::DumpCsv(string fileName) { // **************************************************************************** bool ResultDatabase::IsFileEmpty(string fileName) { - bool fileEmpty; ifstream file(fileName.c_str()); @@ -409,6 +408,7 @@ bool ResultDatabase::IsFileEmpty(string fileName) { if (!file.good()) { return true; } else { + bool fileEmpty; fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof()); file.close(); diff --git a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 8032bd0a20..79d80ce317 100644 --- a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -252,9 +252,9 @@ void RunBenchmark_H2D(ResultDatabase& resultDB) { case MallocUnpinned: if (p_alignedhost) { - delete[] hostMem; - } else { free(hostMem); + } else { + delete[] hostMem; } break; diff --git a/samples/1_Utils/hipCommander/ResultDatabase.cpp b/samples/1_Utils/hipCommander/ResultDatabase.cpp index 1b1ee3a70d..51ced81fae 100644 --- a/samples/1_Utils/hipCommander/ResultDatabase.cpp +++ b/samples/1_Utils/hipCommander/ResultDatabase.cpp @@ -393,7 +393,6 @@ void ResultDatabase::DumpCsv(string fileName) { // **************************************************************************** bool ResultDatabase::IsFileEmpty(string fileName) { - bool fileEmpty; ifstream file(fileName.c_str()); @@ -401,6 +400,7 @@ bool ResultDatabase::IsFileEmpty(string fileName) { if (!file.good()) { return true; } else { + bool fileEmpty; fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof()); file.close(); diff --git a/samples/1_Utils/hipCommander/hipCommander.cpp b/samples/1_Utils/hipCommander/hipCommander.cpp index 37eb0845b1..e95ecd82a1 100644 --- a/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/samples/1_Utils/hipCommander/hipCommander.cpp @@ -282,7 +282,7 @@ class Command { // HCC optimizes away fully NULL kernel calls, so run one that is nearly null: class ModuleKernelCommand : public Command { public: - ModuleKernelCommand(CommandStream* cmdStream, const std::vector args) + ModuleKernelCommand(CommandStream* cmdStream, const std::vector& args) : Command(cmdStream, args), _stream(cmdStream->currentStream()) { hipModule_t module; HIPCHECK(hipModuleLoad(&module, FILENAME)); @@ -316,7 +316,7 @@ class ModuleKernelCommand : public Command { class KernelCommand : public Command { public: enum Type { Null, VectorAdd }; - KernelCommand(CommandStream* cmdStream, const std::vector args, Type kind) + KernelCommand(CommandStream* cmdStream, const std::vector& args, Type kind) : Command(cmdStream, args), _kind(kind), _stream(cmdStream->currentStream()){}; ~KernelCommand(){}; @@ -390,7 +390,7 @@ class CopyCommand : public Command { }; - void dealloc(void* p, MemType memType) { + static void dealloc(void* p, MemType memType) { if (memType == Device) { HIPCHECK(hipFree(p)); } else if (memType == PinnedHost) { @@ -433,7 +433,7 @@ class StreamSyncCommand : public Command { StreamSyncCommand(CommandStream* cmdStream, const std::vector& args) : Command(cmdStream, args), _stream(cmdStream->currentStream()){}; - const char* help() { return "synchronizes the current stream"; }; + static const char* help() { return "synchronizes the current stream"; }; void run() override { HIPCHECK(hipStreamSynchronize(_stream)); }; @@ -537,8 +537,8 @@ CopyCommand::CopyCommand(CommandStream* cmdStream, const std::vectorcurrentStream()) { + _stream(cmdStream->currentStream()), + _kind(kind) { switch (kind) { case hipMemcpyDeviceToHost: _srcType = Device; diff --git a/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp b/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp index ca312c0146..b4c8487b67 100644 --- a/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp +++ b/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp @@ -48,17 +48,14 @@ __global__ void incrementKernel(int32_t* in, int32_t* out, int32_t value, size_t int main() { int32_t incrementValue = 10; - // Host pointers - int32_t* hInput = nullptr; - int32_t* hOutput = nullptr; // Device pointers int32_t* dInput = nullptr; int32_t* dOutput = nullptr; size_t NBytes = SIZE * sizeof(int32_t); - - hInput = static_cast(malloc(NBytes)); - hOutput = static_cast(malloc(NBytes)); + // Host pointers + int32_t* hInput = static_cast(malloc(NBytes)); + int32_t* hOutput = static_cast(malloc(NBytes)); HIP_STATUS_CHECK(hipMalloc(&dInput, NBytes)); HIP_STATUS_CHECK(hipMalloc(&dOutput, NBytes)); @@ -95,4 +92,4 @@ int main() { std::cout << "success\n"; } return 0; -} \ No newline at end of file +} diff --git a/samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp b/samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp index a93eb141ac..a3c3f8f164 100644 --- a/samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp +++ b/samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp @@ -58,8 +58,8 @@ void run_test2() { HIP_ASSERT(hipFree(A_d)); HIP_ASSERT(hipFree(B_d)); - free(A_h); - free(B_h); + delete [] A_h; + delete [] B_h; std::cout << "Test Passed!\n"; } diff --git a/tests/catch/stress/memory/hipMallocManagedStress.cc b/tests/catch/stress/memory/hipMallocManagedStress.cc index 0e7d80edb6..7e992cd152 100644 --- a/tests/catch/stress/memory/hipMallocManagedStress.cc +++ b/tests/catch/stress/memory/hipMallocManagedStress.cc @@ -1,5 +1,5 @@ /* - Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. + 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 @@ -19,6 +19,15 @@ // The following test case allocation, host access, device access of HMM // memory from size 1 to 10KB +/* Test Case Description: + 1) Testing allocation, host access, device access of HMM + memory from size 1 to 10KB + 2) The following test case tests the behavior of kernel with a HMM memory + and hipMalloc memory + 3) The following test case tests when the same Hmm memory is used for + launching multiple different kernels will results in any issue + 4) Testing the allocation of/scenarios around max possible memory + */ #include #include @@ -35,6 +44,86 @@ __global__ void KrnlWth2MemTypesC(unsigned char *Hmm, unsigned char *Dptr, } static bool IfTestPassed = true; +// Kernel functions +__global__ void KrnlWth2MemTypes(int *Hmm, int *Dptr, size_t n) { + size_t index = blockIdx.x * blockDim.x + threadIdx.x; + for (size_t i = index; i < n; i++) { + Hmm[i] = Dptr[i] + 10; + } +} + +__global__ void KernelMulAdd_MngdMem(int *Hmm, size_t n) { + size_t index = blockIdx.x * blockDim.x + threadIdx.x; + size_t stride = blockDim.x * gridDim.x; + for (size_t i = index; i < n; i += stride) { + Hmm[i] = Hmm[i] * 2 + 10; + } +} + +__global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) { + size_t index = blockIdx.x * blockDim.x + threadIdx.x; + size_t stride = blockDim.x * gridDim.x; + for (size_t i = index; i < n; i += stride) { + Hmm[i] = Dptr[i] * 10; + } +} +static bool IfTestPassed = true; + +static void LaunchKrnl4(size_t NumElms, int InitVal) { + int *Hmm = NULL, *Dptr = NULL, blockSize = 64, DataMismatch = 0; + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, (sizeof(int) * NumElms))); + HIP_CHECK(hipMalloc(&Dptr, (sizeof(int) * NumElms))); + int *Hstptr = reinterpret_cast(new int[NumElms]); + for (size_t i = 0; i < NumElms; ++i) { + Hstptr[i] = InitVal; + } + HIP_CHECK(hipMemcpy(Dptr, Hstptr, (NumElms * sizeof(int)), + hipMemcpyHostToDevice)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + KrnlWth2MemTypes<<>>(Hmm, Dptr, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + for (size_t i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + INFO("Data Mismatch observed after the Kernel: KrnlWth2MemTypes!!\n"); + REQUIRE(false); + } + DataMismatch = 0; + KernelMul_MngdMem<<>>(Hmm, Dptr, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // Verifying the result + for (size_t i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal * 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n"); + REQUIRE(false); + } + DataMismatch = 0; + KernelMulAdd_MngdMem<<>>(Hmm, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // Verifying the result + + for (size_t i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal * 10 * 2 + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n"); + REQUIRE(false); + } + delete[] Hstptr; +} + static int HmmAttrPrint() { int managed = 0; INFO("The following are the attribute values related to HMM for" @@ -104,3 +193,139 @@ TEST_CASE("Unit_hipMallocManaged_MultiSize") { } } +// The following test case tests the behavior of kernel with a HMM memory and +// hipMalloc memory + +TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") { + IfTestPassed = true; + int *Hmm = NULL, *Dptr = NULL, InitVal = 123; + size_t NumElms = (1024 * 1024); + int *Hptr = new int[NumElms], blockSize = 64, DataMismatch = 0; + int managed = HmmAttrPrint(); + if (managed == 1) { + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, sizeof(int) * NumElms)); + HIP_CHECK(hipMalloc(&Dptr, sizeof(int) * NumElms)); + for (size_t i = 0; i < NumElms; ++i) { + Hmm[i] = 0; + Hptr[i] = InitVal; + } + HIP_CHECK(hipMemcpy(Dptr, Hptr, sizeof(int) * NumElms, + hipMemcpyHostToDevice)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + KrnlWth2MemTypes<<>>(Hmm, Dptr, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // Verifying the results + for (size_t k = 0; k < NumElms; ++k) { + if (Hmm[k] != (InitVal + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("DataMismatch observed!\n"); + IfTestPassed = false; + } + + HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipFree(Dptr)); + delete[] Hptr; + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +// The following test case tests when the same Hmm memory is used for +// launching multiple different kernels will results in any issue +TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") { + int managed = HmmAttrPrint(); + if (managed) { + int InitVal = 123, NumElms = (1024 * 1024); + LaunchKrnl4(NumElms, InitVal); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +// Testing the allocation of/scenarios around max possible memory +TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") { + int managed = HmmAttrPrint(); + if (managed == 1) { + bool IfTestPassed = true; + hipError_t err; + void *Hmm = NULL; + size_t totalDevMem = 0, freeDevMem = 0; + int NumDevs = 0; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + // Testing allocation of extreme and unusual mem values + for (int i = 0; i < NumDevs; i++) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMemGetInfo(&freeDevMem, &totalDevMem)); + err = hipMallocManaged(&Hmm, 1, hipMemAttachGlobal); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating memory on GPU: " << i); + WARN(" size 1 with"); + WARN(" hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachGlobal); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max free memory on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachGlobal); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max (free - 1) memory on "); + WARN("GPU: " << i); + WARN(" using hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, 1, hipMemAttachHost); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating memory size 1 on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachHost); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max free memory on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachHost); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max (freeDevMem - 1) memory" + " on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + } + REQUIRE(IfTestPassed); + } else { + SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result"); + } +} diff --git a/tests/catch/unit/deviceLib/CMakeLists.txt b/tests/catch/unit/deviceLib/CMakeLists.txt index 347e6d73c9..2f6819f43f 100644 --- a/tests/catch/unit/deviceLib/CMakeLists.txt +++ b/tests/catch/unit/deviceLib/CMakeLists.txt @@ -10,6 +10,7 @@ set(TEST_SRC popc.cc ldg.cc threadfence_system.cc + hipTestDeviceSymbol.cc ) # skipped for windows compiler issue - Illegal instruction detected diff --git a/tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc b/tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc new file mode 100644 index 0000000000..501a0805cb --- /dev/null +++ b/tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc @@ -0,0 +1,141 @@ +/* +Copyright (c) 2021 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 Case Description: Calling hipMemcpyTo/FromSymbolAsync() using user + declared stream obj and hipStreamPerThread*/ + +#include +#define NUM 1024 +#define SIZE 1024 * 4 + +__device__ int globalIn[NUM]; +__device__ int globalOut[NUM]; + +__global__ void Assign(int* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + Out[tid] = globalIn[tid]; + globalOut[tid] = globalIn[tid]; +} + +__device__ __constant__ int globalConst[NUM]; + +__global__ void checkAddress(int* addr, bool* out) { + *out = (globalConst == addr); +} + +TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") { + int *A, *Am, *B, *Ad, *C, *Cm; + A = new int[NUM]; + B = new int[NUM]; + C = new int[NUM]; + for (int i = 0; i < NUM; i++) { + A[i] = -1 * i; + B[i] = 0; + C[i] = 0; + } + + HIP_CHECK(hipMalloc((void**)&Ad, SIZE)); + HIP_CHECK(hipHostMalloc((void**)&Am, SIZE)); + HIP_CHECK(hipHostMalloc((void**)&Cm, SIZE)); + for (int i = 0; i < NUM; i++) { + Am[i] = -1 * i; + Cm[i] = 0; + } + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (int i = 0; i < NUM; i++) { + assert(Am[i] == B[i]); + assert(Am[i] == Cm[i]); + } + + for (int i = 0; i < NUM; i++) { + A[i] = -2 * i; + B[i] = 0; + } + + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, + hipMemcpyHostToDevice)); + hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, + hipMemcpyDeviceToHost)); + for (int i = 0; i < NUM; i++) { + assert(A[i] == B[i]); + assert(A[i] == C[i]); + } + + for (int i = 0; i < NUM; i++) { + A[i] = -3 * i; + B[i] = 0; + } + SECTION("Calling hipMemcpyTo/FromSymbol using user declared stream obj") { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, + hipMemcpyHostToDevice, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, + hipMemcpyDeviceToHost, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } + for (int i = 0; i < NUM; i++) { + assert(A[i] == B[i]); + assert(A[i] == C[i]); + } + + bool *checkOkD; + bool checkOk = false; + size_t symbolSize = 0; + int *symbolAddress; + HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst))); + HIP_CHECK(hipGetSymbolAddress((void**) &symbolAddress, HIP_SYMBOL(globalConst))); + HIP_CHECK(hipMalloc((void**)&checkOkD, sizeof(bool))); + hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD); + HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(checkOkD)); + HIP_ASSERT(checkOk); + HIP_ASSERT((symbolSize == SIZE)); + + HIP_CHECK(hipHostFree(Am)); + HIP_CHECK(hipHostFree(Cm)); + HIP_CHECK(hipFree(Ad)); + delete[] A; + delete[] B; + delete[] C; +} diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 557ddeca06..afcfb2028a 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -1,4 +1,4 @@ -# Copyright (c) 2021 Advanced Micro Devices, Inc. All Rights Reserved. +# 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 @@ -25,6 +25,19 @@ set(TEST_SRC hipGraph.cc hipSimpleGraphWithKernel.cc hipGraphAddMemcpyNode.cc + hipGraphClone.cc + hipGraphInstantiateWithFlags.cc + hipGraphAddHostNode.cc + hipGraphAddMemcpyNodeFromSymbol.cc + hipGraphChildGraphNodeGetGraph.cc + hipGraphNodeFindInClone.cc + hipGraphExecHostNodeSetParams.cc + hipGraphAddMemcpyNodeToSymbol.cc + hipGraphExecMemsetNodeSetParams.cc + hipGraphMemcpyNodeSetParamsToSymbol.cc + hipGraphDestroyNode.cc + hipGraphGetNodes.cc + hipGraphGetRootNodes.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphAddHostNode.cc b/tests/catch/unit/graph/hipGraphAddHostNode.cc new file mode 100644 index 0000000000..596a57d755 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphAddHostNode.cc @@ -0,0 +1,312 @@ +/* +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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/** +Testcase Scenarios of hipGraphAddHostNode API: + +Functional: +1. Creates graph, Adds HostNode which updates the variable and validates the result +2. Create graph, Add Graphnodes and clones the graph. Add Hostnode to the cloned graph + and validate the result +3. Creates graph which performs the square of number in the kernel function and the result + is validated in the callback function of hipGraphAddHostNode API + +Negative: + +1) Pass pGraphNode as nullptr and verify api doesn’t crash, returns error code. +2) Pass graph as nullptr and verify api doesn’t crash, returns error code. +3) Pass pNodeParams as nullptr and verify api doesn’t crash, returns error code. +4) Pass hipHostNodeParams::hipHostFn_t as nullptr and verify api doesn't crash, returns error code. +*/ + +#include +#include + +#define SIZE 1024 + +static int *B_h; +static int *D_h; + +static void callbackfunc(void *A_h) { + int *A = reinterpret_cast(A_h); + for (int i = 0; i < SIZE; i++) { + A[i] = i; + } +} + +static void __global__ vector_square(int *B_d, int *D_d) { + for (int i = 0; i < SIZE; i++) { + D_d[i] = B_d[i] * B_d[i]; + } +} +static void vectorsquare_callback(void* ptr) { + // The callback func is not working with zero parameters + // Temporary fix for adding the below 2 lines and ticket + // has been raised for the same. + int *A = reinterpret_cast(ptr); + A++; + for (int i = 0; i < SIZE; i++) { + if (D_h[i] != B_h[i] * B_h[i]) { + INFO("Validation failed " << D_h[i] << B_h[i]); + REQUIRE(false); + } + } +} +/* +This testcase verifies the negative scenarios of +hipGraphAddHostNode API +*/ +TEST_CASE("Unit_hipGraphAddHostNode_Negative") { + constexpr size_t N = 1024; + hipGraph_t graph; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + + SECTION("Passing nullptr to graph node") { + REQUIRE(hipGraphAddHostNode(nullptr, graph, + nullptr, + 0, &hostParams) == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to graph") { + REQUIRE(hipGraphAddHostNode(&hostNode, nullptr, + nullptr, + 0, &hostParams) == hipErrorInvalidValue); + } + +#if HT_NVIDIA + SECTION("Passing nullptr to host params") { + REQUIRE(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, nullptr) == hipErrorInvalidValue); + } +#endif + + SECTION("Passing nullptr to host func") { + hostParams.fn = nullptr; + REQUIRE(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, &hostParams) == hipErrorInvalidValue); + } + + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This testcase verifies hipGraphAddHostNode API in cloned graph +Creates graph, Add graph nodes and clone the graph +Add HostNode to the cloned graph and validate the result +*/ +TEST_CASE("Unit_hipGraphAddHostNode_ClonedGraphwithHostNode") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, memcpyH2D_C, + memcpyD2H_AC; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, + 0, C_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr, + 0, A_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + hipGraph_t clonedgraph; + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + HIP_CHECK(hipGraphAddHostNode(&hostNode, clonedgraph, + nullptr, + 0, &hostParams)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C, + &memcpyD2H_AC, 1)); + + // Instantiate and launch the cloned graph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify execution result + for (size_t i = 0; i < N; i++) { + if (A_h[i] != static_cast(i)) { + INFO("Validation failed i " << i << "C_h[i] "<< C_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, C_d, A_h, nullptr, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(clonedgraph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/* +This testcase verifies the square of number by +creating graph, Add kernel node which does the square +of number and the result is validated byhipGrahAddHostNode API +*/ +TEST_CASE("Unit_hipGraphAddHostNode_VectorSquare") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *A_h{nullptr}, *B_d{nullptr}, *D_d{nullptr}; + int *param = reinterpret_cast(sizeof(int));; + HipTest::initArrays(&A_d, &B_d, &D_d, + &A_h, &B_h, &D_h, N, false); + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_B, memcpyH2D_D, memcpyD2H_D, kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = vectorsquare_callback; + hostParams.userData = param; + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, + 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_D, graph, nullptr, + 0, D_d, D_h, + Nbytes, hipMemcpyHostToDevice)); + + void* kernelArgs2[] = {&B_d, &D_d}; + kernelNodeParams.func = reinterpret_cast(vector_square); + kernelNodeParams.gridDim = dim3(1); + kernelNodeParams.blockDim = dim3(1); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_D, graph, nullptr, + 0, D_h, D_d, + Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, &hostParams)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, + 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_D, &kernel_vecAdd, + 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, + &memcpyD2H_D, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_D, + &hostNode, 1)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + HipTest::freeArrays(A_d, B_d, D_d, A_h, B_h, D_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} +/* +This testcase verifies the following scenario +Create graph, calls the host function and updates +the parameters in the callback function and +validates it. +*/ +TEST_CASE("Unit_hipGraphAddHostNode_BasicFunc") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, memcpyD2H_AC, memcpyH2D_C; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, + 0, C_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr, + 0, A_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, &hostParams)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_AC, + &hostNode, 1)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify execution result + for (size_t i = 0; i < N; i++) { + if (A_h[i] != static_cast(i)) { + INFO("Validation failed i " << i << "A_h[i] "<< A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, C_d, A_h, nullptr, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} diff --git a/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc b/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc index 5fcf1ccc40..6fcb3fb90e 100644 --- a/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc +++ b/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc @@ -19,18 +19,22 @@ THE SOFTWARE. /** Testcase Scenarios : - 1) Add multiple Memcpy nodes to graph and verify node execution is - working as expected. +1) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D). +Memcpy nodes are added and assigned to default device. +2) Perform memcpy operation for 1D, 2D and 3D arrays on default device and verify the results. +3) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D). +Memcpy nodes are added and assigned to Peer device. +4) Perform memcpy operation for 1D, 2D and 3D arrays on Peer device and verify the results. */ #include #include -/** - * Functional Test adds memcpy nodes of types H2D, D2D and D2H to graph - * and verifies execution sequence by launching graph. - */ -TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { +#define ZSIZE 32 +#define YSIZE 32 +#define XSIZE 32 + +void validateMemcpyNode3DArray(bool peerAccess = false) { constexpr int width{10}, height{10}, depth{10}; hipArray *devArray1, *devArray2; hipChannelFormatKind formatKind = hipChannelFormatKindSigned; @@ -42,6 +46,7 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { hipStream_t streamForGraph; hipGraphExec_t graphExec; + HIP_CHECK(hipSetDevice(0)); int *hData = reinterpret_cast(malloc(size)); int *hOutputData = reinterpret_cast(malloc(size)); @@ -69,6 +74,12 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { make_hipExtent(width, height, depth), hipArrayDefault)); HIP_CHECK(hipGraphCreate(&graph, 0)); + // For peer access test, Memory is allocated on device(0) + // while memcpy nodes are allocated and assigned to peer device(1) + if (peerAccess) { + HIP_CHECK(hipSetDevice(1)); + } + // Host to Device memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); myparams.srcPos = make_hipPos(0, 0, 0); @@ -79,7 +90,6 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { myparams.dstArray = devArray1; myparams.kind = hipMemcpyHostToDevice; - HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams)); dependencies.push_back(memcpyNode); @@ -126,3 +136,247 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { free(hData); free(hOutputData); } + +void validateMemcpyNode2DArray(bool peerAccess = false) { + int harray2D[YSIZE][XSIZE]{}; + int harray2Dres[YSIZE][XSIZE]{}; + constexpr int width{XSIZE}, height{YSIZE}; + hipArray *devArray1, *devArray2; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparams; + hipGraph_t graph; + hipGraphNode_t memcpyNode; + std::vector dependencies; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Initialize 2D object + for (int i = 0; i < YSIZE; i++) { + for (int j = 0; j < XSIZE; j++) { + harray2D[i][j] = i + j + 1; + } + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + // Allocate 2D device array by passing depth(0) + HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc, + make_hipExtent(width, height, 0), hipArrayDefault)); + HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc, + make_hipExtent(width, height, 0), hipArrayDefault)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // For peer access test, Memory is allocated on device(0) + // while memcpy nodes are allocated and assigned to peer device(1) + if (peerAccess) { + HIP_CHECK(hipSetDevice(1)); + } + + // Host to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, height, 1); + myparams.srcPtr = make_hipPitchedPtr(harray2D, width * sizeof(int), + width, height); + myparams.dstArray = devArray1; + myparams.kind = hipMemcpyHostToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams)); + dependencies.push_back(memcpyNode); + + // Device to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.srcArray = devArray1; + myparams.dstArray = devArray2; + myparams.extent = make_hipExtent(width, height, 1); + myparams.kind = hipMemcpyDeviceToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + dependencies.clear(); + dependencies.push_back(memcpyNode); + + // Device to host + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, height, 1); + myparams.dstPtr = make_hipPitchedPtr(harray2Dres, width * sizeof(int), + width, height); + myparams.srcArray = devArray2; + myparams.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Validate result + for (int i = 0; i < YSIZE; i++) { + for (int j = 0; j < XSIZE; j++) { + if (harray2D[i][j] != harray2Dres[i][j]) { + INFO("harray2D: " << harray2D[i][j] << "harray2Dres: " + << harray2Dres[i][j] << " mismatch at (i,j) : " << i << j); + REQUIRE(false); + } + } + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + hipFreeArray(devArray1); + hipFreeArray(devArray2); +} + +void validateMemcpyNode1DArray(bool peerAccess = false) { + int harray1D[XSIZE]{}; + int harray1Dres[XSIZE]{}; + constexpr int width{XSIZE}; + hipArray *devArray1, *devArray2; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparams; + hipGraph_t graph; + hipGraphNode_t memcpyNode; + std::vector dependencies; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Initialize 1D object + for (int i = 0; i < XSIZE; i++) { + harray1D[i] = i + 1; + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + // Allocate 1D device array by passing depth(0), height(0) + HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc, + make_hipExtent(width, 0, 0), hipArrayDefault)); + HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc, + make_hipExtent(width, 0, 0), hipArrayDefault)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // For peer access test, Memory is allocated on device(0) + // while memcpy nodes are allocated and assigned to peer device(1) + if (peerAccess) { + HIP_CHECK(hipSetDevice(1)); + } + + // Host to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, 1, 1); + myparams.srcPtr = make_hipPitchedPtr(harray1D, width * sizeof(int), + width, 1); + myparams.dstArray = devArray1; + myparams.kind = hipMemcpyHostToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams)); + dependencies.push_back(memcpyNode); + + // Device to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.srcArray = devArray1; + myparams.dstArray = devArray2; + myparams.extent = make_hipExtent(width, 1, 1); + myparams.kind = hipMemcpyDeviceToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + dependencies.clear(); + dependencies.push_back(memcpyNode); + + // Device to host + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, 1, 1); + myparams.dstPtr = make_hipPitchedPtr(harray1Dres, width * sizeof(int), + width, 1); + myparams.srcArray = devArray2; + myparams.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Validate result + for (int i = 0; i < XSIZE; i++) { + if (harray1D[i] != harray1Dres[i]) { + INFO("harray1D: " << harray1D[i] << " harray1Dres: " << harray1Dres[i] + << " mismatch at : " << i); + REQUIRE(false); + } + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + hipFreeArray(devArray1); + hipFreeArray(devArray2); +} + +/** + * Basic Functional Tests adds memcpy nodes of types H2D, D2D and D2H to graph + * and verifies execution sequence by launching graph on default device. + * Tests also verify memcpy node addition with 1D, 2D and 3D objects. + */ +TEST_CASE("Unit_hipGraphAddMemcpyNode_BasicFunctional") { + SECTION("Memcpy with 3D array on default device") { + validateMemcpyNode3DArray(); + } + + SECTION("Memcpy with 2D array on default device") { + validateMemcpyNode2DArray(); + } + + SECTION("Memcpy with 1D array on default device") { + validateMemcpyNode1DArray(); + } +} + +/** + * Peer access tests adds and assigns memcpy nodes of types H2D, D2D and D2H + * to peer device. Memory allocations happen on device(0) and memcpy operations + * are performed from device(1). + * Tests also verify memcpy node addition with 1D, 2D and 3D objects. + */ +TEST_CASE("Unit_hipGraphAddMemcpyNode_PeerAccessFunctional") { + int numDevices{}, peerAccess{}; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + } + + if (!peerAccess) { + WARN("Skipping test as peer device access is not found!"); + return; + } + + SECTION("Memcpy with 3D array on peer device") { + validateMemcpyNode3DArray(true); + } + + SECTION("Memcpy with 2D array on peer device") { + validateMemcpyNode2DArray(true); + } + + SECTION("Memcpy with 1D array on peer device") { + validateMemcpyNode1DArray(true); + } +} diff --git a/tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc b/tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc new file mode 100644 index 0000000000..7d9cbaffed --- /dev/null +++ b/tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc @@ -0,0 +1,435 @@ +/* +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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/** +Testcase Scenarios of hipGraphAddMemcpyNodeFromSymbol API: + +Functional : + +1. Allocate global symbol memory, add the MemcpyNodeFromSymbol + node to the graph and verify for different memory kinds +2. Allocate const memory add the MemcpyNodeFromSymbol node to + the graph and verify for different memory kinds +3. Allocate global symbol memory and device memory in GPU-0 + and perform MemcpyToSymbol from peer GPU by adding it to the graph node. +4. Allocate const symbol memory and device memory in GPU-0 + and perform MemcpyToSymbol from peer GPU by adding it to the graph node. +5. Allocate global memory, Add MemcpyFromSymbolNode,KernelNode and memcpynode and validating + the behaviour + +Negative : + +1) Pass nullptr to graph node +2) Pass nullptr to graph +3) Pass nullptr to dependencies +4) Pass invalid numDependencies +5) Pass nullptr to dst +6) Pass nullptr to symbol +7) Pass invalid count +8) Pass offset+count greater than allocated size +9) Pass unintialized graph +*/ + +#include +#include +#include +#define SIZE 256 + +__device__ int globalIn[SIZE]; +__device__ int globalOut[SIZE]; +__device__ __constant__ int globalConst[SIZE]; + +__global__ void MemcpyFromSymbolKernel(int* B_d) { + for (int i = 0 ; i < SIZE; i++) { + globalIn[i] = B_d[i]; + } +} + +/* This testcase verifies negative scenarios of + hipGraphAddMemcpyNodeFromSymbol API */ +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_Negative") { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}, *B_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphNode_t memcpyToSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + // Adding MemcpyNodeToSymbol + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + +#if HT_NVIDIA + hipGraphNode_t memcpyFromSymbolNode; + SECTION("Passing nullptr to graph") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, nullptr, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to graph node") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(nullptr, graph, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing size > 1 and dependencies as nullptr") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + nullptr, + 1, + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing invalid dependencies size") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + 10, + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to dst") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + nullptr, + HIP_SYMBOL(globalIn), Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to source") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_d, + nullptr, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidSymbol); + } + + SECTION("Passing offset+size > max size") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 10, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing Max count") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + std::numeric_limits::max(), 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Pass Unintialized graph") { + hipGraph_t unint_graph; + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, unint_graph, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } +#endif + + HipTest::freeArrays(A_d, B_d, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This function is used to verify the following scenarios +1. Create global variable, allocate Memory in GPU-0 and create dependency graph of + hipGraphAddMemcpyNodeFromSymbol API in GPU-1 and validate the result +2. Allocate global memory, Create dependency graph and validate the result on GPU-0 +3. Allocate global const memory, Create dependency graph and validate the result on GPU-0 +4. Create global const variable, allocate Memory in GPU-0 and create dependency graph of + hipGraphAddMemcpyNodeFromSymbol API in GPU-1 and validate the result +*/ + +void hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(bool device_ctxchg = false, + bool const_device_var = + false) { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + if (device_ctxchg) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipDeviceEnablePeerAccess(0, 0)); + } + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + // Adding MemcpyNodeToSymbol + if (const_device_var) { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalConst), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + + } else { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + } + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + + // Adding MemcpyNodeFromSymbol + if (const_device_var) { + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_h, + HIP_SYMBOL(globalConst), + Nbytes, 0, + hipMemcpyDeviceToHost)); + } else { + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_h, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToHost)); + } + + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Validating the result + for (int i = 0; i < SIZE; i++) { + if (B_h[i] != A_h[i]) { + WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This testcase verifies allocating global symbol memory, +add the MemcpyNodeFromSymbol node to the graph and +erifying the result +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemory") { + hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(false, false); +} + +/* +This testcase verifies allocating global const symbol memory, +add the MemcpyNodeFromSymbol node to the graph and +verifying the result +*/ + +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalConstMemory") { + hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(false, true); +} + +/* +This testcase verifies allocating global symbol memory and device variables +in GPU-0 and add the MemcpyNodeFromSymbol node to the graph and +verifying the result in GPU-1 +*/ +#if HT_NVIDIA +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryPeerDevice") { + int numDevices = 0; + int canAccessPeer = 0; + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(true, false); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} + +/* +This testcase verifies allocating global const symbol memory and device variables +in GPU-0 and add the MemcpyNodeFromSymbol node to the graph and +verifying the result in GPU-1 +*/ + +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalConstMemoryPeerDevice") { + int numDevices = 0; + int canAccessPeer = 0; + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(true, true); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} +#endif +/* +This testcaser verifies allocating global memory, +Add MemcpyFromSymbolNode,KernelNode and memcpynode and validating +the behaviour +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryWithKernel") { + constexpr size_t Nbytes = SIZE * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, SIZE); + hipGraphNode_t memcpyfromsymbolkernel, memcpyD2H_B; + hipKernelNodeParams kernelNodeParams{}; + int *A_d{nullptr}, *B_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice)); + dependencies.clear(); + dependencies.push_back(memcpyFromSymbolNode); + + // Adding Kernel node + void* kernelArgs1[] = {&B_d}; + kernelNodeParams.func = + reinterpret_cast(MemcpyFromSymbolKernel); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&memcpyfromsymbolkernel, graph, + dependencies.data(), dependencies.size(), + &kernelNodeParams)); + dependencies.clear(); + dependencies.push_back(memcpyfromsymbolkernel); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_B, graph, dependencies.data(), + dependencies.size(), B_h, B_d, + Nbytes, hipMemcpyDeviceToHost)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Validating the result + for (int i = 0; i < SIZE; i++) { + if (B_h[i] != A_h[i]) { + WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, B_d, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc b/tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc new file mode 100644 index 0000000000..5a5b08f77e --- /dev/null +++ b/tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc @@ -0,0 +1,400 @@ +/* +Copyright (c) 2021 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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/** +Testcase Scenarios of hipGraphAddMemcpyNodeToSymbol API: + +Functional : + +1. Allocate global symbol memory, add the MemcpyNodeToSymbol + node to the graph and verify for different memory kinds +2. Allocate const memory add the MemcpyNodeToSymbol node to + the graph and verify for different memory kinds +3. Allocate global symbol memory and device memory in GPU-0 + and perform MemcpyToSymbol from peer GPU by adding it to the graph node. +4. Allocate const symbol memory and device memory in GPU-0 + and perform MemcpyToSymbol from peer GPU by adding it to the graph node. +5. Allocate global memory, Add MemcpyToSymbolNode,KernelNode and memcpynode and validating + the behaviour + +Negative : + +1) Pass nullptr to graph node +2) Pass nullptr to graph +3) Pass nullptr to dependencies +4) Pass invalid numDependencies +5) Pass nullptr to dst +6) Pass nullptr to symbol +7) Pass invalid count +8) Pass offset+count greater than allocated size +9) Pass unintialized graph +*/ + +#include +#include +#include +#define SIZE 256 + +__device__ int globalIn[SIZE]; +__device__ __constant__ int globalConst[SIZE]; + +__global__ void MemcpyToSymbolKernel(int* B_d) { + for (int i = 0 ; i < SIZE; i++) { + B_d[i] = globalIn[i]; + } +} + +/* This testcase verifies negative scenarios of + hipGraphAddMemcpyNodeToSymbol API */ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_Negative") { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphNode_t memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); +#if HT_NVIDIA + hipGraphNode_t memcpyToSymbolNode; + SECTION("Passing nullptr to graph") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, nullptr, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_h, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to graph node") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(nullptr, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing size > 1 and dependencies as nullptr") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + nullptr, + 1, + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing invalid dependencies size") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + 10, + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to dst") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + nullptr, + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidSymbol); + } + + SECTION("Passing nullptr to source") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + nullptr, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing offset+size > max size") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 10, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing Max count") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, + std::numeric_limits::max(), 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Pass Unintialized graph") { + hipGraph_t unint_graph; + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, unint_graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } +#endif + + HipTest::freeArrays(A_d, nullptr, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This function is used to verify the following scenarios +1. Create global variable, allocate Memory in GPU-0 and create dependency graph of + hipGraphAddMemcpyNodeToSymbol API in GPU-1 and validate the result +2. Allocate global memory, Create dependency graph and validate the result on GPU-0 +3. Allocate global const memory, Create dependency graph and validate the result on GPU-0 +4. Create global const variable, allocate Memory in GPU-0 and create dependency graph of + hipGraphAddMemcpyNodeToSymbol API in GPU-1 and validate the result +*/ +void hipGraphAddMemcpyNodeToSymbol_GlobalMemory(bool device_ctxchg = false, + bool const_device_var = false) { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + if (device_ctxchg) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipDeviceEnablePeerAccess(0, 0)); + } + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + // Adding MemcpyNodeToSymbol + + if (const_device_var) { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalConst), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + } else { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + } + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + // Adding MemcpyNodeFromSymbol + if (const_device_var) { + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_h, + HIP_SYMBOL(globalConst), + Nbytes, 0, hipMemcpyDeviceToHost)); + } else { + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_h, + HIP_SYMBOL(globalIn), + Nbytes, 0, hipMemcpyDeviceToHost)); + } + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Validating the result + for (int i = 0; i < SIZE; i++) { + if (B_h[i] != A_h[i]) { + WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This testcase verifies allocating global symbol memory, +add the MemcpyNodeToSymbol node to the graph and +erifying the result +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalMemory") { + hipGraphAddMemcpyNodeToSymbol_GlobalMemory(false, false); +} + +/* +This testcase verifies allocating global const symbol memory, +add the MemcpyNodeToSymbol node to the graph and +verifying the result +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalConstMemory") { + hipGraphAddMemcpyNodeToSymbol_GlobalMemory(false, true); +} + +#if HT_NVIDIA +/* +This testcase verifies allocating global symbol memory and device variables +in GPU-0 and add the MemcpyNodeToSymbol node to the graph and +verifying the result in GPU-1 +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalMemoryPeerDevice") { + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + if (canAccessPeer) { + hipGraphAddMemcpyNodeToSymbol_GlobalMemory(true, false); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} +/* +This testcase verifies allocating global const symbol memory and device variables +in GPU-0 and add the MemcpyNodeToSymbol node to the graph and +verifying the result in GPU-1 +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalConstMemoryPeerDevice") { + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + if (canAccessPeer) { + hipGraphAddMemcpyNodeToSymbol_GlobalMemory(true, true); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} +#endif +/* +This testcaser verifies allocating global memory, +Add MemcpyToSymbolNode,KernelNode and memcpynode and validating +the behaviour +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_MemcpyToSymbolNodeWithKernel") { + constexpr size_t Nbytes = SIZE * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, SIZE); + hipGraphNode_t memcpytosymbolkernel, memcpyD2H_B; + hipKernelNodeParams kernelNodeParams{}; + int *A_d{nullptr}, *B_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyToSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + // Adding Kernel node + void* kernelArgs1[] = {&B_d}; + kernelNodeParams.func = + reinterpret_cast(MemcpyToSymbolKernel); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&memcpytosymbolkernel, graph, + dependencies.data(), dependencies.size(), + &kernelNodeParams)); + dependencies.clear(); + dependencies.push_back(memcpytosymbolkernel); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_B, graph, dependencies.data(), + dependencies.size(), B_h, B_d, + Nbytes, hipMemcpyDeviceToHost)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Validating the result + for (int i = 0; i < SIZE; i++) { + if (B_h[i] != A_h[i]) { + WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, B_d, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc b/tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc new file mode 100644 index 0000000000..4a975dd3c0 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc @@ -0,0 +1,163 @@ +/*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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/** +Testcase Scenarios of hipGraphChildGraphNodeGetGraph API: + +Functional Scenarios: +1. Get the child graph node from the original graph and execute it + +Negative Scenarios: +1. Pass nullptr to graph +2. Pass nullptr to graphnode +3. Pass uninitialized graph node +4. Pass orginial graph node instead of child graph node +**/ + +#include +#include +#include + +/* +This testcase verifies the following scenario +Create graph, add multiple child nodes and gets the +graph of one of the child nodes using hipGraphChildGraphNodeGetGraph API +executes it and validates the results +*/ +TEST_CASE("Unit_hipGraphChildGraphNodeGetGraph_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + size_t NElem{N}; + constexpr auto threadsPerBlock = 256; + hipGraph_t graph, childgraph1, childgraph2; + hipGraphExec_t graphExec; + hipKernelNodeParams kernelNodeParams{}; + hipGraphNode_t kernel_vecAdd; + int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + 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)); + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, childGraphNode1, + childGraphNode2, memcpyD2H_C; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&childgraph1, 0)); + HIP_CHECK(hipGraphCreate(&childgraph2, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, childgraph1, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, childgraph2, nullptr, + 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph, + nullptr, 0, childgraph1)); + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode2, graph, + nullptr, 0, childgraph2)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, + 0, C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + HIP_CHECK(hipGraphAddDependencies(graph, &childGraphNode1, + &childGraphNode2, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &childGraphNode2, + &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + hipGraph_t Getgraph; + HIP_CHECK(hipGraphChildGraphNodeGetGraph(childGraphNode1, &Getgraph)); + // Instantiate and launch the child graph + HIP_CHECK(hipGraphInstantiate(&graphExec, Getgraph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + + // Verify child graph execution result + HIP_CHECK(hipMemcpy(C_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < N; i++) { + if (A_h[i] != C_h[i]) { + INFO("Validation failed " << A_h[i] << C_h[i]); + REQUIRE(false); + } + } + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(childgraph2)); + HIP_CHECK(hipGraphDestroy(childgraph1)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/* +This testcase verifies the negative scenarios +of hipGraphChildGraphNodeGetGraph API +*/ +TEST_CASE("Unit_hipGraphChildGraphNodeGetGraph_Negative") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph, childgraph1; + int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, childGraphNode1; + HIP_CHECK(hipGraphCreate(&childgraph1, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, childgraph1, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph, + nullptr, 0, childgraph1)); + + hipGraph_t Getgraph; + SECTION("nullptr to child node") { + REQUIRE((hipGraphChildGraphNodeGetGraph(nullptr, &Getgraph)) + == hipErrorInvalidValue); + } +#if HT_NVIDIA + SECTION("nullptr to graph") { + REQUIRE((hipGraphChildGraphNodeGetGraph(childGraphNode1, nullptr)) + == hipErrorInvalidValue); + } + + SECTION("Passing parent instead of child graph node") { + REQUIRE((hipGraphChildGraphNodeGetGraph(memcpyH2D_A, &Getgraph)) + == hipErrorInvalidValue); + } + + SECTION("Passing unintialized node") { + hipGraphNode_t unint_node; + REQUIRE((hipGraphChildGraphNodeGetGraph(unint_node, &Getgraph)) + == hipErrorInvalidValue); + } +#endif + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphDestroy(childgraph1)); +} diff --git a/tests/catch/unit/graph/hipGraphClone.cc b/tests/catch/unit/graph/hipGraphClone.cc new file mode 100644 index 0000000000..4a6cbc6ec9 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphClone.cc @@ -0,0 +1,319 @@ +/* +Copyright (c) 2021 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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/* +Testcase Scenarios of hipGraphClone API: + +Negative: + +1. Pass nullptr to cloned graph +2. pass nullptr to original graph + +Functional: + +1. Clone the graph,Instantiate and execute the cloned graph +2. Clone the graph and modify the original graph and ensure that the + cloned graph is not modified +3. Create graph on one GPU device and clone it from peer GPU device +4. Create graph in one thread and clone it from multiple threads. +*/ + +#include +#include +#include + +#define NUM_THREADS 10 + +/* This test covers the negative scenarios of + hipGraphClone API */ + +TEST_CASE("Unit_hipGraphClone_Negative") { + SECTION("Passing nullptr to Cloned graph") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + REQUIRE(hipGraphClone(nullptr, graph) == hipErrorInvalidValue); + HIP_CHECK(hipGraphDestroy(graph)); + } + + SECTION("Passing nullptr to original graph") { + hipGraph_t clonedGraph; + REQUIRE(hipGraphClone(&clonedGraph, nullptr) == hipErrorInvalidValue); + } +} +/* +This function creates the graph with dependencies +then performs device context change and clones the cloned graph +Executes the cloned graph and validates the result +*/ +void hipGraphClone_DeviceContextChange() { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph, clonedgraph; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + hipGraphNode_t memcpyH2D_A, memcpyD2H_A; + int *A_d{nullptr}, *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, N, false); + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_A, graph, nullptr, 0, B_h, A_d, + Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &memcpyD2H_A, 1)); + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + // Instantiate and launch the original graph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + for (size_t i = 0; i < N; i++) { + if (A_h[i] != B_h[i]) { + INFO("Validation failed A_h[i] " << A_h[i] << " B_h[i] " << B_h[i]); + REQUIRE(false); + } + } + HipTest::freeArrays(A_d, nullptr, nullptr, A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(clonedgraph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} +/* +This function does the following +1. Creates the graph with multiple dependencies + clones the graph and validates the result. +2. Creates the graph, clones the graph and modifies + the existing graph and execute the cloned graph + to ensure that cloned graph is not modified +*/ +void hipGraphClone_Func(bool ModifyOrigGraph = false) { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph, clonedgraph; + hipGraphNode_t memset_A, memset_B, memsetKer_C; + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C, memcpyD2D_C, + memcpyD2H_C_new; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + hipStream_t streamForGraph; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + hipGraphExec_t graphExec; + hipMemsetParams memsetParams{}; + int memsetVal{}; + size_t NElem{N}; + + 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)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, + &memsetParams)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(B_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0, + &memsetParams)); + + void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast(&NElem)}; + kernelNodeParams.func = + reinterpret_cast(HipTest::memsetReverse); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0, + &kernelNodeParams)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + + if (ModifyOrigGraph) { + // Modify Original graph by adding new dependency + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2D_C, graph, nullptr, 0, + C_d, B_d, + Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C_new, graph, nullptr, 0, + C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2D_C, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2D_C, + &memcpyD2H_C_new, 1)); + + // Instantiate and launch the original graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + for (size_t i= 0; i < NElem; i++) { + if (C_h[i] != B_h[i]) { + INFO("Validation failed C_h is " << C_h[i] << + "B_h is " << B_h[i]); + REQUIRE(false); + } + } + } + + // Instantiate and launch the cloned graph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, C_h, N); + 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(hipGraphDestroy(clonedgraph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/* +This testcase verifies following scenarios +1. Clones the graph and verify the result +2. Clones the graph, Modify the original graph and + validate the result of the cloned graph +3. Device context change for cloned graph +*/ +TEST_CASE("Unit_hipGraphClone_Functional") { + SECTION("hipGraphClone Basic Functionality") { + hipGraphClone_Func(); + } + SECTION("hipGraphClone Modify Original graph") { + hipGraphClone_Func(true); + } + + SECTION("hipGraphClone Device context change") { + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + hipGraphClone_DeviceContextChange(); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } + } +} + +/* +This testcase creates the graph with dependencies +then creates multiple threads and clones the graph +in each thread and executes the cloned graph +hipGraphClone is failing in CUDA in multi threaded +scenario so excluded for nvidia +*/ +#if HT_AMD +TEST_CASE("Unit_hipGraphClone_MultiThreaded") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphNode_t memcpyH2D_A, memcpyD2H_A; + int *A_d{nullptr}, *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, N, false); + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_A, graph, nullptr, 0, B_h, A_d, + Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &memcpyD2H_A, 1)); + std::vector threads; + auto lambdaFunc = [&](){ + hipGraph_t clonedgraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + // Instantiate and launch the cloned graph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, + nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + for (size_t i = 0; i < N; i++) { + if (A_h[i] != B_h[i]) { + INFO("Validation failed A_h[i] " << A_h[i] << " B_h[i] " << B_h[i]); + REQUIRE(false); + } + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(clonedgraph)); + }; + for (int i = 0; i < NUM_THREADS; i++) { + std::thread t(lambdaFunc); + threads.push_back(std::move(t)); + } + for (auto &t : threads) { + t.join(); + } + HipTest::freeArrays(A_d, nullptr, nullptr, A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); +} +#endif diff --git a/tests/catch/unit/graph/hipGraphDestroyNode.cc b/tests/catch/unit/graph/hipGraphDestroyNode.cc new file mode 100644 index 0000000000..4c544a3f68 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphDestroyNode.cc @@ -0,0 +1,139 @@ +/* +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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/* +Testcase Scenarios of hipGraphDestroyNode API: + +Negative :: + +1) Pass nullptr to graph node + +Functional :: + +1) Create Node and destroy the node +2) Create graph with dependencies and destroy one of the dependency node + before executing the graph. +*/ + +#include +#include +#include + + +/* This test covers the negative scenarios of + hipGraphDestroyNode API */ +TEST_CASE("Unit_hipGraphDestroyNode_Negative") { + SECTION("Passing nullptr to graph Node") { + REQUIRE(hipGraphDestroyNode(nullptr) == hipErrorInvalidValue); + } +} + +/* This test covers the basic functionality of + hipGraphDestroyNode API where we create and destroy + the node +*/ +TEST_CASE("Unit_hipGraphDestroyNode_BasicFunctionality") { + char *pOutBuff_d{}; + constexpr size_t size = 1024; + hipGraph_t graph{}; + hipGraphNode_t memsetNode{}; + + HIP_CHECK(hipMalloc(&pOutBuff_d, size)); + hipMemsetParams memsetParams{}; + memsetParams.dst = reinterpret_cast(pOutBuff_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = size * sizeof(char); + memsetParams.height = 1; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + REQUIRE(hipGraphDestroyNode(memsetNode) == hipSuccess); + HIP_CHECK(hipFree(pOutBuff_d)); +} + +/* +This testcase verifies the following scenario where +graph is created with dependencies and one of the dependency is +destroyed before execute the graph +*/ +TEST_CASE("Unit_hipGraphDestroyNode_DestroyDependencyNode") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph; + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyH2D_B2Copies, memcpyD2H_C; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + hipGraphExec_t graphExec; + size_t NElem{N}; + hipStream_t streamForGraph; + + 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(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B2Copies, graph, nullptr, + 0, B_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, B_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B2Copies, &kernel_vecAdd, + 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + + // Destroy one of the dependency node + HIP_CHECK(hipGraphDestroyNode(memcpyH2D_B)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, C_h, B_h, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/tests/catch/unit/graph/hipGraphExecHostNodeSetParams.cc b/tests/catch/unit/graph/hipGraphExecHostNodeSetParams.cc new file mode 100644 index 0000000000..6f07b23269 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphExecHostNodeSetParams.cc @@ -0,0 +1,276 @@ +/* +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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/** +Testcase Scenarios of hipGraphExecHostNodeSetParams API: + +Functional: +1. Creates graph, Adds HostNode, update hostNode params using hipGraphExecHostNodeSetParams API + and validates the result +2. Create graph, Add Graphnodes and clones the graph. Add Hostnode to the cloned graph, update + hostNode params using hipGraphExecHostNodeSetParams API and validate the result + +Negative: + +1) Pass hGraphExec as nullptr and verify api doen't crash, returns error code. +2) Pass node as nullptr and verify api doen't crash, returns error code. +3) Pass pNodeParams as nullptr and verify api doesn't crash, returns error code. +3) Pass hipHostNodeParams::hipHostFn_t as nullptr and verify api doesn't crash, returns error code. +4) Pass unintialized host params and verify api doesn't crash, returns error code. +5) Pass unintialized graph and verify api doesn't crash, returns error code. +6) Pass nullptr to hostfunc and verify api doesn't crash, returns error code. +*/ + +#include +#include + +#define SIZE 1024 + +void callbackfunc(void *A_h) { + int *A = reinterpret_cast(A_h); + for (int i = 0; i < SIZE; i++) { + A[i] = i; + } +} + +void callbackfunc_setparams(void *B_h) { + int *B = reinterpret_cast(B_h); + for (int i = 0; i < SIZE; i++) { + B[i] = i * i; + } +} + +/* +This testcase verifies the negative scenarios of +hipGraphExecHostNodeSetParams API +*/ +TEST_CASE("Unit_hipGraphExecHostNodeSetParams_Negative") { +constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, memcpyD2H_AC, memcpyH2D_C; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, + 0, C_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr, + 0, A_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, &hostParams)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_AC, + &hostNode, 1)); + + hipHostNodeParams sethostParams = {0, 0}; + sethostParams.fn = callbackfunc_setparams; + sethostParams.userData = C_h; + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); +#if HT_NVIDIA + SECTION("Passing nullptr to graphExec") { + REQUIRE(hipGraphExecHostNodeSetParams(nullptr, hostNode, &sethostParams) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to hostParams") { + REQUIRE(hipGraphExecHostNodeSetParams(graphExec, hostNode, nullptr) + == hipErrorInvalidValue); + } +#endif + SECTION("Passing nullptr to graph") { + REQUIRE(hipGraphExecHostNodeSetParams(graphExec, nullptr, &sethostParams) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to host func") { + sethostParams.fn = nullptr; + REQUIRE(hipGraphExecHostNodeSetParams(graphExec, hostNode, &sethostParams) + == hipErrorInvalidValue); + } + + + SECTION("Passing unintialized hostParams") { + hipHostNodeParams unintParams = {0, 0}; + REQUIRE(hipGraphExecHostNodeSetParams(graphExec, hostNode, &unintParams) + == hipErrorInvalidValue); + } + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This testcase verifies hipGraphExecHostNodeSetParams API in cloned graph +Creates graph, Add graph nodes and clone the graph +Add HostNode to the cloned graph,update the host params using +hipGraphExecHostNodeSetParams API and validates the result +*/ +TEST_CASE("Unit_hipGraphExecHostNodeSetParams_ClonedGraphwithHostNode") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, memcpyH2D_C, + memcpyD2H_AC; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, + 0, C_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr, + 0, A_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + hipGraph_t clonedgraph; + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + HIP_CHECK(hipGraphAddHostNode(&hostNode, clonedgraph, + nullptr, + 0, &hostParams)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C, + &memcpyD2H_AC, 1)); + + hipHostNodeParams sethostParams = {0, 0}; + sethostParams.fn = callbackfunc_setparams; + sethostParams.userData = C_h; + + + // Instantiate and launch the cloned graph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphExecHostNodeSetParams(graphExec, hostNode, &sethostParams)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify execution result + for (size_t i = 0; i < N; i++) { + if (C_h[i] != static_cast(i * i)) { + INFO("Validation failed i " << i << "C_h[i] "<< C_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, C_d, A_h, nullptr, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/* +This testcase verifies the following scenario +Create graph, Adds host node to the graph, +updates the host params using hipGraphExecHostNodeSetParams API +and validates the result +*/ +TEST_CASE("Unit_hipGraphExecHostNodeSetParams_BasicFunc") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, memcpyD2H_AC, memcpyH2D_C; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, + 0, C_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr, + 0, A_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, &hostParams)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_AC, + &hostNode, 1)); + + hipHostNodeParams sethostParams = {0, 0}; + sethostParams.fn = callbackfunc_setparams; + sethostParams.userData = C_h; + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphExecHostNodeSetParams(graphExec, hostNode, &sethostParams)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify execution result + for (size_t i = 0; i < N; i++) { + if (C_h[i] != static_cast(i * i)) { + INFO("Validation failed i " << i << "C_h[i] "<< C_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, C_d, A_h, nullptr, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} diff --git a/tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc b/tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc new file mode 100644 index 0000000000..7e51eb296b --- /dev/null +++ b/tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc @@ -0,0 +1,201 @@ +/* +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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : +Functional- +1) Instantiate a graph with memset node, obtain executable graph and update the + hipMemsetParams node params with set. Make sure they are taking effect. +Negative- +1) Pass hGraphExec as nullptr and verify api returns error code. +2) Pass graph node as nullptr and verify api returns error code. +3) Pass different hipGraphNode_t which was not used in graphExec and verify api returns error code. +4) Pass Pass different Graph which was not used in graphExec and verify api returns error code. +5) Pass pNodeParams as nullptr and verify api returns error code. +6) Pass pNodeParams as empty structure object and verify api returns error code. +7) Pass hipMemsetParams::dst as nullptr, api should return error code. +8) Pass hipMemsetParams::element size other than 1, 2, or 4 and check api should return error code. +9) Pass hipMemsetParams::height as zero and check api should return error code. +*/ + +#include +#include + +/* Test verifies hipGraphExecMemsetNodeSetParams API Negative scenarios. + */ +TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Negative") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(char); + constexpr size_t val = 0; + char *devData, *hOutputData; + + HIP_CHECK(hipMalloc(&devData, Nbytes)); + hOutputData = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hOutputData != nullptr); + memset(hOutputData, 0, Nbytes); + + hipGraph_t graph; + hipError_t ret; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + hipGraphNode_t memsetNode; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + hipMemsetParams mParams{}; + memset(&mParams, 0, sizeof(mParams)); + mParams.dst = reinterpret_cast(devData); + mParams.value = val; + mParams.pitch = 0; + mParams.elementSize = sizeof(char); + mParams.width = Nbytes; + mParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &mParams)); + + std::vector dependencies; + dependencies.push_back(memsetNode); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + SECTION("Pass hGraphExec as nullptr") { + ret = hipGraphExecMemsetNodeSetParams(nullptr, memsetNode, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hGraphNode as nullptr") { + ret = hipGraphExecMemsetNodeSetParams(graphExec, nullptr, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass different hGraphNode which was not used in graphExec") { + hipGraphNode_t memsetNode1{}; + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode1, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass different Graph which was not used in graphExec") { + hipGraph_t graph1; + HIP_CHECK(hipGraphCreate(&graph1, 0)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph1, nullptr, 0, &mParams)); + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK(hipGraphDestroy(graph1)); + } + SECTION("Pass pNodeParams as nullptr") { + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } +#if HT_NVIDIA + SECTION("Pass pNodeParams as empty structure object") { + hipMemsetParams mParmTemp{}; + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParmTemp); + REQUIRE(hipErrorInvalidValue == ret); + } +#endif + SECTION("Pass hipMemsetParams::dst as nullptr") { + mParams.dst = nullptr; + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } +#if HT_NVIDIA + SECTION("Pass hipMemsetParams::element size other than 1, 2, or 4") { + mParams.dst = reinterpret_cast(devData); + mParams.elementSize = 9; + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hipMemsetParams::height as zero") { + mParams.elementSize = sizeof(char); + mParams.height = 0; + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } +#endif + + free(hOutputData); + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/* Test verifies hipGraphExecMemsetNodeSetParams API Functional scenarios. + */ +TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(char); + constexpr size_t val = 0; + constexpr size_t updateVal = 2; + char *devData, *devData1, *hOutputData, *hOutputData1; + + HIP_CHECK(hipMalloc(&devData, Nbytes)); + HIP_CHECK(hipMalloc(&devData1, Nbytes)); + hOutputData = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hOutputData != nullptr); + memset(hOutputData, updateVal, Nbytes); + hOutputData1 = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hOutputData1 != nullptr); + memset(hOutputData1, 0, Nbytes); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + hipGraphNode_t memsetNode; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = val; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + + std::vector dependencies; + dependencies.push_back(memsetNode); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData1); + memsetParams.value = updateVal; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + + REQUIRE(hipSuccess == hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, + &memsetParams)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + HIP_CHECK(hipMemcpy(hOutputData1, devData1, Nbytes, hipMemcpyDeviceToHost)); + HipTest::checkArray(hOutputData, hOutputData1, Nbytes, 1); + + free(hOutputData); + free(hOutputData1); + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipFree(devData1)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} diff --git a/tests/catch/unit/graph/hipGraphGetNodes.cc b/tests/catch/unit/graph/hipGraphGetNodes.cc new file mode 100644 index 0000000000..18b3df1722 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphGetNodes.cc @@ -0,0 +1,219 @@ +/* +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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/** +Testcase Scenarios +------------------ +Functional :: +1) Add nodes to graph and get nodes. Verify the added nodes are present in returned list. +2) Pass nodes as nullptr and verify numNodes returns actual number of nodes added to graph. +3) If numNodes passed is greater than the actual number of nodes, the remaining entries in nodes +will be set to NULL, and the number of nodes actually obtained will be returned in numNodes. + +Argument Validation :: +1) Pass graph as nullptr and verify api returns error code. +2) Pass numNodes as nullptr and other params as valid values. Expect api to return error code. +3) When there are no nodes in graph, expect numNodes to be set to zero. +4) Pass numNodes less than actual number of nodes. Expect api to populate requested number of node entries +and does update numNodes. +*/ + +#include +#include +#include + +/** + * Functional Test for hipGraphGetNodes API fetching node list + */ +TEST_CASE("Unit_hipGraphGetNodes_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + constexpr auto addlEntries = 4; + hipGraph_t graph; + hipGraphNode_t memcpyNode, kernelNode; + hipKernelNodeParams kernelNodeParams{}; + hipStream_t streamForGraph; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + std::vector dependencies, nodelist; + hipGraphExec_t graphExec; + size_t NElem{N}; + + 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, NULL, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + nodelist.push_back(memcpyNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + nodelist.push_back(memcpyNode); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, dependencies.data(), + dependencies.size(), &kernelNodeParams)); + dependencies.clear(); + dependencies.push_back(kernelNode); + nodelist.push_back(kernelNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), + dependencies.size(), C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + nodelist.push_back(memcpyNode); + + // Get numNodes by passing nodes as nullptr. + // verify : numNodes is set to actual number of nodes added + size_t numNodes{}; + HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes)); + INFO("Num of nodes returned by GetNodes : " << numNodes); + REQUIRE(numNodes == nodelist.size()); + + // Request for extra/additional nodes. + // verify : totNodes is reset to actual number of nodes + // verify : additional entries in nodes are set to nullptr + size_t totNodes = numNodes + addlEntries; + int numBytes = sizeof(hipGraphNode_t) * totNodes; + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(nodes != nullptr); + HIP_CHECK(hipGraphGetNodes(graph, nodes, &totNodes)); + REQUIRE(totNodes == nodelist.size()); + for (auto i = numNodes; i < numNodes + addlEntries; i++) { + REQUIRE(nodes[i] == nullptr); + } + + // Verify added nodes are present in the node entries returned + for (auto Node : nodelist) { + bool found = false; + for (size_t i = 0; i < numNodes; i++) { + if (Node == nodes[i]) { + found = true; + break; + } + } + + if (!found) { + INFO("Added node " << Node << " not present in returned list"); + REQUIRE(false); + } + } + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + 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(streamForGraph)); + free(nodes); +} + +/** + * Test performs api parameter validation by passing various values + * as input and output parameters and validates the behavior. + * Test will include both negative and positive scenarios. + */ +TEST_CASE("Unit_hipGraphGetNodes_ParamValidation") { + hipStream_t stream{nullptr}; + hipGraph_t graph{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 1000000; + size_t Nbytes = N * sizeof(float), numNodes{}; + float *A_d, *C_d; + float *A_h, *C_h; + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(C_h != nullptr); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(C_d != nullptr); + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, stream, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes)); + INFO("Num of nodes returned by GetNodes : " << numNodes); + + int numBytes = sizeof(hipGraphNode_t) * numNodes; + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(nodes != nullptr); + + SECTION("graph as nullptr") { + hipError_t ret = hipGraphGetNodes(nullptr, nodes, &numNodes); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("numNodes as nullptr") { + hipError_t ret = hipGraphGetNodes(graph, nodes, nullptr); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("no nodes in graph") { + hipGraph_t emptyGraph{}; + HIP_CHECK(hipGraphCreate(&emptyGraph, 0)); + HIP_CHECK(hipGraphGetNodes(emptyGraph, nullptr, &numNodes)); + REQUIRE(numNodes == 0); + } + + SECTION("numNodes less than actual number of nodes") { + size_t numPartNodes = numNodes - 1; + hipGraphNodeType nodeType; + HIP_CHECK(hipGraphGetNodes(graph, nodes, &numPartNodes)); + + // verify numPartNodes is unchanged + REQUIRE(numPartNodes == numNodes - 1); + // verify partial node list returned has valid nodes + for (size_t i = 0; i < numPartNodes; i++) { + HIP_CHECK(hipGraphNodeGetType(nodes[i], &nodeType)); + REQUIRE(nodeType >= 0); + REQUIRE(nodeType < hipGraphNodeTypeCount); + } + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); + free(A_h); + free(C_h); + free(nodes); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} diff --git a/tests/catch/unit/graph/hipGraphGetRootNodes.cc b/tests/catch/unit/graph/hipGraphGetRootNodes.cc new file mode 100644 index 0000000000..c2484fc8b5 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphGetRootNodes.cc @@ -0,0 +1,239 @@ +/* +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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/** +Testcase Scenarios +------------------ +Functional :: + 1) Add nodes to graph with and without dependencies, verify the api returns list of + root nodes (i.e., nodes without dependencies). + 2) Pass nodes as nullptr and verify api returns actual number of root nodes added to graph. + 3) If NumRootNodes passed is greater than the actual number of root nodes, the remaining entries in + nodes list will be set to NULL, and the number of nodes actually obtained will be returned in NumRootNodes. + +Argument Validation :: + 1) Pass graph as nullptr and verify api returns error code. + 2) Pass numRootNodes as nullptr and other params as valid values. Expect api to return error code. + 3) When there are no nodes in graph, expect numRootNodes to be set to zero. + 4) Pass numRootNodes less than actual number of nodes. Expect api to populate requested number of node entries + and does update numRootNodes. +*/ + +#include +#include +#include + +/** + * Functional Test for API fetching root node list + */ +TEST_CASE("Unit_hipGraphGetRootNodes_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + constexpr auto addlEntries = 5; + hipGraph_t graph; + + hipGraphNode_t memcpyNode, kernelNode; + hipKernelNodeParams kernelNodeParams{}; + hipStream_t streamForGraph; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + std::vector dependencies, rootnodelist; + hipGraphExec_t graphExec; + size_t NElem{N}; + + 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, NULL, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + rootnodelist.push_back(memcpyNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + rootnodelist.push_back(memcpyNode); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, dependencies.data(), + dependencies.size(), &kernelNodeParams)); + dependencies.clear(); + dependencies.push_back(kernelNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), + dependencies.size(), C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + // Get numRootNodes by passing rootnodes list as nullptr. + // verify : numRootNodes is set to actual number of root nodes added + size_t numRootNodes{}; + HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes)); + INFO("Num of nodes returned by GetRootNodes : " << numRootNodes); + REQUIRE(numRootNodes == rootnodelist.size()); + + // Request for extra/additional nodes. + // verify : totNodes is reset to actual number of root nodes present + // verify : additional entries in rootnodes list are set to nullptr + size_t totNodes = numRootNodes + addlEntries; + int numBytes = sizeof(hipGraphNode_t) * totNodes; + hipGraphNode_t* rootnodes = + reinterpret_cast(malloc(numBytes)); + REQUIRE(rootnodes != nullptr); + HIP_CHECK(hipGraphGetRootNodes(graph, rootnodes, &totNodes)); + REQUIRE(totNodes == rootnodelist.size()); + for (auto i = numRootNodes; i < numRootNodes + addlEntries; i++) { + REQUIRE(rootnodes[i] == nullptr); + } + + // Verify added nodes(without dependencies) are present + // in the root nodes fetched. + for (auto Node : rootnodelist) { + bool found = false; + for (size_t i = 0; i < numRootNodes; i++) { + if (Node == rootnodes[i]) { + found = true; + break; + } + } + + if (!found) { + INFO("Returned root node " << Node << " not present in added list"); + REQUIRE(false); + } + } + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + 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(streamForGraph)); + free(rootnodes); +} + +/** + * Test performs api parameter validation by passing various values + * as input and output parameters and validates the behavior. + * Test will include both negative and positive scenarios. + */ +TEST_CASE("Unit_hipGraphGetRootNodes_ParamValidation") { + hipStream_t stream1{nullptr}, stream2{nullptr}, mstream{nullptr}; + hipEvent_t memsetEvent1, memsetEvent2, forkStreamEvent; + hipGraph_t graph{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 1000000; + size_t Nbytes = N * sizeof(float), numRootNodes{}; + float *A_d, *C_d; + float *A_h, *C_h; + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(C_h != nullptr); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(C_d != nullptr); + + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&mstream)); + HIP_CHECK(hipEventCreate(&memsetEvent1)); + HIP_CHECK(hipEventCreate(&memsetEvent2)); + HIP_CHECK(hipEventCreate(&forkStreamEvent)); + HIP_CHECK(hipStreamBeginCapture(mstream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(forkStreamEvent, mstream)); + HIP_CHECK(hipStreamWaitEvent(stream1, forkStreamEvent, 0)); + HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0)); + HIP_CHECK(hipMemsetAsync(A_d, 0, Nbytes, stream1)); + HIP_CHECK(hipEventRecord(memsetEvent1, stream1)); + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream2)); + HIP_CHECK(hipEventRecord(memsetEvent2, stream2)); + HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent1, 0)); + HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent2, 0)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mstream)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, mstream, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream)); + HIP_CHECK(hipStreamEndCapture(mstream, &graph)); + HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes)); + INFO("Num of nodes returned by GetRootNodes : " << numRootNodes); + int numBytes = sizeof(hipGraphNode_t) * numRootNodes; + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(nodes != nullptr); + + SECTION("graph as nullptr") { + hipError_t ret = hipGraphGetRootNodes(nullptr, nodes, &numRootNodes); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("numRootNodes as nullptr") { + hipError_t ret = hipGraphGetRootNodes(graph, nodes, nullptr); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("no nodes in graph") { + hipGraph_t emptyGraph{}; + HIP_CHECK(hipGraphCreate(&emptyGraph, 0)); + HIP_CHECK(hipGraphGetRootNodes(emptyGraph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == 0); + } + + SECTION("numRootNodes less than actual number of nodes") { + size_t numPartNodes = numRootNodes - 1; + hipGraphNodeType nodeType; + HIP_CHECK(hipGraphGetRootNodes(graph, nodes, &numPartNodes)); + + // verify numPartNodes is unchanged + REQUIRE(numPartNodes == numRootNodes - 1); + // verify partial node list returned has valid nodes + for (size_t i = 0; i < numPartNodes; i++) { + HIP_CHECK(hipGraphNodeGetType(nodes[i], &nodeType)); + REQUIRE(nodeType >= 0); + REQUIRE(nodeType < hipGraphNodeTypeCount); + } + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(mstream)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipEventDestroy(memsetEvent1)); + HIP_CHECK(hipEventDestroy(memsetEvent2)); + free(A_h); + free(C_h); + free(nodes); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} diff --git a/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc b/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc new file mode 100644 index 0000000000..b8fabb69ee --- /dev/null +++ b/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc @@ -0,0 +1,304 @@ +/* +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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/* +hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t graph, unsigned long long flags); +Testcase Scenarios of hipGraphInstantiateWithFlags API: + +Negative: +1) Pass nullptr to pGraphExec +2) Pass nullptr to graph +4) Pass invalid flag + +Functional: + +1) Create dependencies graph and instantiate the graph +2) Create graph in one GPU device and instantiate, launch in peer GPU device +3) Create stream capture graph and instantite the graph +4) Create stream capture graph in one GPU device and instantite the graph launch + in peer GPU device + +Mapping is missing for NVIDIA platform hence skipping the testcases +*/ + + +#include +#include +#include + +constexpr size_t N = 1000000; +#if HT_AMD +/* This test covers the negative scenarios of + hipGraphInstantiateWithFlags API */ +TEST_CASE("Unit_hipGraphInstantiateWithFlags_Negative") { +#if HT_NVIDIA + SECTION("Passing nullptr pGraphExec") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + REQUIRE(hipGraphInstantiateWithFlags(nullptr, + graph, 0) == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to graph") { + hipGraphExec_t graphExec; + REQUIRE(hipGraphInstantiateWithFlags(&graphExec, + nullptr, 0) == hipErrorInvalidValue); + } + + SECTION("Passing Invalid flag") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphExec_t graphExec; + REQUIRE(hipGraphInstantiateWithFlags(&graphExec, graph, 10) != hipSuccess); + } +#endif +} +/* +This function verifies the following scenarios +1. Creates dependency graph, Instantiates the graph with flags and verifies it +2. Creates graph on one GPU-1 device and instantiates the graph on peer GPU device +*/ +void GraphInstantiateWithFlags_DependencyGraph(bool ctxt_change = false) { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph; + hipGraphNode_t memset_A, memset_B, memsetKer_C; + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + hipGraphExec_t graphExec; + hipMemsetParams memsetParams{}; + int memsetVal{}; + size_t NElem{N}; + + 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)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, + &memsetParams)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(B_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0, + &memsetParams)); + + void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast(&NElem)}; + kernelNodeParams.func = + reinterpret_cast(HipTest::memsetReverse); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0, + &kernelNodeParams)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + + if (ctxt_change) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipDeviceEnablePeerAccess(0, 0)); + } + // Instantiate and launch the cloned graph + HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, C_h, N); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/* +This function verifies the following scenarios +1. Creates stream capture graph, Instantiates the graph with flags and verifies it +2. Creates graph on one GPU-1 device and instantiates the graph on peer GPU device +*/ +void GraphInstantiateWithFlags_StreamCapture(bool deviceContextChg = false) { + float *A_d, *C_d; + float *A_h, *C_h; + size_t Nbytes = N * sizeof(float); + hipStream_t stream; + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(C_h != nullptr); + + // Fill with Phi + i + for (size_t i = 0; i < N; i++) { + A_h[i] = 1.618f + i; + } + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(C_d != nullptr); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + + HIP_CHECK(hipStreamCreate(&stream)); + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, stream, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + + if (deviceContextChg) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipDeviceEnablePeerAccess(0, 0)); + } + + // Validate end capture is successful + REQUIRE(graph != nullptr); + HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, 0)); + REQUIRE(graphExec != nullptr); + + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + + // Validate the computation + for (size_t i = 0; i < N; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + UNSCOPED_INFO("A and C not matching at " << i); + REQUIRE(false); + } + } + + HIP_CHECK(hipStreamDestroy(stream)); + free(A_h); + free(C_h); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} +/* +This testcase verifies hipGraphInstantiateWithFlags API +by creating dependency graph and instantiate, launching and verifying +the result +*/ +TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraph") { + GraphInstantiateWithFlags_DependencyGraph(); +} +/* +This testcase verifies hipGraphInstantiateWithFlags API +by creating dependency graph on GPU-0 and instantiate, launching and verifying +the result on GPU-1 +*/ +#if HT_NVIDIA +TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraphDeviceCtxtChg") { + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + GraphInstantiateWithFlags_DependencyGraph(true); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} +#endif +/* +This testcase verifies hipGraphInstantiateWithFlags API +by creating capture graph and instantiate, launching and verifying +the result +*/ +TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCapture") { + GraphInstantiateWithFlags_StreamCapture(); +} + +/* +This testcase verifies hipGraphInstantiateWithFlags API +by creating capture graph on GPU-0 and instantiate, launching and verifying +the result on GPU-1 +*/ +TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCaptureDeviceContextChg") { + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + GraphInstantiateWithFlags_StreamCapture(true); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} +#endif diff --git a/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc new file mode 100644 index 0000000000..804c502a0c --- /dev/null +++ b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc @@ -0,0 +1,107 @@ +#include +#include +#include +#define SIZE 256 + +__device__ int globalIn[SIZE], globalOut[SIZE]; +__device__ __constant__ int globalConst[SIZE]; + + +/* This testcase verifies negative scenarios of + hipGraphMemcpyNodeSetParamsToSymbol API */ +TEST_CASE("Unit_hipGraphMemcpyNodeSetParamsToSymbol_Negative") { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}, *B_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipError_t ret; + hipGraphNode_t memcpyToSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + + SECTION("Pass GraphNode as nullptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(nullptr, + HIP_SYMBOL(globalIn), + B_d, Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass symbol ptr as nullptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + nullptr, + B_d, Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidSymbol == ret); + } + SECTION("Pass src ptr as nullptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + nullptr, Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass count as zero") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + B_d, 0, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + + SECTION("Pass count more than allocated size for source and dstn ptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + B_d, Nbytes+8, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass offset+count greater than allocated size") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + B_d, Nbytes, 10, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass same symbol pointer as source ptr and destination ptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass 2 different symbol pointer as source ptr and dstn ptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + HIP_SYMBOL(globalOut), + Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Copy from host ptr to device ptr but pass kind as different") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + A_h, + Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + + HipTest::freeArrays(A_d, B_d, nullptr, A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); +} \ No newline at end of file diff --git a/tests/catch/unit/graph/hipGraphNodeFindInClone.cc b/tests/catch/unit/graph/hipGraphNodeFindInClone.cc new file mode 100644 index 0000000000..174c0ebcb9 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphNodeFindInClone.cc @@ -0,0 +1,241 @@ +/* +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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/* +Testcase Scenarios of hipGraphNodeFindInClone API: + +Negative: + +1) Pass nullptr to graph node +2) pass nullptr to original graph node +3) pass nullptr to clonedGraph +4) Pass original graph in place of the cloned graph +5) Pass invalid originalNode +6) Destroy the graph node in the original graph + and try to get the deleted graph node + from the cloned graph +7) Clone the graph,Add node to Original graph + and try to find the original node in the cloned graph + + +Functional: + +1) Get the graph node from the cloned graph corresponding to the original node +2) Create and clone the graph, modify the original graph and clone the graph again, + then try to find the newly added graph node from the cloned graph + +*/ + +#include +#include +#include +#include + + +/* This test covers the negative scenarios of + hipGraphNodeFindInClone API */ + +TEST_CASE("Unit_hipGraphNodeFindInClone_Negative") { + hipGraph_t graph; + hipGraph_t clonedgraph; + hipGraphNode_t graphnode, newnode; + hipGraphNode_t clonedgraphnode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + + int *A_d, *A_h, *B_d, *B_h; + HipTest::initArrays(&A_d, &B_d, nullptr, &A_h, + &B_h, nullptr, 1024, false); + HIP_CHECK(hipGraphAddMemcpyNode1D(&graphnode, graph, nullptr, 0, A_d, A_h, + 1024, hipMemcpyHostToDevice)); + // Cloned the graph + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&newnode, graph, nullptr, 0, B_d, B_h, + 1024, hipMemcpyHostToDevice)); + + SECTION("Passing nullptr to Cloned graph") { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, graphnode, nullptr) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to original graph") { + REQUIRE(hipGraphNodeFindInClone(nullptr, graphnode, clonedgraph) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to graph node") { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, nullptr, clonedgraph) + == hipErrorInvalidValue); + } +#if HT_NVIDIA + SECTION("Pass uncloned graph") { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, graphnode, graph) + == hipErrorInvalidValue); + } + + SECTION("Destroy the graph node and find in cloned graph") { + HIP_CHECK(hipGraphDestroyNode(graphnode)); + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, graphnode, + clonedgraph) + == hipErrorInvalidValue); + } +#endif + + SECTION("Pass invalid original graphnode") { + hipGraphNode_t unintialized_graphnode{nullptr}; + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, unintialized_graphnode, + graph) + == hipErrorInvalidValue); + } + + SECTION("Find node in cloned graph which is only present in original graph") { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, newnode, + clonedgraph) == hipErrorInvalidValue); + } + + + HipTest::freeArrays(A_d, B_d, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(clonedgraph)); +} + + +void hipGraphNodeFindInClone_Func(bool ModifyOrigGraph = false) { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph, clonedgraph; + hipGraphNode_t memset_A, memset_B, memsetKer_C; + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C, memcpyD2D_C, + memcpyD2H_C_new; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + hipMemsetParams memsetParams{}; + int memsetVal{}; + size_t NElem{N}; + + 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)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, + &memsetParams)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(B_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0, + &memsetParams)); + + void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast(&NElem)}; + kernelNodeParams.func = + reinterpret_cast(HipTest::memsetReverse); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0, + &kernelNodeParams)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + + + if (ModifyOrigGraph) { + // Cloned the graph + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + // Modify Original graph by adding new dependency + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2D_C, graph, nullptr, 0, + C_d, B_d, + Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C_new, graph, nullptr, 0, + C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2D_C, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2D_C, + &memcpyD2H_C_new, 1)); + } + // Cloned the graph + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + hipGraphNode_t clonedgraphnode; + if (ModifyOrigGraph) { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, + memcpyD2H_C_new, clonedgraph) + == hipSuccess); + } else { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, + memcpyH2D_A, clonedgraph) + == hipSuccess); + } + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(clonedgraph)); +} + +TEST_CASE("Unit_hipGraphNodeFindInClone_Functional") { + SECTION("hipGraphNodeFindInClone Basic Functionality") { + hipGraphNodeFindInClone_Func(); + } + SECTION("hipGraphNodeFindInClone Modify Original graph") { + hipGraphNodeFindInClone_Func(true); + } +} diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index b450921321..23bbd07aae 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -35,6 +35,9 @@ set(TEST_SRC hipMemPrefetchAsyncExtTsts.cc hipMemAdviseMmap.cc hipMemCoherencyTst.cc + hipMallocManaged.cc + hipMemRangeGetAttribute.cc + hipHmmOvrSubscriptionTst.cc ) else() set(TEST_SRC @@ -69,6 +72,9 @@ set(TEST_SRC hipMallocManagedFlagsTst.cc hipMemPrefetchAsyncExtTsts.cc hipMemAdviseMmap.cc + hipMallocManaged.cc + hipMemRangeGetAttribute.cc + hipHmmOvrSubscriptionTst.cc ) endif() diff --git a/tests/catch/unit/memory/hipHmmOvrSubscriptionTst.cc b/tests/catch/unit/memory/hipHmmOvrSubscriptionTst.cc new file mode 100644 index 0000000000..de5d48ee4b --- /dev/null +++ b/tests/catch/unit/memory/hipHmmOvrSubscriptionTst.cc @@ -0,0 +1,213 @@ +/* +Copyright (c) 2021-Present 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 Case Description: This test case tests the working of OverSubscription + feature which is part of HMM.*/ + +#include +#ifdef __linux__ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#endif +#include + +#define INIT_VAL 2.5 +#define NUM_ELMS 268435456 // 268435456 * 4 = 1GB +#define ITERATIONS 10 +#define ONE_GB 1024 * 1024 * 1024 + +static void GetTotGpuMem(int *TotMem); +static void DisplayHmmFlgs(int *Signal); +// Kernel function +__global__ void Square(int n, float *x) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) { + x[i] = x[i] + 10; + } +} + +static void OneGBMemTest(int dev) { + int DataMismatch = 0; + float *HmmAG = nullptr; + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + // Testing hipMemAttachGlobal Flag + HIP_CHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float), + hipMemAttachGlobal)); + + // Initializing HmmAG memory + for (int i = 0; i < NUM_ELMS; i++) { + HmmAG[i] = INIT_VAL; + } + + int blockSize = 256; + int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize; + dim3 dimGrid(numBlocks, 1, 1); + dim3 dimBlock(blockSize, 1, 1); + HIP_CHECK(hipSetDevice(dev)); + for (int i = 0; i < ITERATIONS; ++i) { + Square<<>>(NUM_ELMS, HmmAG); + } + HIP_CHECK(hipStreamSynchronize(strm)); + for (int j = 0; j < NUM_ELMS; ++j) { + if (HmmAG[j] != (INIT_VAL + ITERATIONS * 10)) { + DataMismatch++; + break; + } + } + if (DataMismatch != 0) { + WARN("Data Mismatch observed when kernel launched on device: " << dev); + REQUIRE(false); + } + HIP_CHECK(hipFree(HmmAG)); + HIP_CHECK(hipStreamDestroy(strm)); +} + +static void GetTotGpuMem(int *TotMem) { + size_t FreeMem, TotGpuMem; + HIP_CHECK(hipMemGetInfo(&FreeMem, &TotGpuMem)); + TotMem[0] = (TotGpuMem/(ONE_GB)); + TotMem[1] = 1; +} + +static void DisplayHmmFlgs(int *Signal) { + int managed = 0; + WARN("The following are the attribute values related to HMM for" + " device 0:\n"); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); + WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeConcurrentManagedAccess, 0)); + WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccess, 0)); + WARN("hipDeviceAttributePageableMemoryAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0)); + WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" + << managed); + + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, + 0)); + WARN("hipDeviceAttributeManagedMemory: " << managed); + + // Checking for Vega20 or MI100 + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + char *p = NULL; + p = strstr(prop.gcnArchName, "gfx906"); + if (p) { + WARN("This system has MI60 gpu hence OverSubscription test will be"); + WARN(" skipped"); + Signal[2] = 1; + } + p = strstr(prop.gcnArchName, "gfx908"); + if (p) { + WARN("This system has MI100 gpu hence OverSubscription test will be"); + WARN(" skipped"); + Signal[2] = 1; + } + Signal[1] = managed; + Signal[0] = 1; +} + +TEST_CASE("Unit_HMM_OverSubscriptionTst") { + int HmmEnabled = 0; + // The following Shared Mem is to get Max GPU Mem + // The size requested is for three ints + // 1) To get Max GPU Mem in GB + // 2) To Signal parent that req. info is available to consume + // 3) To know if MI60 or MI100 gpu are there in the system + key_t key = ftok("shmTotMem", 66); + int shmid = shmget(key, (3 * sizeof(int)), 0666|IPC_CREAT); + int *TotGpuMem = reinterpret_cast(shmat(shmid, NULL, 0)); + TotGpuMem[0] = 0; TotGpuMem[1] = 0; + // The following function DisplayHmmFlgs() displays the flag values related + // to HMM and also sends us ManagedMemory attribute value + if (fork() == 0) { + DisplayHmmFlgs(TotGpuMem); + exit(1); + } + while (TotGpuMem[0] == 0) { + sleep(2); + } + // The following if block will skip test if either of MI60 or MI100 is found + if (TotGpuMem[2] == 1) { + SUCCEED("Test is skipped!!"); + REQUIRE(true); + } else { + HmmEnabled = TotGpuMem[1]; + + // Re-setting the shared memory values for further usage + TotGpuMem[0] = 0; + TotGpuMem[1] = 0; + + std::list PidLst; + // The following function gets the MaxGpu memory in GBs and also launches + // OverSubscription test + if (HmmEnabled) { + if ((setenv("HSA_XNACK", "1", 1)) != 0) { + WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); + REQUIRE(false); + } + if (fork() == 0) { + GetTotGpuMem(TotGpuMem); + } + while (TotGpuMem[1] == 0) { + sleep(2); + } + int NumGB = TotGpuMem[0], TotalThreads = (NumGB + 10); + WARN("Launching " << TotalThreads); + WARN(" processes to test OverSubscription."); + pid_t pid; + for (int k = 0; k < TotalThreads; ++k) { + pid = fork(); + PidLst.push_back(pid); + if (pid == 0) { + OneGBMemTest(0); + exit(10); + } + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } + int status; + for (pid_t pd : PidLst) { + waitpid(pd, &status, 0); + if (!(WIFEXITED(status))) { + REQUIRE(false); + } + } + } + shmdt(TotGpuMem); + shmctl(shmid, IPC_RMID, NULL); +} diff --git a/tests/catch/unit/memory/hipMallocManaged.cc b/tests/catch/unit/memory/hipMallocManaged.cc index e8bc69997c..d40aa924a1 100644 --- a/tests/catch/unit/memory/hipMallocManaged.cc +++ b/tests/catch/unit/memory/hipMallocManaged.cc @@ -17,12 +17,11 @@ THE SOFTWARE. */ -/* - List of Test cases: - 1) Unit_hipMallocManaged_Basic - 2) Unit_hipMallocManaged_MultiSize - 3) Unit_hipMallocManaged_MultiKrnlHmmAccess - 4) Unit_hipMallocManaged_KrnlWth2MemTypes +/* Test Case Description: + 1) This testcase verifies the hipMallocManaged basic scenario - supported on + all devices + 2) This testcase verifies the hipMallocManaged basic scenario - supported + only on HMM enabled devices */ #include @@ -32,12 +31,6 @@ // Kernel functions -__global__ void KrnlWth2MemTypes(int *Hmm, int *Dptr, size_t n) { - size_t index = blockIdx.x * blockDim.x + threadIdx.x; - for (size_t i = index; i < n; i++) { - Hmm[i] = Dptr[i] + 10; - } -} __global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) { size_t index = blockIdx.x * blockDim.x + threadIdx.x; @@ -64,9 +57,6 @@ __global__ void KrnlWth2MemTypesC(unsigned char *Hmm, unsigned char *Dptr, } } -// The following variable will be used to get the result of computation -// from multiple threads -static bool IfTestPassed = true; static int HmmAttrPrint() { int managed = 0; @@ -93,62 +83,6 @@ static int HmmAttrPrint() { } -static void LaunchKrnl4(size_t NumElms, int InitVal) { - int *Hmm = NULL, *Dptr = NULL, blockSize = 64, DataMismatch = 0; - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, (sizeof(int) * NumElms))); - HIP_CHECK(hipMalloc(&Dptr, (sizeof(int) * NumElms))); - int *Hstptr = reinterpret_cast(new int[NumElms]); - for (size_t i = 0; i < NumElms; ++i) { - Hstptr[i] = InitVal; - } - HIP_CHECK(hipMemcpy(Dptr, Hstptr, (NumElms * sizeof(int)), - hipMemcpyHostToDevice)); - dim3 dimBlock(blockSize, 1, 1); - dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); - KrnlWth2MemTypes<<>>(Hmm, Dptr, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - for (size_t i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - INFO("Data Mismatch observed after the Kernel: KrnlWth2MemTypes!!\n"); - REQUIRE(false); - } - DataMismatch = 0; - KernelMul_MngdMem<<>>(Hmm, Dptr, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // Verifying the result - for (size_t i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal * 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n"); - REQUIRE(false); - } - DataMismatch = 0; - KernelMulAdd_MngdMem<<>>(Hmm, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // Verifying the result - - for (size_t i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal * 10 * 2 + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n"); - REQUIRE(false); - } - delete[] Hstptr; -} - - static size_t N{4 * 1024 * 1024}; static unsigned blocksPerCU{6}; @@ -241,104 +175,3 @@ TEST_CASE("Unit_hipMallocManaged_Advanced") { } } - -// The following test case tests the behavior of kernel with a HMM memory and -// hipMalloc memory - -TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") { - IfTestPassed = true; - int *Hmm = NULL, *Dptr = NULL, InitVal = 123; - size_t NumElms = (1024 * 1024); - int *Hptr = new int[NumElms], blockSize = 64, DataMismatch = 0; - int managed = HmmAttrPrint(); - if (managed == 1) { - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, sizeof(int) * NumElms)); - HIP_CHECK(hipMalloc(&Dptr, sizeof(int) * NumElms)); - for (size_t i = 0; i < NumElms; ++i) { - Hmm[i] = 0; - Hptr[i] = InitVal; - } - HIP_CHECK(hipMemcpy(Dptr, Hptr, sizeof(int) * NumElms, - hipMemcpyHostToDevice)); - dim3 dimBlock(blockSize, 1, 1); - dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); - KrnlWth2MemTypes<<>>(Hmm, Dptr, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // Verifying the results - for (size_t k = 0; k < NumElms; ++k) { - if (Hmm[k] != (InitVal + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - WARN("DataMismatch observed!\n"); - IfTestPassed = false; - } - - HIP_CHECK(hipFree(Hmm)); - HIP_CHECK(hipFree(Dptr)); - delete[] Hptr; - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - -// The following test case tests when the same Hmm memory is used for -// launching multiple different kernels will results in any issue -TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") { - int managed = HmmAttrPrint(); - if (managed) { - int InitVal = 123, NumElms = (1024 * 1024); - LaunchKrnl4(NumElms, InitVal); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - - -// The following test case allocation, host access, device access of HMM -// memory from size 1 to 10KB - -TEST_CASE("Unit_hipMallocManaged_MultiSize") { - IfTestPassed = true; - int managed = HmmAttrPrint(); - if (managed == 1) { - unsigned char *Hmm1 = NULL, *Hmm2 = NULL; - int InitVal = 100, blockSize = 64, DataMismatch = 0; - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - dim3 dimBlock(blockSize, 1, 1); - for (int i = 1; i < (1024*1024); ++i) { - HIP_CHECK(hipMallocManaged(&Hmm1, i)); - HIP_CHECK(hipMallocManaged(&Hmm2, i)); - for (int j = 0; j < i; ++j) { - Hmm1[j] = InitVal; - } - dim3 dimGrid((i + blockSize -1)/blockSize, 1, 1); - KrnlWth2MemTypesC<<>>(Hmm2, Hmm1, i); - HIP_CHECK(hipStreamSynchronize(strm)); - // Verifying the results - for (int k = 0; k < i; ++k) { - if (Hmm2[k] != (InitVal + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - WARN("DataMismatch observed!\n"); - IfTestPassed = false; - } - DataMismatch = 0; - HIP_CHECK(hipFree(Hmm1)); - HIP_CHECK(hipFree(Hmm2)); - REQUIRE(IfTestPassed); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} diff --git a/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc b/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc index e0d13b764e..8ac1c37853 100644 --- a/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc +++ b/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc @@ -262,79 +262,3 @@ TEST_CASE("Unit_hipMallocManaged_AccessMultiStream") { } } -TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") { - int managed = HmmAttrPrint(); - if (managed == 1) { - bool IfTestPassed = true; - hipError_t err; - void *Hmm = NULL; - size_t totalDevMem = 0, freeDevMem = 0; - int NumDevs = 0; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - // Testing allocation of extreme and unusual mem values - for (int i = 0; i < NumDevs; i++) { - HIP_CHECK(hipSetDevice(i)); - HIP_CHECK(hipMemGetInfo(&freeDevMem, &totalDevMem)); - err = hipMallocManaged(&Hmm, 1, hipMemAttachGlobal); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating memory on GPU: " << i); - WARN(" size 1 with"); - WARN(" hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachGlobal); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating max free memory on GPU: " << i); - WARN(" with hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachGlobal); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating max (free - 1) memory on "); - WARN("GPU: " << i); - WARN(" using hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - err = hipMallocManaged(&Hmm, 1, hipMemAttachHost); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating memory size 1 on GPU: " << i); - WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachHost); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating max free memory on GPU: " << i); - WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachHost); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating max (freeDevMem - 1) memory" - " on GPU: " << i); - WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - } - REQUIRE(IfTestPassed); - } else { - SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result"); - } -} diff --git a/tests/catch/unit/memory/hipMemRangeGetAttribute.cc b/tests/catch/unit/memory/hipMemRangeGetAttribute.cc new file mode 100644 index 0000000000..64fc8e6be4 --- /dev/null +++ b/tests/catch/unit/memory/hipMemRangeGetAttribute.cc @@ -0,0 +1,408 @@ +/* +Copyright (c) 2021 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 WARRANNTY 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 +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 +THE SOFTWARE. +*/ + +/* Test Case Description: + Scenario-1: The following function tests the count parameter(last param) to + hipMemRangeGetAttribute api by passing possible extreme values. + Curently the only way to test if count param working properly is to verify + the first parameter of hipMemRangeGetAttribute() api has value 1 stored + + Scenario-2: This test case checks the behavior of hipMemRangeGetAttribute() with + AccessedBy flag is consistent with cuda's counter part + + Scenario-3: Allocate 4 * page size of memory with the flag hipMemAttachGloal. Advise + AccessedBy, ReadMostly and PreferredLocation to first half(2*pageSz) of the + memory and probe the for the flags which are set earlier using + hipMemRangeGetAttribute() api for the full size(4*PageSz). + + + Scenario-4: The following scenarios tests that probing the attributes which are not set + by hipMemAdvise() but being probed using hipMemRangeGetAttribute() should + not result in a crash + + Scenario-5: The following scenario is a simple test which does the following: + Allocate Hmm memory --> hipMemPrefetchAsync() to device 0 and then + probe LastPrefetchLocation attribute using hipMemRangeGetAttribute + + Scenario-6: The following Test Case does negative tests on hipMemRangeGetAttribute()*/ + +#include +#include +#ifdef __linux__ + #include + #include +#endif + +static bool CheckError(hipError_t err, int LineNo) { + if (err == hipSuccess) { + WARN("Error expected but received hipSuccess at line no.:" + << LineNo); + return false; + } else { + return true; + } +} + + +static int HmmAttrPrint() { + int managed = 0; + WARN("The following are the attribute values related to HMM for" + " device 0:\n"); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); + WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeConcurrentManagedAccess, 0)); + WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccess, 0)); + WARN("hipDeviceAttributePageableMemoryAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0)); + WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" + << managed); + + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, + 0)); + WARN("hipDeviceAttributeManagedMemory: " << managed); + return managed; +} + +// The following function tests the count parameter(last param) to +// hipMemRangeGetAttribute api by passing possible extreme values. +// Curently the only way to test if count param working properly is to verify +// the first parameter of hipMemRangeGetAttribute() api has value 1 stored +TEST_CASE("Unit_hipMemRangeGetAttribute_TstCountParam") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + int MEM_SIZE = 4096, RND_NUM = 9999, FLG_READMOSTLY_ENBLD = 1; + bool IfTestPassed = true; + int data = RND_NUM, *devPtr = nullptr; + size_t TotGpuMem, TotGpuFreeMem; + HIP_CHECK(hipMemGetInfo(&TotGpuFreeMem, &TotGpuMem)); + + HIP_CHECK(hipMallocManaged(&devPtr, MEM_SIZE, hipMemAttachGlobal)); + HIP_CHECK(hipMemAdvise(devPtr, MEM_SIZE, hipMemAdviseSetReadMostly, 0)); + HIP_CHECK(hipMemRangeGetAttribute(reinterpret_cast(&data), + sizeof(int), + hipMemRangeAttributeReadMostly, + devPtr, MEM_SIZE)); + if (data != FLG_READMOSTLY_ENBLD) { + WARN("hipMemRangeGetAttribute() api didnt return expected value!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipFree(devPtr)); + HIP_CHECK(hipMallocManaged(&devPtr, TotGpuFreeMem, hipMemAttachGlobal)); + HIP_CHECK(hipMemAdvise(devPtr, TotGpuFreeMem, hipMemAdviseSetReadMostly, + 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, + devPtr, TotGpuFreeMem)); + + if (data != FLG_READMOSTLY_ENBLD) { + WARN("hipMemRangeGetAttribute() api didnt return expected value!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipFree(devPtr)); + HIP_CHECK(hipMallocManaged(&devPtr, (TotGpuFreeMem - 1), + hipMemAttachGlobal)); + HIP_CHECK(hipMemAdvise(devPtr, (TotGpuFreeMem - 1), + hipMemAdviseSetReadMostly, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, + devPtr, (TotGpuFreeMem - 1))); + + if (data != FLG_READMOSTLY_ENBLD) { + WARN("hipMemRangeGetAttribute() api didnt return expected value!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipFree(devPtr)); + + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +/* The following Test Case does negative tests on hipMemRangeGetAttribute()*/ + +TEST_CASE("Unit_hipMemRangeGetAttribute_NegativeTests") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + int MEM_SIZE = 4096, RND_NUM = 9999; + float *devPtr = nullptr; + int NumDevs; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + int data = RND_NUM; + int *OutData = new int[NumDevs]; + for (int m = 0; m < NumDevs; ++m) { + OutData[m] = RND_NUM; + } + HIP_CHECK(hipMallocManaged(&devPtr, MEM_SIZE, hipMemAttachGlobal)); + HIP_CHECK(hipMemAdvise(devPtr, MEM_SIZE, hipMemAdviseSetReadMostly, 0)); + + // checking the behavior with dataSize 0 + SECTION("checking the behavior with dataSize 0") { + REQUIRE(CheckError(hipMemRangeGetAttribute(&data, 0, + hipMemRangeAttributeReadMostly, + devPtr, MEM_SIZE), __LINE__)); + } + // checking the behavior with dataSize > 4 and even + SECTION("checking the behavior with dataSize > 4 and even") { + REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 6, + hipMemRangeAttributeReadMostly, + devPtr, MEM_SIZE), __LINE__)); + } + // checking the behavior with dataSize > 4 and odd + SECTION("checking the behavior with dataSize > 4 and odd") { + REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 7, + hipMemRangeAttributeReadMostly, + devPtr, MEM_SIZE), __LINE__)); + } + // checking the behavior with dataSize which is not multiple of 4 + SECTION("checking the behavior with dataSize which is not multiple of 4") { + REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 27, + hipMemRangeAttributeReadMostly, + devPtr, MEM_SIZE), __LINE__)); + } + // checking the behaviour with devPtr(4th param) as NULL + SECTION("checking the behaviour with devPtr(4th param) as NULL") { + REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, + NULL, MEM_SIZE), __LINE__)); + } + // checking the behaviour with count(5th param) as 0 + SECTION("checking the behaviour with count(5th param) as 0") { + REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, + devPtr, 0), __LINE__)); + } + // checking the behavior with invalid attribute (3rd param) as 0 + // as it is attribute hence avoiding the negative tests with 3rd param + + // checking the behaviour of the api with ptr allocated using + // hipHostMalloc + void *ptr = nullptr; + SECTION("Checking behavior with hipHostMalloc ptr") { + HIP_CHECK(hipHostMalloc(&ptr, MEM_SIZE, 0)); + REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, + ptr, MEM_SIZE), __LINE__)); + HIP_CHECK(hipHostFree(ptr)); + } + HIP_CHECK(hipFree(devPtr)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +/* This test case checks the behavior of hipMemRangeGetAttribute() with + AccessedBy flag is consistent with cuda's counter part*/ +TEST_CASE("Unit_hipMemRangeGetAttribute_AccessedBy1") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int Ngpus = 0, *Hmm = NULL, MEM_SZ = 4096, RND_NUM = 999; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + int *OutData = new int[Ngpus]; + for (int i = 0; i < Ngpus; ++i) { + OutData[Ngpus] = RND_NUM; + } + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SZ)); + HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus, + hipMemRangeAttributeAccessedBy, + Hmm, MEM_SZ)); + if (OutData[0] != 0) { + WARN("Didn't receive expected value at line: " << __LINE__); + REQUIRE(false); + } + for (int i = 1; i < Ngpus; ++i) { + if (OutData[i] != -2) { + WARN("Didn't receive expected value at line: " << __LINE__); + REQUIRE(false); + } + } + if (Ngpus >= 2) { + for (int i = 0; i < Ngpus; ++i) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ, hipMemAdviseSetAccessedBy, i)); + } + // checking the behavior with dataSize less than the number of gpus + // This should not result in segfault. + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*(Ngpus-1), + hipMemRangeAttributeAccessedBy, + Hmm, MEM_SZ)); + // OutData should have stored the gpu ordinals for which AccessedBy is + // assigned except for the last element which should have -2 stored + // so as to be consistent with cuda's behavior + for (int i = 0; i < (Ngpus - 1); ++i) { + if (OutData[i] != i) { + WARN("Didn't receive expected value at line: " << __LINE__); + REQUIRE(false); + } + } + if (OutData[Ngpus - 1] != -2) { + WARN("Didn't receive expected value at line: " << __LINE__); + REQUIRE(false); + } + } + HIP_CHECK(hipFree(Hmm)); + delete[] OutData; + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + + + +/* Allocate 4 * page size of memory with the flag hipMemAttachGloal. Advise + AccessedBy, ReadMostly and PreferredLocation to first half(2*pageSz) of the + memory and probe the for the flags which are set earlier using + hipMemRangeGetAttribute() api for the full size(4*PageSz).*/ +/* Need to discuss the difference in behavior w.r.t cuda*/ + +TEST_CASE("Unit_hipMemRangeGetAttribte_3") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int Ngpus = 0, *Hmm = NULL, MEM_SZ = 4096*4, RND_NUM = 999; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + int *OutData = new int[Ngpus]; + for (int i = 0; i < Ngpus; ++i) { + OutData[Ngpus] = RND_NUM; + } + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SZ)); + HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus, + hipMemRangeAttributeAccessedBy, + (Hmm), MEM_SZ)); + + HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetReadMostly, 0)); + // The Api called below should not fail + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributeReadMostly, + (Hmm), MEM_SZ)); + + HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetPreferredLocation, 0)); + // The api called below should not fail + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributePreferredLocation, + (Hmm), MEM_SZ)); + HIP_CHECK(hipFree(Hmm)); + delete[] OutData; + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* The following scenarios tests that probing the attributes which are not set + by hipMemAdvise() but being probed using hipMemRangeGetAttribute() should + not result in a crash*/ + +TEST_CASE("Unit_hipMemRangeGetAttribute_4") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, PageSz = 4096, Ngpus, RND_NUM = 999; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + int *OutData = new int[Ngpus]; + for (int i = 0; i < Ngpus; ++i) { + OutData[i] = RND_NUM; + } + HIP_CHECK(hipMallocManaged(&Hmm, 4*PageSz)); + SECTION("Set ReadMostly & probe other flags") { + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetReadMostly, 0)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus, + hipMemRangeAttributeAccessedBy, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributePreferredLocation, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetReadMostly, 0)); + } + SECTION("Set AccessedBy & probe other flags") { + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributeReadMostly, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributePreferredLocation, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetAccessedBy, 0)); + } + SECTION("Set AccessedBy & probe other flags") { + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetPreferredLocation, + 0)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributeReadMostly, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus, + hipMemRangeAttributeAccessedBy, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetPreferredLocation, + 0)); + } + HIP_CHECK(hipFree(Hmm)); + delete[] OutData; + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* The following scenario is a simple test which does the following: + Allocate Hmm memory --> hipMemPrefetchAsync() to device 0 and then + probe LastPrefetchLocation attribute using hipMemRangeGetAttribute*/ + +TEST_CASE("Unit_hipMemRangeGetAttribute_PrefetchAndGtAttr") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int Ngpus = 0, *Hmm = NULL, RND_NUM = 999; + size_t PageSz = 4096; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + + int *OutData = new int[Ngpus]; + for (int i = 0; i < Ngpus; ++i) { + OutData[Ngpus] = RND_NUM; + } + HIP_CHECK(hipMallocManaged(&Hmm, PageSz*4)); + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, PageSz*4, 0, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributeLastPrefetchLocation, + Hmm, PageSz*4)); + HIP_CHECK(hipStreamDestroy(strm)); + HIP_CHECK(hipFree(Hmm)); + if (OutData[0] != 0) { + WARN("Didnt receive expected value at line: " << __LINE__); + delete[] OutData; + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + diff --git a/tests/catch/unit/memory/hipMemcpy2DAsync.cc b/tests/catch/unit/memory/hipMemcpy2DAsync.cc index c5733a095e..10924687df 100644 --- a/tests/catch/unit/memory/hipMemcpy2DAsync.cc +++ b/tests/catch/unit/memory/hipMemcpy2DAsync.cc @@ -83,22 +83,40 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", "" // Initialize the data HipTest::setDefaultData(NUM_W*NUM_H, A_h, B_h, C_h); + SECTION("Calling Async apis with stream object created by user") { + // Host to Device + HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType), + COLUMNS*sizeof(TestType), ROWS, + hipMemcpyHostToDevice, stream)); - // Host to Device - HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType), - COLUMNS*sizeof(TestType), ROWS, - hipMemcpyHostToDevice, stream)); + // Performs D2D on same GPU device + HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d, + pitch_A, COLUMNS*sizeof(TestType), + ROWS, hipMemcpyDeviceToDevice, stream)); - // Performs D2D on same GPU device - HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d, - pitch_A, COLUMNS*sizeof(TestType), - ROWS, hipMemcpyDeviceToDevice, stream)); + // hipMemcpy2DAsync Device to Host + HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B, + COLUMNS*sizeof(TestType), ROWS, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling Async apis with hipStreamPerThread") { + // Host to Device + HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType), + COLUMNS*sizeof(TestType), ROWS, + hipMemcpyHostToDevice, hipStreamPerThread)); - // hipMemcpy2DAsync Device to Host - HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B, - COLUMNS*sizeof(TestType), ROWS, - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); + // Performs D2D on same GPU device + HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d, + pitch_A, COLUMNS*sizeof(TestType), + ROWS, hipMemcpyDeviceToDevice, hipStreamPerThread)); + + // hipMemcpy2DAsync Device to Host + HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B, + COLUMNS*sizeof(TestType), ROWS, + hipMemcpyDeviceToHost, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } // Validating the result REQUIRE(HipTest::checkArray(A_h, B_h, COLUMNS, ROWS) == true); diff --git a/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc b/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc index a8c9ac2944..6603243585 100644 --- a/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc +++ b/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc @@ -62,10 +62,18 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Basic") { HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, width, NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); + SECTION("Calling hipMemcpy2DFromArrayAsync() with user declared stream obj") { + HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpy2DFromArrayAsync() with hipStreamPerThread") { + HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); // Cleaning the memory diff --git a/tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc b/tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc index fecc031f52..5145f28aa1 100644 --- a/tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc +++ b/tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc @@ -58,11 +58,18 @@ TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Basic") { HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); HIP_CHECK(hipStreamCreate(&stream)); - - HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, - width, NUM_H, - hipMemcpyHostToDevice, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); + SECTION("Calling hipMemcpy2DToArrayAsync() with user declared stream obj") { + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpy2DToArrayAsync() with hipStreamPerThread") { + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, 0, 0, width, NUM_H, hipMemcpyDeviceToHost)); diff --git a/tests/catch/unit/memory/hipMemcpy3DAsync.cc b/tests/catch/unit/memory/hipMemcpy3DAsync.cc index fb2fa8e1dc..6ebc480844 100644 --- a/tests/catch/unit/memory/hipMemcpy3DAsync.cc +++ b/tests/catch/unit/memory/hipMemcpy3DAsync.cc @@ -602,8 +602,14 @@ void Memcpy3DAsync::simple_Memcpy3DAsync() { #else myparms.kind = hipMemcpyHostToDevice; #endif - REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); - HIP_CHECK(hipStreamSynchronize(stream)); + SECTION("Calling hipMemcpy3DAsync() using user declared stream obj") { + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpy3DAsync() using hipStreamPerThread") { + REQUIRE(hipMemcpy3DAsync(&myparms, hipStreamPerThread) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } // Array to Array memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); diff --git a/tests/catch/unit/memory/hipMemcpyPeerAsync.cc b/tests/catch/unit/memory/hipMemcpyPeerAsync.cc index c2e67bb802..c3a0b31501 100644 --- a/tests/catch/unit/memory/hipMemcpyPeerAsync.cc +++ b/tests/catch/unit/memory/hipMemcpyPeerAsync.cc @@ -149,11 +149,20 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") { // Copying data from GPU-0 to GPU-1 and performing vector addition HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes, - stream)); - HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes, - stream)); - HIP_CHECK(hipStreamSynchronize(stream)); + SECTION("Calling hipMemcpyPerAsync() using user defined stream obj") { + HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes, + stream)); + HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpyPerAsync() using hipStreamPerThread") { + HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes, + hipStreamPerThread)); + HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes, + hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0, static_cast(X_d), static_cast(Y_d), Z_d, numElements*sizeof(int)); diff --git a/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc b/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc index 479e163c26..89390bb908 100644 --- a/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc +++ b/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc @@ -61,7 +61,6 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { size_t elements = NUM_W * NUM_H; unsigned blocks{}; int validateCount{}; - hipStream_t stream; blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, @@ -81,21 +80,42 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { } HIP_CHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipStreamCreate(&stream)); + SECTION("Using User created stream") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + for (size_t k = 0; k < ITER; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, stream, B_d, C_d, elements); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, + hipMemcpyDeviceToHost)); + for (size_t p = 0 ; p < elements ; p++) { + if (A_h[p] == memsetval) { + validateCount+= 1; + } + } + } + HIP_CHECK(hipStreamDestroy(stream)); + } + SECTION("Using hipStreamPerThread stream") { + for (size_t k = 0; k < ITER; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, hipStreamPerThread, B_d, C_d, elements); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, + hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, + hipMemcpyDeviceToHost)); - for (size_t k = 0; k < ITER; k++) { - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, stream, B_d, C_d, elements); - - HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, - hipMemcpyDeviceToHost)); - - for (size_t p = 0 ; p < elements ; p++) { - if (A_h[p] == memsetval) { - validateCount+= 1; + for (size_t p = 0 ; p < elements ; p++) { + if (A_h[p] == memsetval) { + validateCount+= 1; + } } } } @@ -104,7 +124,6 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); free(A_h); free(B_h); - HIP_CHECK(hipStreamDestroy(stream)); } diff --git a/tests/catch/unit/memory/hipMemset3DFunctional.cc b/tests/catch/unit/memory/hipMemset3DFunctional.cc index b47798487b..35d5696bc4 100644 --- a/tests/catch/unit/memory/hipMemset3DFunctional.cc +++ b/tests/catch/unit/memory/hipMemset3DFunctional.cc @@ -184,11 +184,17 @@ static void testMemsetMaxValue(bool bAsync) { HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); if (bAsync) { - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipStreamDestroy(stream)); + SECTION("Using user created stream") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + } + SECTION("Using hipStreamPerThread") { + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } } else { HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); } @@ -236,7 +242,7 @@ static void seekAndSet3DArraySlice(bool bAsync) { // select random slice for memset unsigned int seed = time(nullptr); - int slice_index = HipTest::RAND_R(&seed) % ZSIZE_S; + int slice_index = rand_r(&seed) % ZSIZE_S; INFO("memset3d for sliceindex " << slice_index); diff --git a/tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc b/tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc index adbd4a3964..270561f811 100644 --- a/tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc +++ b/tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc @@ -21,7 +21,6 @@ * Test for checking order of execution of device kernel and * hipMemsetAsync apis on all gpus */ - #include #include #include @@ -83,16 +82,26 @@ class MemSetKernelTest { } }; -static bool testhipMemsetAsyncWithKernel() { +static bool testhipMemsetAsyncWithKernel(bool UseStrmPerThrd) { MemSetKernelTest obj; constexpr char memsetval = 0x42; obj.memAllocate(memsetval); - for (int k = 0 ; k < ITER ; k++) { - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); - HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream)); - HIP_CHECK(hipStreamSynchronize(obj.stream)); + for (int k = 0 ; k < ITER ; ++k) { + if (UseStrmPerThrd) { // will use hipStreamPerThread stream object + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, hipStreamPerThread, obj.B_d, + obj.C_d, N); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal, N, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } else { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, + N); + HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream)); + HIP_CHECK(hipStreamSynchronize(obj.stream)); + } HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); obj.validateExecutionOrder(); @@ -109,7 +118,7 @@ static bool testhipMemsetD32AsyncWithKernel() { hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)obj.C_d , obj.memSetVal, - N, obj.stream)); + N, obj.stream)); HIP_CHECK(hipStreamSynchronize(obj.stream)); HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); @@ -161,7 +170,7 @@ static bool testhipMemsetD8AsyncWithKernel() { */ TEST_CASE("Unit_hipMemsetAsync_VerifyExecutionWithKernel") { int numDevices = 0; - bool ret; + bool ret, UseStrmPerThrd = false; blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); @@ -172,7 +181,13 @@ TEST_CASE("Unit_hipMemsetAsync_VerifyExecutionWithKernel") { HIP_CHECK(hipSetDevice(devNum)); SECTION("hipMemsetAsync With Kernel") { - ret = testhipMemsetAsyncWithKernel(); + UseStrmPerThrd = false; + ret = testhipMemsetAsyncWithKernel(UseStrmPerThrd); + REQUIRE(ret == true); + } + SECTION("hipMemsetAsync With Kernel using hipStreamPerThread stream obj") { + UseStrmPerThrd = true; + ret = testhipMemsetAsyncWithKernel(UseStrmPerThrd); REQUIRE(ret == true); } diff --git a/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_down.cpp b/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_down.cpp new file mode 100644 index 0000000000..aa7b00e290 --- /dev/null +++ b/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_down.cpp @@ -0,0 +1,276 @@ +/* +Copyright (c) 2020 - 2021 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 Description: +/* This test implements sum reduction kernel, first with each threads own rank + as input and comparing the sum with expected sum output derieved from n(n-1)/2 + formula. + This sample tests functionality of intrinsics provided by thread_block_tile type, + shfl_down and shfl_xor. +*/ + +#include "test_common.h" +#include +#include +#include + +using namespace cooperative_groups; + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) +#define WAVE_SIZE 32 + +__device__ int reduction_kernel_shfl_down(coalesced_group const& g, int val) { + int sz = g.size(); + + for (int i = sz / 2; i > 0; i >>= 1) { + val += g.shfl_down(val, i); + } + + // Choose the 0'th indexed thread that holds the reduction value to return + if (g.thread_rank() == 0) { + return val; + } + // Rest of the threads return no useful values + else { + return -1; + } +} + +__global__ void kernel_shfl_down (int * dPtr, int *dResults, int lane_delta, int cg_sizes) { + int id = threadIdx.x + blockIdx.x * blockDim.x; + + if (id % cg_sizes == 0) { + coalesced_group const& g = coalesced_threads(); + int rank = g.thread_rank(); + int val = dPtr[rank]; + dResults[rank] = g.shfl_down(val, lane_delta); + return; + } +} + +__global__ void kernel_cg_group_partition(int* result, unsigned int tileSz, int cg_sizes) { + + int id = threadIdx.x + blockIdx.x * blockDim.x; + if (id % cg_sizes == 0) { + coalesced_group threadBlockCGTy = coalesced_threads(); + int input, outputSum, expectedSum; + + // Choose a leader thread to print the results + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)threadBlockCGTy.size() / tileSz, tileSz); + } + + threadBlockCGTy.sync(); + + coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz); + int threadRank = tiledPartition.thread_rank(); + + input = tiledPartition.thread_rank(); + + // (n-1)(n)/2 + expectedSum = ((tileSz - 1) * tileSz / 2); + + outputSum = reduction_kernel_shfl_down(tiledPartition, input); + + if (tiledPartition.thread_rank() == 0) { + printf( + " Sum of all ranks 0..%d in this tiledPartition group using shfl_down is %d (expected " + "%d)\n", + tiledPartition.size() - 1, outputSum, expectedSum); + result[threadBlockCGTy.thread_rank() / (tileSz)] = outputSum; + } + return; + } +} + +void verifyResults(int* ptr, int expectedResult, int numTiles) { + for (int i = 0; i < numTiles; i++) { + if (ptr[i] != expectedResult) { + printf(" Results do not match! "); + } + } +} + +void compareResults(int* cpu, int* gpu, int size) { + for (unsigned int i = 0; i < size / sizeof(int); i++) { + if (cpu[i] != gpu[i]) { + printf(" results do not match."); + } + } +} + +void printResults(int* ptr, int size) { + for (int i = 0; i < size; i++) { + std::cout << ptr[i] << " "; + } + std::cout << '\n'; +} + +static void test_group_partition(unsigned int tileSz) { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = 32; + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + int numTiles = ((blockSize * threadsPerBlock) / i) / tileSz; + int expectedSum = ((tileSz - 1) * tileSz / 2); + int* expectedResult = new int[numTiles]; + + // numTiles = 0 when partitioning is possible. The below statement is to avoid + // out-of-bounds error and still evaluate failure case. + numTiles = (numTiles == 0) ? 1 : numTiles; + + for (int i = 0; i < numTiles; i++) { + expectedResult[i] = expectedSum; + } + + int* dResult = NULL; + int* hResult = NULL; + + hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault); + memset(hResult, 0, numTiles * sizeof(int)); + + hipMalloc(&dResult, numTiles * sizeof(int)); + + + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dResult, tileSz, i); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + + + hipMemcpy(hResult, dResult, sizeof(int) * numTiles, hipMemcpyDeviceToHost); + + verifyResults(hResult, expectedSum, numTiles); + + // Free all allocated memory on host and device + hipFree(dResult); + hipFree(hResult); + delete[] expectedResult; + + printf("\n...PASSED.\n\n"); + } +} + +static void test_shfl_down() { + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + int totalThreads = blockSize * threadsPerBlock; + int group_size = totalThreads / i; + int group_size_in_bytes = group_size * sizeof(int); + + int* hPtr = NULL; + int* dPtr = NULL; + int* dResults = NULL; + int lane_delta = rand() % group_size; + std::cout << "Testing coalesced_groups shfl_down with lane_delta " << lane_delta << "and group size " + << WAVE_SIZE << '\n' << std::endl; + + int arrSize = blockSize * threadsPerBlock * sizeof(int); + + hipHostMalloc(&hPtr, arrSize); + // Fill up the array + for (int i = 0; i < WAVE_SIZE; i++) { + hPtr[i] = rand() % 1000; + } + + int* cpuResultsArr = (int*)malloc(group_size_in_bytes); + for (int i = 0; i < group_size; i++) { + cpuResultsArr[i] = (i + lane_delta >= group_size) ? hPtr[i] : hPtr[i + lane_delta]; + } + //printf("Array passed to GPU for computation\n"); + //printResults(hPtr, WAVE_SIZE); + hipMalloc(&dPtr, group_size_in_bytes); + hipMalloc(&dResults, group_size_in_bytes); + + hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice); + // Launch Kernel + hipLaunchKernelGGL(kernel_shfl_down, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i); + hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + //printf("GPU results: \n"); + //printResults(hPtr, WAVE_SIZE); + //printf("Printing cpu to be verified array\n"); + //printResults(cpuResultsArr, WAVE_SIZE); + + compareResults(hPtr, cpuResultsArr, group_size_in_bytes); + std::cout << "Results verified!\n"; + + hipFree(hPtr); + hipFree(dPtr); + free(cpuResultsArr); + } +} + + + +int main() { + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + if (!deviceProperties.cooperativeLaunch) { + std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n"; + if (hip_skip_tests_enabled()) { + return hip_skip_retcode(); + } else { + passed(); + } + return 0; + } + + // Test shfl_down with random group sizes + for (int i = 0; i < 100; i++) { + test_shfl_down(); + } + + std::cout << "Testing static tiled_partition for different tile sizes using shfl_down" + << std::endl; + + int testNo = 1; + std::vector tileSizes = {2, 4, 8, 16, 32}; + for (auto i : tileSizes) { + std::cout << "TEST " << testNo << ":" << '\n' << std::endl; + test_group_partition(i); + testNo++; + } + + passed(); +} diff --git a/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_up.cpp b/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_up.cpp new file mode 100644 index 0000000000..62a05ec244 --- /dev/null +++ b/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_up.cpp @@ -0,0 +1,260 @@ +/* +Copyright (c) 2020 - 2021 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 Description: +/* This test implements prefix sum(scan) kernel, first with each threads own rank + as input and comparing the sum with expected serial summation output on CPU. + + This sample tests functionality of intrinsics provided by coalesced_group, + shfl_up. +*/ +#include "test_common.h" +#include +#include +#include + +using namespace cooperative_groups; + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) +#define WAVE_SIZE 32 +__device__ int prefix_sum_kernel(coalesced_group const& g, int val) { + int sz = g.size(); + for (int i = 1; i < sz; i <<= 1) { + int temp = g.shfl_up(val, i); + + if (g.thread_rank() >= i) { + val += temp; + } + } + return val; +} + +__global__ void kernel_shfl_up (int * dPtr, int *dResults, int lane_delta, int cg_sizes) { + int id = threadIdx.x + blockIdx.x * blockDim.x; + + if (id % cg_sizes == 0) { + coalesced_group g = coalesced_threads(); + int rank = g.thread_rank(); + int val = dPtr[rank]; + dResults[rank] = g.shfl_up(val, lane_delta); + return; + } + +} + +__global__ void kernel_cg_group_partition(int* dPtr, unsigned int tileSz, int cg_sizes) { + + int id = threadIdx.x + blockIdx.x * blockDim.x; + if (id % cg_sizes == 0) { + coalesced_group threadBlockCGTy = coalesced_threads(); + int input, outputSum; + + // we pass its own thread rank as inputs + input = threadBlockCGTy.thread_rank(); + + // Choose a leader thread to print the results + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)threadBlockCGTy.size() / tileSz, tileSz); + } + + threadBlockCGTy.sync(); + + coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz); + + input = tiledPartition.thread_rank(); + + outputSum = prefix_sum_kernel(tiledPartition, input); + + // Update the result array with the corresponsing prefix sum + dPtr[threadBlockCGTy.thread_rank()] = outputSum; + return; + } +} + +void serialScan(int* ptr, int size) { + // Fill up the array + for (int i = 0; i < size; i++) { + ptr[i] = i; + } + + int acc = 0; + for (int i = 0; i < size; i++) { + acc = acc + ptr[i]; + ptr[i] = acc; + } +} + +void printResults(int* ptr, int size) { + for (int i = 0; i < size; i++) { + std::cout << ptr[i] << " "; + } + std::cout << '\n'; +} + +void verifyResults(int* cpu, int* gpu, int size) { + for (unsigned int i = 0; i < size / sizeof(int); i++) { + if (cpu[i] != gpu[i]) { + printf(" Results do not match."); + } + } +} + +static void test_group_partition(unsigned tileSz) { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + int* hPtr = NULL; + int* dPtr = NULL; + int* cpuPrefixSum = NULL; + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + int arrSize = blockSize * threadsPerBlock * sizeof(int); + + hipHostMalloc(&hPtr, arrSize); + hipMalloc(&dPtr, arrSize); + + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dPtr, tileSz, i); + hipMemcpy(hPtr, dPtr, arrSize, hipMemcpyDeviceToHost); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + + cpuPrefixSum = new int[tileSz]; + serialScan(cpuPrefixSum, tileSz); + //std::cout << "\nPrefix sum results on CPU\n"; + //printResults(cpuPrefixSum, tileSz); + + //std::cout << "\nPrefix sum results on GPU\n"; + //printResults(hPtr, tileSz); + std::cout << "\n"; + verifyResults(hPtr, cpuPrefixSum, tileSz); + std::cout << "Results verified!\n"; + + delete[] cpuPrefixSum; + hipFree(hPtr); + hipFree(dPtr); + } +} + +static void test_shfl_up() { + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + hipError_t err; + int blockSize = 1; + + int threadsPerBlock = WAVE_SIZE; + int totalThreads = blockSize * threadsPerBlock; + int group_size = totalThreads / i; + int group_size_in_bytes = group_size * sizeof(int); + + int* hPtr = NULL; + int* dPtr = NULL; + int* dResults = NULL; + int lane_delta = (rand() % group_size); + + std::cout << "Testing coalesced_groups shfl_up with lane_delta " << lane_delta + << " and group size " << WAVE_SIZE << '\n' << std::endl; + + int arrSize = blockSize * threadsPerBlock * sizeof(int); + + hipHostMalloc(&hPtr, arrSize); + // Fill up the array + for (int i = 0; i < WAVE_SIZE; i++) { + hPtr[i] = rand() % 1000; + } + //printResults(hPtr, WAVE_SIZE); + + int* cpuResultsArr = (int*)malloc(group_size_in_bytes); + for (int i = 0; i < group_size; i++) { + cpuResultsArr[i] = (i <= (lane_delta - 1)) ? hPtr[i] : hPtr[i - lane_delta]; + } + + //printf("Printing cpu results arr\n"); + //printResults(cpuResultsArr, WAVE_SIZE); + + hipMalloc(&dPtr, group_size_in_bytes); + hipMalloc(&dResults, group_size_in_bytes); + + hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice); + // Launch Kernel + hipLaunchKernelGGL(kernel_shfl_up, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i); + hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + //printf("GPU computation array :\n"); + //printResults(hPtr, WAVE_SIZE); + + verifyResults(hPtr, cpuResultsArr, group_size_in_bytes); + std::cout << "Results verified!\n"; + + hipFree(hPtr); + hipFree(dPtr); + free(cpuResultsArr); + } +} + +int main() { + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + if (!deviceProperties.cooperativeLaunch) { + std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n"; + if (hip_skip_tests_enabled()) { + return hip_skip_retcode(); + } else { + passed(); + } + return 0; + } + + for (int i = 0; i < 100; i++) { + test_shfl_up(); + } + + std::cout << "Testing coalesced_groups partitioning and shfl_up" << '\n' << std::endl; + + int testNo = 1; + std::vector tileSizes = {2, 4, 8, 16, 32}; + for (auto i : tileSizes) { + std::cout << "TEST " << testNo << ":" << '\n' << std::endl; + test_group_partition(i); + testNo++; + } + passed(); +} + +/* Kogge-Stone algorithm */ \ No newline at end of file diff --git a/tests/src/runtimeApi/cooperativeGrps/simple_coalesced_groups.cpp b/tests/src/runtimeApi/cooperativeGrps/simple_coalesced_groups.cpp new file mode 100644 index 0000000000..9f2775b5dd --- /dev/null +++ b/tests/src/runtimeApi/cooperativeGrps/simple_coalesced_groups.cpp @@ -0,0 +1,583 @@ +/* +Copyright (c) 2020 - 2021 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 Description: +/* This test implements sum reduction kernel, first with each threads own rank + as input and comparing the sum with expected sum output derieved from n(n-1)/2 + formula. The second part, partitions this parent group into child subgroups + a.k.a tiles using using tiled_partition() collective operation. This can be called + with a static tile size, passed in templated non-type variable-tiled_partition, + or in runtime as tiled_partition(thread_group parent, tileSz). This test covers both these + cases. + This test tests functionality of cg group partitioning, (static and dynamic) and its respective + API's size(), thread_rank(), and sync(). +*/ + +#include "test_common.h" +#include +#include +#include + +using namespace cooperative_groups; + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +#define NUM_ELEMS 10000000 +#define NUM_THREADS_PER_BLOCK 512 +#define WAVE_SIZE 32 + +/* Test coalesced group's functionality. + * + */ + +__device__ int atomicAggInc(int *ptr) { + coalesced_group g = coalesced_threads(); + int prev; + // elect the first active thread to perform atomic add + if (g.thread_rank() == 0) { + prev = atomicAdd(ptr, g.size()); + } + // broadcast previous value within the warp + // and add each active thread’s rank to it + prev = g.thread_rank() + g.shfl(prev, 0); + return prev; +} + +__global__ void kernel_shfl (int * dPtr, int *dResults, int srcLane, int cg_sizes) { + int id = threadIdx.x + blockIdx.x * blockDim.x; + if (id % cg_sizes == 0) { + coalesced_group const& g = coalesced_threads(); + int rank = g.thread_rank(); + int val = dPtr[rank]; + dResults[rank] = g.shfl(val, srcLane); + return; + } +} + +__global__ void kernel_shfl_any_to_any (int *randVal, int *dsrcArr, int *dResults, int cg_sizes) { + + int id = threadIdx.x + blockIdx.x * blockDim.x; + + if (id % cg_sizes == 0) { + coalesced_group const& g = coalesced_threads(); + int rank = g.thread_rank(); + int val = randVal[rank]; + dResults[rank] = g.shfl(val, dsrcArr[rank]); + return; + } + +} + +__global__ void filter_arr(int *dst, int *nres, const int *src, int n) { + int id = threadIdx.x + blockIdx.x * blockDim.x; + + for (int i = id; i < n; i += gridDim.x * blockDim.x) { + if (src[i] > 0) dst[atomicAggInc(nres)] = src[i]; + } +} + +/* Parallel reduce kernel. + * + * Step complexity: O(log n) + * Work complexity: O(n) + * + * Note: This kernel works only with power of 2 input arrays. + */ +__device__ int reduction_kernel(coalesced_group g, int* x, int val) { + int lane = g.thread_rank(); + int sz = g.size(); + + for (int i = g.size() / 2; i > 0; i /= 2) { + // use lds to store the temporary result + x[lane] = val; + // Ensure all the stores are completed. + g.sync(); + + if (lane < i) { + val += x[lane + i]; + } + // It must work on one tiled thread group at a time, + // and it must make sure all memory operations are + // completed before moving to the next stride. + // sync() here just does that. + g.sync(); + } + + // Choose the 0'th indexed thread that holds the reduction value to return + if (g.thread_rank() == 0) { + return val; + } + // Rest of the threads return no useful values + else { + return -1; + } +} + +__global__ void kernel_cg_coalesced_group_partition(unsigned int tileSz, int* result, + bool isGlobalMem, int* globalMem, int cg_sizes) { + + int id = threadIdx.x + blockIdx.x * blockDim.x; + if (id % cg_sizes == 0) { + coalesced_group threadBlockCGTy = coalesced_threads(); + int threadBlockGroupSize = threadBlockCGTy.size(); + + int* workspace = NULL; + + if (isGlobalMem) { + workspace = globalMem; + } else { + // Declare a shared memory + extern __shared__ int sharedMem[]; + workspace = sharedMem; + } + + int input, outputSum, expectedOutput; + + // input to reduction, for each thread, is its' rank in the group + input = threadBlockCGTy.thread_rank(); + + expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2; + + outputSum = reduction_kernel(threadBlockCGTy, workspace, input); + + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Sum of all ranks 0..%d in coalesced_group is %d\n\n", + (int)threadBlockCGTy.size() - 1, outputSum); + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)threadBlockCGTy.size() / tileSz, tileSz); + } + + threadBlockCGTy.sync(); + + coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz); + + // This offset allows each group to have its own unique area in the workspace array + int workspaceOffset = threadBlockCGTy.thread_rank() - tiledPartition.thread_rank(); + + outputSum = reduction_kernel(tiledPartition, workspace + workspaceOffset, input); + + if (tiledPartition.thread_rank() == 0) { + printf( + " Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread " + "rank: %d\n", + tiledPartition.size() - 1, outputSum, input); + + result[input / (tileSz)] = outputSum; + } + return; + } +} + +__global__ void kernel_coalesced_active_groups() { + thread_block threadBlockCGTy = this_thread_block(); + int threadBlockGroupSize = threadBlockCGTy.size(); + + // input to reduction, for each thread, is its' rank in the group + int input = threadBlockCGTy.thread_rank(); + + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Creating odd and even set of active thread groups based on branch divergence\n\n"); + } + + threadBlockCGTy.sync(); + + // Group all active odd threads + if (threadBlockCGTy.thread_rank() % 2) { + coalesced_group activeOdd = coalesced_threads(); + + if (activeOdd.thread_rank() == 0) { + printf(" ODD: Size of odd set of active threads is %d." + " Corresponding parent thread_rank is %d.\n\n", + activeOdd.size(), threadBlockCGTy.thread_rank()); + } + } + else { // Group all active even threads + coalesced_group activeEven = coalesced_threads(); + + if (activeEven.thread_rank() == 0) { + printf(" EVEN: Size of even set of active threads is %d." + " Corresponding parent thread_rank is %d.", + activeEven.size(), threadBlockCGTy.thread_rank()); + } + } + return; +} + +void printResults(int* ptr, int size) { + for (int i = 0; i < size; i++) { + std::cout << ptr[i] << " "; + } + std::cout << '\n'; +} + +void compareResults(int* cpu, int* gpu, int size) { + for (unsigned int i = 0; i < size / sizeof(int); i++) { + if (cpu[i] != gpu[i]) { + printf(" results do not match."); + } + } +} + +static void test_active_threads_grouping() { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + // Launch Kernel + hipLaunchKernelGGL(kernel_coalesced_active_groups, blockSize, threadsPerBlock, 0, 0); + + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + printf("\n...PASSED.\n\n"); +} + +// Search if the sum exists in the expected results array +void verifyResults(int* hPtr, int* dPtr, int size) { + int i = 0, j = 0; + for (i = 0; i < size; i++) { + for (j = 0; j < size; j++) { + if (hPtr[i] == dPtr[j]) { + break; + } + } + if (j == size) { + printf(" Result verification failed!"); + } + } +} + + +static void test_group_partition(unsigned int tileSz, bool useGlobalMem) { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + int numTiles = ((blockSize * threadsPerBlock) / i) / tileSz; + + // numTiles = 0 when partitioning is possible. The below statement is to avoid + // out-of-bounds error and still evaluate failure case. + numTiles = (numTiles == 0) ? 1 : numTiles; + + // Build an array of expected reduction sum output on the host + // based on the sum of their respective thread ranks to use for verification + int* expectedSum = new int[numTiles]; + int temp = 0, sum = 0; + for (int i = 1; i <= numTiles; i++) { + sum = temp; + temp = (((tileSz * i) - 1) * (tileSz * i)) / 2; + expectedSum[i-1] = temp - sum; + } + + int* dResult = NULL; + hipMalloc(&dResult, sizeof(int) * numTiles); + + int* globalMem = NULL; + if (useGlobalMem) { + hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int)); + } + + int* hResult = NULL; + hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault); + memset(hResult, 0, numTiles * sizeof(int)); + + // Launch Kernel + if (useGlobalMem) { + hipLaunchKernelGGL(kernel_cg_coalesced_group_partition, blockSize, threadsPerBlock, 0, 0, tileSz, + dResult, useGlobalMem, globalMem, i); + + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + } else { + hipLaunchKernelGGL(kernel_cg_coalesced_group_partition, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, tileSz, dResult, useGlobalMem, globalMem, i); + + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + } + + hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost); + verifyResults(expectedSum, hResult, numTiles); + // Free all allocated memory on host and device + hipFree(dResult); + hipFree(hResult); + if (useGlobalMem) { + hipFree(globalMem); + } + delete[] expectedSum; + + printf("\n...PASSED.\n\n"); + } +} +static void test_shfl_any_to_any() { + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + int totalThreads = blockSize * threadsPerBlock; + int group_size = (totalThreads + i - 1) / i; + int group_size_in_bytes = group_size * sizeof(int); + + int* hPtr = NULL; + int* dPtr = NULL; + int* dsrcArr = NULL; + int* dResults = NULL; + int* srcArr = (int*)malloc(group_size_in_bytes); + int* srcArrCpu = (int*)malloc(group_size_in_bytes); + + std::cout << "Testing coalesced_groups shfl any-to-any\n" < cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + int totalThreads = blockSize * threadsPerBlock; + int group_size = (totalThreads + i - 1) / i; + int group_size_in_bytes = group_size * sizeof(int); + + int* hPtr = NULL; + int* dPtr = NULL; + int* dResults = NULL; + int srcLane = rand() % 1000; + int srcLaneCpu = 0; + std::cout << "Testing coalesced_groups shfl with srcLane " << srcLane << '\n' + << " and group size " << i <(malloc(sizeof(int) * NUM_ELEMS)); + + // Generate input data. + for (int i = 0; i < NUM_ELEMS; i++) { + data_to_filter[i] = rand() % numOfBuckets; + } + + + hipMalloc(&d_data_to_filter, sizeof(int) * NUM_ELEMS); + hipMalloc(&d_filtered_data, sizeof(int) * NUM_ELEMS); + hipMalloc(&d_nres, sizeof(int)); + + hipMemcpy(d_data_to_filter, data_to_filter, + sizeof(int) * NUM_ELEMS, hipMemcpyHostToDevice); + hipMemset(d_nres, 0, sizeof(int)); + + dim3 dimBlock(NUM_THREADS_PER_BLOCK, 1, 1); + dim3 dimGrid((NUM_ELEMS / NUM_THREADS_PER_BLOCK) + 1, 1, 1); + + filter_arr<<>>(d_filtered_data, d_nres, d_data_to_filter, + NUM_ELEMS); + + + hipMemcpy(&nres, d_nres, sizeof(int), hipMemcpyDeviceToHost); + + filtered_data = reinterpret_cast(malloc(sizeof(int) * nres)); + + hipMemcpy(filtered_data, d_filtered_data, sizeof(int) * nres, + hipMemcpyDeviceToHost); + + int *host_filtered_data = + reinterpret_cast(malloc(sizeof(int) * NUM_ELEMS)); + + // Generate host output with host filtering code. + int host_flt_count = 0; + for (int i = 0; i < NUM_ELEMS; i++) { + if (data_to_filter[i] > 0) { + host_filtered_data[host_flt_count++] = data_to_filter[i]; + } + } + + printf("\nWarp Aggregated Atomics %s \n", + (host_flt_count == nres) ? "PASSED" : "FAILED"); + + // Now, testing shfl collective + std::cout << "Now testing shfl collective as a broadcast" << '\n' << std::endl; + + for (int i = 0; i < 100; i++) { + test_shfl_broadcast(); + } + + + // Now, testing shfl collective + std::cout << "Now testing shfl operations any-to-any member lanes" << '\n' << std::endl; + + for (int i = 0; i < 100; i++) { + test_shfl_any_to_any(); + } + + // Now, pass a already coalesced_group that was partitioned + /* Test coalesced group partitioning */ + std::cout << "Now testing coalesced_groups partitioning" << '\n' << std::endl; + + int testNo = 1; + for (int memTy = 0; memTy < 2; memTy++) { + std::vector tileSizes = {2, 4, 8, 16, 32}; + for (auto i : tileSizes) { + std::cout << "TEST " << testNo << ":" << '\n' << std::endl; + test_group_partition(i, memTy); + testNo++; + } + } + + std::cout << "Now grouping active threads based on branch divergence" << '\n' << std::endl; + test_active_threads_grouping(); + + passed(); + return 0; +} diff --git a/tests/src/runtimeApi/graph/hipGraph.cpp b/tests/src/runtimeApi/graph/hipGraph.cpp index 427e583181..fc32512208 100644 --- a/tests/src/runtimeApi/graph/hipGraph.cpp +++ b/tests/src/runtimeApi/graph/hipGraph.cpp @@ -105,6 +105,25 @@ bool hipWithoutGraphs(float* inputVec_h, float* inputVec_d, double* outputVec_d, return true; } +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 = (callBackData_t*)(data); + // checkCudaErrors(tmp->status); + double* result = (double*)(tmp->data); + char* function = (char*)(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 +} + bool hipGraphsUsingStreamCapture(float* inputVec_h, float* inputVec_d, double* outputVec_d, double* result_d, size_t inputSize, size_t numOfBlocks) { hipStream_t stream1, stream2, stream3, streamForGraph; @@ -237,6 +256,16 @@ bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d, nodeDependencies.clear(); nodeDependencies.push_back(memcpyNode); hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0}; + hostParams.fn = myHostNodeCallback; + callBackData_t hostFnData; + hostFnData.data = &result_h; + hostFnData.fn_name = "hipGraphsManual"; + hostParams.userData = &hostFnData; + + HIPCHECK(hipGraphAddHostNode(&hostNode, graph, nodeDependencies.data(), nodeDependencies.size(), + &hostParams)); + hipGraphExec_t graphExec; hipGraphNode_t* nodes = NULL; size_t numNodes = 0; @@ -266,8 +295,8 @@ bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d, for (int i = 0; i < inputSize; i++) { result_h_cpu += inputVec_h[i]; } - if (result_h_cpu != result_h) { - printf("Final reduced sum = %lf %lf\n", result_h_cpu, result_h); + if (result_h_cpu != result_gpu) { + printf("Final reduced sum = %lf %lf\n", result_h_cpu, result_gpu); return false; } return true; @@ -304,4 +333,4 @@ int main(int argc, char** argv) { failed("Failed during hipGraph with capture\n"); } passed(); -} \ No newline at end of file +}