diff --git a/projects/hip-tests/catch/unit/device/CMakeLists.txt b/projects/hip-tests/catch/unit/device/CMakeLists.txt index 4daa65f2a7..59f64028f6 100644 --- a/projects/hip-tests/catch/unit/device/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/device/CMakeLists.txt @@ -33,6 +33,9 @@ set(TEST_SRC if(HIP_PLATFORM MATCHES "amd" AND BUILD_SHARED_LIBS) set(AMD_SRC hipGetProcAddressDevMgmt.cc) + if(UNIX) + set(AMD_SRC ${AMD_SRC} hipGetProcAddressIpcApis.cc) + endif() set(TEST_SRC ${TEST_SRC} ${AMD_SRC}) endif() diff --git a/projects/hip-tests/catch/unit/device/hipGetProcAddressDevMgmt.cc b/projects/hip-tests/catch/unit/device/hipGetProcAddressDevMgmt.cc index 6d373c672d..bd9196b7f9 100644 --- a/projects/hip-tests/catch/unit/device/hipGetProcAddressDevMgmt.cc +++ b/projects/hip-tests/catch/unit/device/hipGetProcAddressDevMgmt.cc @@ -19,6 +19,7 @@ THE SOFTWARE. #include #include #include +#include /** * @addtogroup hipGetProcAddress hipGetProcAddress @@ -50,29 +51,33 @@ void CreateMemPool(int device, hipMemPool_t &mem_pool) { * - HIP_VERSION >= 6.2 */ TEST_CASE("Unit_hipGetProcAddress_ValidateDeviceApis") { - void* hipGetDeviceCount_ptr; - void* hipRuntimeGetVersion_ptr; - void* hipDeviceGetLimit_ptr; - void* hipDeviceSetLimit_ptr; - void* hipDeviceComputeCapability_ptr; - void* hipDeviceGet_ptr; - void* hipDeviceGetPCIBusId_ptr; - void* hipDeviceGetByPCIBusId_ptr; - void* hipDeviceGetDefaultMemPool_ptr; - void* hipDeviceGetName_ptr; - void* hipDeviceGetUuid_ptr; - void* hipGetDeviceFlags_ptr; - void* hipSetDeviceFlags_ptr; - void* hipDeviceReset_ptr; - void* hipDriverGetVersion_ptr; - void* hipDeviceGetCacheConfig_ptr; - void* hipDeviceSetCacheConfig_ptr; - void* hipDeviceTotalMem_ptr; - void* hipGetDeviceProperties_ptr; - void* hipChooseDevice_ptr; - void* hipDeviceSetSharedMemConfig_ptr; - void* hipDeviceGetSharedMemConfig_ptr; - void* hipDeviceGetAttribute_ptr; + void* hipGetDeviceCount_ptr = nullptr; + void* hipRuntimeGetVersion_ptr = nullptr; + void* hipDeviceGetLimit_ptr = nullptr; + void* hipDeviceSetLimit_ptr = nullptr; + void* hipDeviceComputeCapability_ptr = nullptr; + void* hipDeviceGet_ptr = nullptr; + void* hipDeviceGetPCIBusId_ptr = nullptr; + void* hipDeviceGetByPCIBusId_ptr = nullptr; + void* hipDeviceGetDefaultMemPool_ptr = nullptr; + void* hipDeviceGetName_ptr = nullptr; + void* hipDeviceGetUuid_ptr = nullptr; + void* hipGetDeviceFlags_ptr = nullptr; + void* hipSetDeviceFlags_ptr = nullptr; + void* hipDeviceReset_ptr = nullptr; + void* hipDriverGetVersion_ptr = nullptr; + void* hipDeviceGetCacheConfig_ptr = nullptr; + void* hipDeviceSetCacheConfig_ptr = nullptr; + void* hipDeviceTotalMem_ptr = nullptr; + void* hipGetDeviceProperties_ptr = nullptr; + void* hipGetDevicePropertiesR0000_ptr = nullptr; + void* hipGetDevicePropertiesR0600_ptr = nullptr; + void* hipChooseDevice_ptr = nullptr; + void* hipChooseDeviceR0000_ptr = nullptr; + void* hipChooseDeviceR0600_ptr = nullptr; + void* hipDeviceSetSharedMemConfig_ptr = nullptr; + void* hipDeviceGetSharedMemConfig_ptr = nullptr; + void* hipDeviceGetAttribute_ptr = nullptr; int currentHipVersion = 0; HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); @@ -134,8 +139,20 @@ TEST_CASE("Unit_hipGetProcAddress_ValidateDeviceApis") { HIP_CHECK(hipGetProcAddress("hipGetDeviceProperties", &hipGetDeviceProperties_ptr, currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipGetDevicePropertiesR0000", + &hipGetDevicePropertiesR0000_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipGetDevicePropertiesR0600", + &hipGetDevicePropertiesR0600_ptr, + currentHipVersion, 0, nullptr)); HIP_CHECK(hipGetProcAddress("hipChooseDevice", &hipChooseDevice_ptr, currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipChooseDeviceR0000", + &hipChooseDeviceR0000_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipChooseDeviceR0600", + &hipChooseDeviceR0600_ptr, + currentHipVersion, 0, nullptr)); HIP_CHECK(hipGetProcAddress("hipDeviceSetSharedMemConfig", &hipDeviceSetSharedMemConfig_ptr, currentHipVersion, 0, nullptr)); @@ -201,9 +218,23 @@ TEST_CASE("Unit_hipGetProcAddress_ValidateDeviceApis") { hipError_t (*dyn_hipGetDeviceProperties_ptr)(hipDeviceProp_t*, int) = reinterpret_cast (hipGetDeviceProperties_ptr); + hipError_t(*dyn_hipGetDevicePropertiesR0000_ptr) + (hipDeviceProp_tR0000*, int) = + reinterpret_cast + (hipGetDevicePropertiesR0000_ptr); + hipError_t(*dyn_hipGetDevicePropertiesR0600_ptr) + (hipDeviceProp_tR0600*, int) = + reinterpret_cast + (hipGetDevicePropertiesR0600_ptr); hipError_t (*dyn_hipChooseDevice_ptr)(int*, hipDeviceProp_t*) = reinterpret_cast (hipChooseDevice_ptr); + hipError_t (*dyn_hipChooseDeviceR0000_ptr)(int*, hipDeviceProp_tR0000*) = + reinterpret_cast + (hipChooseDevice_ptr); + hipError_t (*dyn_hipChooseDeviceR0600_ptr)(int*, hipDeviceProp_tR0600*) = + reinterpret_cast + (hipChooseDevice_ptr); hipError_t (*dyn_hipDeviceSetSharedMemConfig_ptr)(hipSharedMemConfig) = reinterpret_cast (hipDeviceSetSharedMemConfig_ptr); @@ -336,12 +367,37 @@ TEST_CASE("Unit_hipGetProcAddress_ValidateDeviceApis") { HIP_CHECK(hipGetDeviceProperties(&prop, 0)); HIP_CHECK(dyn_hipGetDeviceProperties_ptr(&prop_ptr, 0)); REQUIRE(prop.major == prop_ptr.major); + + // hipGetDevicePropertiesR0000 API + hipDeviceProp_tR0000 propR0000, propR0000_ptr; + HIP_CHECK(hipGetDevicePropertiesR0000(&propR0000, 0)); + HIP_CHECK(dyn_hipGetDevicePropertiesR0000_ptr(&propR0000_ptr, 0)); + REQUIRE(propR0000.major == propR0000_ptr.major); + + // hipGetDevicePropertiesR0600 API + hipDeviceProp_tR0600 propR0600, propR0600_ptr; + HIP_CHECK(hipGetDevicePropertiesR0600(&propR0600, 0)); + HIP_CHECK(dyn_hipGetDevicePropertiesR0600_ptr(&propR0600_ptr, 0)); + REQUIRE(propR0600.major == propR0600_ptr.major); + // hipChooseDevice API int dev, dev_ptr; HIP_CHECK(hipChooseDevice(&dev, &prop)); HIP_CHECK(dyn_hipChooseDevice_ptr(&dev_ptr, &prop)); REQUIRE(dev == dev_ptr); + // hipChooseDeviceR0000 API + int devR0000, devR0000_ptr; + HIP_CHECK(hipChooseDeviceR0000(&devR0000, &propR0000)); + HIP_CHECK(dyn_hipChooseDeviceR0000_ptr(&devR0000_ptr, &propR0000)); + REQUIRE(devR0000 == devR0000_ptr); + + // hipChooseDeviceR0600 API + int devR0600, devR0600_ptr; + HIP_CHECK(hipChooseDeviceR0600(&devR0600, &propR0600)); + HIP_CHECK(dyn_hipChooseDeviceR0600_ptr(&devR0600_ptr, &propR0600)); + REQUIRE(devR0600 == devR0600_ptr); + // hipDeviceSetSharedMemConfig API HIP_CHECK(hipDeviceSetSharedMemConfig(hipSharedMemBankSizeFourByte)); HIP_CHECK(dyn_hipDeviceSetSharedMemConfig_ptr(hipSharedMemBankSizeFourByte)); @@ -376,11 +432,11 @@ TEST_CASE("Unit_hipGetProcAddress_ValidateDeviceApis") { */ TEST_CASE("Unit_hipGetProcAddress_PeerDeviceAccessAPIs") { - void* hipDeviceCanAccessPeer_ptr; - void* hipSetDevice_ptr; - void* hipGetDevice_ptr; - void* hipDeviceEnablePeerAccess_ptr; - void* hipDeviceDisablePeerAccess_ptr; + void* hipDeviceCanAccessPeer_ptr = nullptr; + void* hipSetDevice_ptr = nullptr; + void* hipGetDevice_ptr = nullptr; + void* hipDeviceEnablePeerAccess_ptr = nullptr; + void* hipDeviceDisablePeerAccess_ptr = nullptr; int currentHipVersion = 0; HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); @@ -467,8 +523,8 @@ bool CheckMemPoolSupport(const int device) { } TEST_CASE("Unit_hipGetProcAddress_SetGetMemPoolAPIs") { - void* hipDeviceSetMemPool_ptr; - void* hipDeviceGetMemPool_ptr; + void* hipDeviceSetMemPool_ptr = nullptr; + void* hipDeviceGetMemPool_ptr = nullptr; int currentHipVersion = 0; HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); HIP_CHECK(hipGetProcAddress("hipDeviceSetMemPool", diff --git a/projects/hip-tests/catch/unit/device/hipGetProcAddressHelpers.hh b/projects/hip-tests/catch/unit/device/hipGetProcAddressHelpers.hh new file mode 100644 index 0000000000..1eb0a56152 --- /dev/null +++ b/projects/hip-tests/catch/unit/device/hipGetProcAddressHelpers.hh @@ -0,0 +1,229 @@ +/* +Copyright (c) 2024 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. +*/ + +#ifndef CATCH_UNIT_DEVICE_HIPGETPROCADDRESSHELPERS_HH_ +#define CATCH_UNIT_DEVICE_HIPGETPROCADDRESSHELPERS_HH_ + +#include + +/** + * Local Function to fill the array with given value + */ +void fillHostArray(int *arr, int size, int value) { + for ( int i = 0; i < size; i++ ) { + arr[i] = value; + } +} + +/** + * Local Function to validate the array with given reference value + */ +bool validateHostArray(int *arr, int size, int refValue) { + for ( int i = 0; i < size; i++ ) { + if ( arr[i] != refValue ) { + return false; + } + } + return true; +} + +/** + * Local Function to fill the character array with given value + */ +void fillCharHostArray(char *arr, int size, int value) { + for ( int i = 0; i < size; i++ ) { + arr[i] = value; + } +} + +/** + * Local Function to validate the array with given reference value + */ +bool validateCharHostArray(char *arr, int size, int refValue) { + for ( int i = 0; i < size; i++ ) { + if ( arr[i] != refValue ) { + return false; + } + } + return true; +} + +/** + * Kernel to validate the array with given reference value + */ +__global__ void verifyArray(int *arr, int size, int refValue, int* status) { + for ( int i = 0; i < size; i++ ) { + if ( arr[i] != refValue ) { + *status = 0; + return; + } + } + *status = 1; + return; +} + +/** + * Local Function to validate the device array with given reference value + */ +bool validateDeviceArray(int *arr, int size, int refValue) { + int *devStatus = nullptr; + HIP_CHECK(hipMalloc(&devStatus, sizeof(int))); + REQUIRE(devStatus != nullptr); + + verifyArray<<<1, 1>>>(arr, size, refValue, devStatus); + int status; + HIP_CHECK(hipMemcpy(&status, devStatus, sizeof(int), hipMemcpyDeviceToHost)); + + HIP_CHECK(hipFree(devStatus)); + + if ( status == 1 ) { + return true; + } else { + return false; + } +} + +/** + * Kernel to validate the character array with given reference value + */ +__global__ void verifyCharArray(char *arr, int size, + int refValue, int* status) { + for ( int i = 0; i < size; i++ ) { + if ( arr[i] != refValue ) { + *status = 0; + return; + } + } + *status = 1; + return; +} + +/** + * Local Function to validate the character device array with + * given reference value + */ +bool validateCharDeviceArray(char *arr, int size, int refValue) { + int *devStatus = nullptr; + HIP_CHECK(hipMalloc(&devStatus, sizeof(int))); + REQUIRE(devStatus != nullptr); + + verifyCharArray<<< 1, 1 >>>(arr, size, refValue, devStatus); + int status; + HIP_CHECK(hipMemcpy(&status, devStatus, sizeof(int), hipMemcpyDeviceToHost)); + + HIP_CHECK(hipFree(devStatus)); + + if ( status == 1 ) { + return true; + } else { + return false; + } +} + +/** + * Kernel to fill the array with given value + */ +__global__ void fillArray(int *arr, int size, int value) { + for ( int i = 0; i < size; i++ ) { + arr[i] = value; + } +} + +/** + * Local Function to fill the device array with given value + */ +void fillDeviceArray(int *arr, int size, int value) { + fillArray<<<1, 1>>>(arr, size, value); +} + +/** + * Local Function to validate the array of different types + */ +template +bool validateArrayT(T *arr, int size, T value) { + for ( int i = 0; i < size; i++ ) { + if ( arr[i] != value ) { + return false; + } + } + return true; +} + +/** + * Kernel to increment 1 value to all the elements in array + */ +__global__ void addOneKernel(int *a, int size) { + int offset = blockDim.x * blockIdx.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = offset; i < size; i+= stride) { + a[i] += 1; + } +} + +/** + * Kernel to increment 2 value to all the elements in array + */ +__global__ void addTwoKernel(int *a, int size) { + int offset = blockDim.x * blockIdx.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = offset; i < size; i+= stride) { + a[i] += 2; + } +} + +/** + * A Simple Kernel useful for validating some kernel related scenarios + */ +__global__ void simpleKernel() { +} + +/** + * In addTenKernel, integer a value increased by 10 + */ +__global__ void addTenKernel(int *a) { + *a = *a + 10; +} + +/** + * Host function adds 10 to the integer a + */ +void addTen(void* data) { + int *a = reinterpret_cast(data); + *a = *a + 10; +} + +/** + * Host function adds 20 to the integer a + */ +void addTwenty(void* data) { + int *a = reinterpret_cast(data); + *a = *a + 20; +} + +/** + * Local function used in call back scenarios + */ +void callBackFunction(hipStream_t stream, hipError_t status, void* userData) { + REQUIRE(stream != nullptr); + REQUIRE(status == hipSuccess); + int *a = reinterpret_cast(userData); + *a += 100; +} + +#endif // CATCH_UNIT_DEVICE_HIPGETPROCADDRESSHELPERS_HH_ diff --git a/projects/hip-tests/catch/unit/device/hipGetProcAddressIpcApis.cc b/projects/hip-tests/catch/unit/device/hipGetProcAddressIpcApis.cc new file mode 100644 index 0000000000..8511c2dcbd --- /dev/null +++ b/projects/hip-tests/catch/unit/device/hipGetProcAddressIpcApis.cc @@ -0,0 +1,252 @@ +/* +Copyright (c) 2024 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. +*/ + +#include +#include +#include +#include +#include +#include +#include +#include "hipGetProcAddressHelpers.hh" + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - memory IPC related APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/device/hipGetProcAddress_IPC_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_IPC_Memory") { + int N = 40; + int Nbytes = N * sizeof(int); + + int fd[2]; + REQUIRE(pipe(fd) == 0); + + auto pid = fork(); + + // Validating hipIpcGetMemHandle API + if ( pid != 0 ) { // parent process + void* hipIpcGetMemHandle_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipIpcGetMemHandle", + &hipIpcGetMemHandle_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipIpcGetMemHandle_ptr)(hipIpcMemHandle_t *, void *) = + reinterpret_cast + (hipIpcGetMemHandle_ptr); + + int* srcHostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(srcHostMem != nullptr); + fillHostArray(srcHostMem, N, 10); + + int* devMemSrc = nullptr; + HIP_CHECK(hipMalloc(&devMemSrc, Nbytes)); + REQUIRE(devMemSrc != nullptr); + HIP_CHECK(hipMemcpy(devMemSrc, srcHostMem, Nbytes, hipMemcpyHostToDevice)); + + hipIpcMemHandle_t handle; + HIP_CHECK(dyn_hipIpcGetMemHandle_ptr(&handle, devMemSrc)); + + REQUIRE(write(fd[1], &handle, sizeof(handle)) >= 0); + REQUIRE(close(fd[1]) == 0); + + REQUIRE(wait(NULL) >= 0); + + HIP_CHECK(hipFree(devMemSrc)); + free(srcHostMem); + } else { // child process + // Validating hipIpcOpenMemHandle, hipIpcCloseMemHandle API's + void* hipIpcOpenMemHandle_ptr = nullptr; + void* hipIpcCloseMemHandle_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipIpcOpenMemHandle", + &hipIpcOpenMemHandle_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipIpcCloseMemHandle", + &hipIpcCloseMemHandle_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipIpcOpenMemHandle_ptr)(void **, hipIpcMemHandle_t, + unsigned int) = + reinterpret_cast + (hipIpcOpenMemHandle_ptr); + hipError_t (*dyn_hipIpcCloseMemHandle_ptr)(void *) = + reinterpret_cast + (hipIpcCloseMemHandle_ptr); + + hipIpcMemHandle_t handle; + REQUIRE(read(fd[0], &handle, sizeof(handle)) >= 0); + REQUIRE(close(fd[0]) == 0); + + int *devPtr = nullptr; + HIP_CHECK(dyn_hipIpcOpenMemHandle_ptr(reinterpret_cast(&devPtr), + handle, hipIpcMemLazyEnablePeerAccess)); + REQUIRE(devPtr != nullptr); + + addOneKernel<<< 1, 1 >>>(devPtr , N); + + int* dstHostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(dstHostMem != nullptr); + + HIP_CHECK(hipMemcpy(dstHostMem, devPtr, Nbytes, hipMemcpyDeviceToHost)); + REQUIRE(validateHostArray(dstHostMem, N, 11) == true); + + HIP_CHECK(dyn_hipIpcCloseMemHandle_ptr(devPtr)); + + free(dstHostMem); + } +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Event IPC related APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/device/hipGetProcAddress_IPC_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_IPC_Event") { + int fd[2]; + REQUIRE(pipe(fd) == 0); + + auto pid = fork(); + + // Validating hipIpcGetEventHandle API + if ( pid != 0 ) { // parent process + void* hipIpcGetEventHandle_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipIpcGetEventHandle", + &hipIpcGetEventHandle_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipIpcGetEventHandle_ptr)(hipIpcEventHandle_t *, + hipEvent_t) = + reinterpret_cast + (hipIpcGetEventHandle_ptr); + + hipEvent_t start = nullptr; + HIP_CHECK(hipEventCreateWithFlags(&start, hipEventInterprocess | + hipEventDisableTiming)); + REQUIRE(start != nullptr); + + hipIpcEventHandle_t handle; + HIP_CHECK(dyn_hipIpcGetEventHandle_ptr(&handle, start)); + + REQUIRE(write(fd[1], &handle, sizeof(hipIpcEventHandle_t)) >= 0); + REQUIRE(close(fd[1]) == 0); + + REQUIRE(wait(NULL) >= 0); + + HIP_CHECK(hipEventDestroy(start)); + } else { // child process + // Validating hipIpcOpenMemHandle API + void* hipIpcOpenEventHandle_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipIpcOpenEventHandle", + &hipIpcOpenEventHandle_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipIpcOpenEventHandle_ptr)(hipEvent_t *, + hipIpcEventHandle_t) = + reinterpret_cast + (hipIpcOpenEventHandle_ptr); + + hipIpcEventHandle_t handle; + REQUIRE(read(fd[0], &handle, sizeof(handle)) >= 0); + REQUIRE(close(fd[0]) == 0); + + hipEvent_t start = nullptr; + HIP_CHECK(dyn_hipIpcOpenEventHandle_ptr(&start, handle)); + REQUIRE(start != nullptr); + + hipEvent_t stop = nullptr; + HIP_CHECK(hipEventCreate(&stop)); + REQUIRE(stop != nullptr); + + int N = 40; + int Nbytes = N * sizeof(int); + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, 10); + + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipEventRecord(start, stream)); + + HIP_CHECK(hipMemcpyAsync(devMem, hostMem, Nbytes, + hipMemcpyHostToDevice, stream)); + addOneKernel<<< 1, 1 >>>(devMem , N); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, Nbytes, + hipMemcpyDeviceToHost, stream)); + + HIP_CHECK(hipEventRecord(stop, stream)); + HIP_CHECK(hipEventSynchronize(stop)); + + REQUIRE(validateHostArray(hostMem, N, 11) == true); + + float time = 0.0f; + HIP_CHECK(hipEventElapsedTime(&time, start, stop)); + REQUIRE(time > 0.0f); + + HIP_CHECK(hipEventDestroy(stop)); + HIP_CHECK(hipStreamDestroy(stream)); + free(hostMem); + HIP_CHECK(hipFree(devMem)); + } +} diff --git a/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt b/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt index 16647b8c51..5e30302b88 100644 --- a/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt @@ -11,6 +11,7 @@ if(HIP_PLATFORM MATCHES "amd") hipExtLaunchKernel.cc hipExtLaunchMultiKernelMultiDevice.cc launch_api.cc + hipGetProcAddressLaunchCbExecCtrlApis.cc ) else() # These functions are currently unimplemented on AMD diff --git a/projects/hip-tests/catch/unit/executionControl/hipGetProcAddressLaunchCbExecCtrlApis.cc b/projects/hip-tests/catch/unit/executionControl/hipGetProcAddressLaunchCbExecCtrlApis.cc new file mode 100644 index 0000000000..886dbf36d8 --- /dev/null +++ b/projects/hip-tests/catch/unit/executionControl/hipGetProcAddressLaunchCbExecCtrlApis.cc @@ -0,0 +1,546 @@ +/* +Copyright (c) 2024 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. +*/ + +#include +#include +#include +#include "../device/hipGetProcAddressHelpers.hh" + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - kernel launch APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/executionControl/hipGetProcAddress_Launch_Callback_ExecCtrl_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_KernelLaunchApis") { + void* hipConfigureCall_ptr = nullptr; + void* hipSetupArgument_ptr = nullptr; + void* hipLaunchByPtr_ptr = nullptr; + void* __hipPushCallConfiguration_ptr = nullptr; + void* __hipPopCallConfiguration_ptr = nullptr; + void* hipExtLaunchKernel_ptr = nullptr; + void* hipDrvMemcpy2DUnaligned_ptr = nullptr; + void* hipLaunchKernel_ptr = nullptr; + void* hipLaunchHostFunc_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipConfigureCall", + &hipConfigureCall_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipSetupArgument", + &hipSetupArgument_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipLaunchByPtr", + &hipLaunchByPtr_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "__hipPushCallConfiguration", + &__hipPushCallConfiguration_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "__hipPopCallConfiguration", + &__hipPopCallConfiguration_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipExtLaunchKernel", + &hipExtLaunchKernel_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipDrvMemcpy2DUnaligned", + &hipDrvMemcpy2DUnaligned_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipLaunchKernel", + &hipLaunchKernel_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipLaunchHostFunc", + &hipLaunchHostFunc_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipConfigureCall_ptr)(dim3, dim3, size_t, hipStream_t) = + reinterpret_cast + (hipConfigureCall_ptr); + hipError_t (*dyn_hipSetupArgument_ptr)(const void *, size_t, size_t) = + reinterpret_cast + (hipSetupArgument_ptr); + hipError_t (*dyn_hipLaunchByPtr_ptr)(const void *) = + reinterpret_cast + (hipLaunchByPtr_ptr); + hipError_t (*dyn__hipPushCallConfiguration_ptr)(dim3, dim3, + size_t, hipStream_t) = + reinterpret_cast + (__hipPushCallConfiguration_ptr); + hipError_t (*dyn__hipPopCallConfiguration_ptr)(dim3 *, dim3 *, + size_t *, hipStream_t *) = + reinterpret_cast + (__hipPopCallConfiguration_ptr); + hipError_t (*dyn_hipExtLaunchKernel_ptr)(const void *, dim3, dim3, void **, + size_t, hipStream_t, hipEvent_t, + hipEvent_t, int) = + reinterpret_cast + (hipExtLaunchKernel_ptr); + hipError_t (*dyn_hipDrvMemcpy2DUnaligned_ptr)(const hip_Memcpy2D *) = + reinterpret_cast + (hipDrvMemcpy2DUnaligned_ptr); + + hipError_t (*dyn_hipLaunchKernel_ptr)(const void *, dim3, dim3, void **, + size_t, hipStream_t) = + reinterpret_cast + (hipLaunchKernel_ptr); + + hipError_t (*dyn_hipLaunchHostFunc_ptr)(hipStream_t, hipHostFn_t, void *) = + reinterpret_cast + (hipLaunchHostFunc_ptr); + + // Validating hipConfigureCall, hipSetupArgument, hipLaunchByPtr API's + { + // hipConfigureCall + dim3 blocksPerGrid(1, 1, 1); + dim3 threadsPerBlock(1, 1, 1); + + HIP_CHECK(dyn_hipConfigureCall_ptr(blocksPerGrid, threadsPerBlock, + 0, nullptr)); + + // hipSetupArgument + int hostInt = 10; + int *devInt = nullptr; + HIP_CHECK(hipMalloc(&devInt, sizeof(int))); + REQUIRE(devInt != nullptr); + HIP_CHECK(hipMemcpy(devInt, &hostInt, sizeof(int), hipMemcpyHostToDevice)); + + HIP_CHECK(dyn_hipSetupArgument_ptr(&devInt, sizeof(int*), 0)); + + // hipLaunchByPtr + HIP_CHECK(dyn_hipLaunchByPtr_ptr(reinterpret_cast(addTenKernel))); + + HIP_CHECK(hipMemcpy(&hostInt, devInt, sizeof(int), hipMemcpyDeviceToHost)); + REQUIRE(hostInt == 20); + + HIP_CHECK(hipFree(devInt)); + } + + // Validating __hipPushCallConfiguration, __hipPopCallConfiguration APIs + { + // __hipPushCallConfiguration + dim3 blocksPerGrid(4, 6, 8); + dim3 threadsPerBlock(20, 30, 40); + HIP_CHECK(dyn__hipPushCallConfiguration_ptr(blocksPerGrid, threadsPerBlock, + 512, nullptr)); + + // __hipPopCallConfiguration + dim3 gridDim, blockDim; + size_t sharedMem = -1; + hipStream_t stream; + + HIP_CHECK(dyn__hipPopCallConfiguration_ptr(&gridDim, &blockDim, + &sharedMem, &stream)); + + REQUIRE(gridDim.x == blocksPerGrid.x); + REQUIRE(gridDim.y == blocksPerGrid.y); + REQUIRE(gridDim.z == blocksPerGrid.z); + REQUIRE(blockDim.x == threadsPerBlock.x); + REQUIRE(blockDim.y == threadsPerBlock.y); + REQUIRE(blockDim.z == threadsPerBlock.z); + REQUIRE(sharedMem == 512); + REQUIRE(stream == nullptr); + } + + // Validating hipExtLaunchKernel API + { + dim3 blocksPerGrid(1, 1, 1); + dim3 threadsPerBlock(1, 1, 1); + + int hostInt = 10; + int *devInt = nullptr; + HIP_CHECK(hipMalloc(&devInt, sizeof(int))); + REQUIRE(devInt != nullptr); + HIP_CHECK(hipMemcpy(devInt, &hostInt, sizeof(int), hipMemcpyHostToDevice)); + + void *kernel_args[] = {&devInt}; + HIP_CHECK(dyn_hipExtLaunchKernel_ptr(reinterpret_cast(addTenKernel), + blocksPerGrid, threadsPerBlock, kernel_args, 0, + nullptr, nullptr, nullptr, 0)); + + HIP_CHECK(hipMemcpy(&hostInt, devInt, sizeof(int), hipMemcpyDeviceToHost)); + REQUIRE(hostInt == 20); + + HIP_CHECK(hipFree(devInt)); + } + + // Validating hipDrvMemcpy2DUnaligned API + { + size_t width = 256; + size_t height = 256; + const int N = width * height; + int value = 10; + size_t pitch; + + // Host to Host + { + char* sHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(sHostMem != nullptr); + fillCharHostArray(sHostMem, N, value); + + char* dHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(dHostMem != nullptr); + + hip_Memcpy2D desc = {}; + desc.srcMemoryType = hipMemoryTypeHost; + desc.srcHost = sHostMem; + desc.srcPitch = width; + + desc.dstMemoryType = hipMemoryTypeHost; + desc.dstHost = dHostMem; + desc.dstPitch = width; + desc.WidthInBytes = width * sizeof(char); + desc.Height = height; + + HIP_CHECK(dyn_hipDrvMemcpy2DUnaligned_ptr(&desc)); + + REQUIRE(validateCharHostArray(dHostMem, N, value) == true); + + free(sHostMem); + free(dHostMem); + } + + // Host to Device + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + + hip_Memcpy2D desc = {}; + desc.srcMemoryType = hipMemoryTypeHost; + desc.srcHost = hostMem; + desc.srcPitch = width; + desc.dstMemoryType = hipMemoryTypeDevice; + desc.dstDevice = hipDeviceptr_t(devMem); + desc.dstPitch = pitch; + desc.WidthInBytes = width*sizeof(char); + desc.Height = height; + + HIP_CHECK(dyn_hipDrvMemcpy2DUnaligned_ptr(&desc)); + + REQUIRE(validateCharDeviceArray(devMem, N, value) == true); + + free(hostMem); + HIP_CHECK(hipFree(devMem)); + } + + // Device To Host + { + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + HIP_CHECK(hipMemset2D(devMem, pitch, value, width, height)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + + hip_Memcpy2D desc = {}; + desc.srcMemoryType = hipMemoryTypeDevice; + desc.srcDevice = (hipDeviceptr_t)devMem; + desc.srcPitch = width; + + desc.dstMemoryType = hipMemoryTypeHost; + desc.dstHost = hostMem; + desc.dstPitch = width; + desc.WidthInBytes = width * sizeof(char); + desc.Height = height; + + HIP_CHECK(dyn_hipDrvMemcpy2DUnaligned_ptr(&desc)); + + REQUIRE(validateCharHostArray(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + + // Device To Device + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + hip_Memcpy2D desc = {}; + desc.srcMemoryType = hipMemoryTypeDevice; + desc.srcDevice = hipDeviceptr_t(sDevMem); + desc.srcPitch = sPitch; + desc.dstMemoryType = hipMemoryTypeDevice; + desc.dstDevice = hipDeviceptr_t(dDevMem); + desc.dstPitch = dPitch; + desc.WidthInBytes = width*sizeof(char); + desc.Height = height; + + HIP_CHECK(dyn_hipDrvMemcpy2DUnaligned_ptr(&desc)); + + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + } + } + + // Validating hipLaunchKernel API + { + const int N = 10; + const int Nbytes = 10 * sizeof(int); + + int *hostArr = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostArr != nullptr); + fillHostArray(hostArr, N, 10); + + int *devArr = nullptr; + HIP_CHECK(hipMalloc(&devArr, Nbytes)); + REQUIRE(devArr != nullptr); + HIP_CHECK(hipMemcpy(devArr, hostArr, Nbytes, hipMemcpyHostToDevice)); + + dim3 blocksPerGrid(1, 1, 1); + dim3 threadsPerBlock(1, 1, N); + + struct kernelArgs { + void *arr; + int size; + }; + kernelArgs kernelArg; + kernelArg.arr = devArr; + kernelArg.size = N; + void *kernel_args[] = {&kernelArg.arr, &kernelArg.size}; + + HIP_CHECK(dyn_hipLaunchKernel_ptr(reinterpret_cast(addOneKernel), + blocksPerGrid, threadsPerBlock, kernel_args, 0, nullptr)); + + HIP_CHECK(hipMemcpy(hostArr, devArr, Nbytes, hipMemcpyDeviceToHost)); + REQUIRE(validateHostArray(hostArr, N, 11) == true); + + free(hostArr); + HIP_CHECK(hipFree(devArr)); + } + + // Validating hipLaunchHostFunc API + { + int data = 30; + hipHostFn_t fn = addTen; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipLaunchHostFunc_ptr(stream, fn, + reinterpret_cast(&data))); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(data == 40); + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - call back activity APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/executionControl/hipGetProcAddress_Launch_Callback_ExecCtrl_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_CallbackActivityAPIs") { + void* hipGetStreamDeviceId_ptr = nullptr; + void* hipApiName_ptr = nullptr; + void* hipKernelNameRef_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipGetStreamDeviceId", + &hipGetStreamDeviceId_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipApiName", + &hipApiName_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipKernelNameRef", + &hipKernelNameRef_ptr, + currentHipVersion, 0, nullptr)); + + int (*dyn_hipGetStreamDeviceId_ptr)(hipStream_t) = + reinterpret_cast + (hipGetStreamDeviceId_ptr); + const char * (*dyn_hipApiName_ptr)(uint32_t) = + reinterpret_cast + (hipApiName_ptr); + const char * (*dyn_hipKernelNameRef_ptr)(const hipFunction_t) = + reinterpret_cast + (hipKernelNameRef_ptr); + + // Validating hipGetStreamDeviceId API + { + int deviceId = 0; + HIP_CHECK(hipSetDevice(deviceId)); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + REQUIRE(hipGetStreamDeviceId(stream) == + dyn_hipGetStreamDeviceId_ptr(stream)); + + REQUIRE(dyn_hipGetStreamDeviceId_ptr(stream) == deviceId); + + HIP_CHECK(hipStreamDestroy(stream)); + } + + // Validating hipApiName API + { + uint32_t id = 1; + REQUIRE(strcmp(hipApiName(id), dyn_hipApiName_ptr(id)) == 0); + } + + // Validating hipKernelNameRef API + { + hipFunction_t funcPointer; + REQUIRE(hipGetFuncBySymbol(&funcPointer, + reinterpret_cast(addTenKernel)) == hipSuccess); + + REQUIRE(strcmp(hipKernelNameRef(funcPointer), + dyn_hipKernelNameRef_ptr(funcPointer)) == 0); + } +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Execution Control APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/executionControl/hipGetProcAddress_Launch_Callback_ExecCtrl_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_ExecutionControlAPIs") { + void* hipFuncSetAttribute_ptr = nullptr; + void* hipFuncSetCacheConfig_ptr = nullptr; + void* hipFuncSetSharedMemConfig_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipFuncSetAttribute", + &hipFuncSetAttribute_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipFuncSetCacheConfig", + &hipFuncSetCacheConfig_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipFuncSetSharedMemConfig", + &hipFuncSetSharedMemConfig_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipFuncSetAttribute_ptr)(const void *, + hipFuncAttribute, int) = + reinterpret_cast + (hipFuncSetAttribute_ptr); + + hipError_t (*dyn_hipFuncSetCacheConfig_ptr)(const void *, + hipFuncCache_t) = + reinterpret_cast + (hipFuncSetCacheConfig_ptr); + + hipError_t (*dyn_hipFuncSetSharedMemConfig_ptr)(const void *, + hipSharedMemConfig) = + reinterpret_cast + (hipFuncSetSharedMemConfig_ptr); + + // Validating hipFuncSetAttribute API + hipFuncAttribute funcAttributes[] = + { hipFuncAttributeMaxDynamicSharedMemorySize, + hipFuncAttributePreferredSharedMemoryCarveout, + hipFuncAttributeMax}; + int value = 90; + for ( auto funcAttribute : funcAttributes ) { + HIP_CHECK(dyn_hipFuncSetAttribute_ptr( + reinterpret_cast(addOneKernel), + funcAttribute, value)); + } + + // Validating hipFuncSetCacheConfig API + hipFuncCache_t funcCacheConfig[] = { hipFuncCachePreferNone, + hipFuncCachePreferShared, + hipFuncCachePreferL1, + hipFuncCachePreferEqual + }; + for ( auto config : funcCacheConfig ) { + HIP_CHECK(dyn_hipFuncSetCacheConfig_ptr( + reinterpret_cast(addOneKernel), config)); + } + + // Validating hipFuncSetSharedMemConfig API + hipSharedMemConfig sharedMemConfig[] = { hipSharedMemBankSizeDefault, + hipSharedMemBankSizeFourByte, + hipSharedMemBankSizeEightByte + }; + for ( auto config : sharedMemConfig ) { + HIP_CHECK(dyn_hipFuncSetSharedMemConfig_ptr( + reinterpret_cast(addOneKernel), config)); + } +} diff --git a/projects/hip-tests/catch/unit/graph/CMakeLists.txt b/projects/hip-tests/catch/unit/graph/CMakeLists.txt index db0ffb4bb2..c9cc8f3c38 100644 --- a/projects/hip-tests/catch/unit/graph/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/graph/CMakeLists.txt @@ -88,6 +88,7 @@ if(HIP_PLATFORM MATCHES "amd") set(AMD_SRC hipStreamCaptureExtModuleLaunchKernel.cc hipStreamBeginCaptureToGraph.cc + hipGetProcAddressGraphApis.cc ) set(TEST_SRC ${TEST_SRC} ${AMD_SRC}) endif() diff --git a/projects/hip-tests/catch/unit/graph/hipGetProcAddressGraphApis.cc b/projects/hip-tests/catch/unit/graph/hipGetProcAddressGraphApis.cc new file mode 100644 index 0000000000..25fbc6cd2d --- /dev/null +++ b/projects/hip-tests/catch/unit/graph/hipGetProcAddressGraphApis.cc @@ -0,0 +1,1544 @@ +/* +Copyright (c) 2024 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. +*/ + +#include +#include +#include +#include +#include "../device/hipGetProcAddressHelpers.hh" + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Graph (stream capture) APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/graph/hipGetProcAddress_Graph_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_GraphAPIs_StreamCapture") { + void* hipStreamBeginCapture_ptr = nullptr; + void* hipStreamIsCapturing_ptr = nullptr; + void* hipStreamEndCapture_ptr = nullptr; + void* hipStreamAddCallback_ptr = nullptr; + void* hipGraphInstantiate_ptr = nullptr; + void* hipGraphLaunch_ptr = nullptr; + void* hipGraphExecDestroy_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipStreamBeginCapture", + &hipStreamBeginCapture_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamIsCapturing", + &hipStreamIsCapturing_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamEndCapture", + &hipStreamEndCapture_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamAddCallback", + &hipStreamAddCallback_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphInstantiate", + &hipGraphInstantiate_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphLaunch", + &hipGraphLaunch_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphExecDestroy", + &hipGraphExecDestroy_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipStreamBeginCapture_ptr)( + hipStream_t, hipStreamCaptureMode) = + reinterpret_cast + (hipStreamBeginCapture_ptr); + + hipError_t (*dyn_hipStreamIsCapturing_ptr)(hipStream_t, + hipStreamCaptureStatus *) = + reinterpret_cast + (hipStreamIsCapturing_ptr); + + hipError_t (*dyn_hipStreamEndCapture_ptr)(hipStream_t, hipGraph_t *) = + reinterpret_cast + (hipStreamEndCapture_ptr); + + hipError_t (*dyn_hipStreamAddCallback_ptr)(hipStream_t, hipStreamCallback_t, + void *, unsigned int) = + reinterpret_cast + (hipStreamAddCallback_ptr); + + hipError_t (*dyn_hipGraphInstantiate_ptr)(hipGraphExec_t *, hipGraph_t, + hipGraphNode_t *, char *, size_t) = + reinterpret_cast + (hipGraphInstantiate_ptr); + + hipError_t (*dyn_hipGraphLaunch_ptr)(hipGraphExec_t, hipStream_t) = + reinterpret_cast + (hipGraphLaunch_ptr); + + hipError_t (*dyn_hipGraphExecDestroy_ptr)(hipGraphExec_t) = + reinterpret_cast + (hipGraphExecDestroy_ptr); + + int N = 40; + int Nbytes = N * sizeof(int); + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, 10); + + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + + hipGraph_t graph = nullptr; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Validating hipStreamBeginCapture API + HIP_CHECK(dyn_hipStreamBeginCapture_ptr(stream, + hipStreamCaptureModeGlobal)); + + // Validating hipStreamIsCapturing API + hipStreamCaptureStatus pCaptureStatus = hipStreamCaptureStatusNone; + HIP_CHECK(dyn_hipStreamIsCapturing_ptr(stream, &pCaptureStatus)); + REQUIRE(pCaptureStatus == hipStreamCaptureStatusActive); + + HIP_CHECK(hipMemcpyAsync(devMem, hostMem, Nbytes, + hipMemcpyHostToDevice, stream)); + addOneKernel<<< 1, 1, 0, stream >>>(devMem , N); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, Nbytes, + hipMemcpyDeviceToHost, stream)); + + // Validating hipStreamEndCapture API + HIP_CHECK(dyn_hipStreamEndCapture_ptr(stream, &graph)); + + // Validating hipStreamAddCallback API + int data = 100; + HIP_CHECK(dyn_hipStreamAddCallback_ptr(stream, callBackFunction, + reinterpret_cast(&data), 0)); + + // Validating hipGraphInstantiate API + hipGraphExec_t graphExec; + HIP_CHECK(dyn_hipGraphInstantiate_ptr(&graphExec, graph, + nullptr, nullptr, 0)); + + // Validating hipGraphLaunch API + HIP_CHECK(dyn_hipGraphLaunch_ptr(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateHostArray(hostMem, N, 11) == true); + REQUIRE(data == 200); + + // Validating hipGraphExecDestroy API + HIP_CHECK(dyn_hipGraphExecDestroy_ptr(graphExec)); + REQUIRE(dyn_hipGraphLaunch_ptr(graphExec, stream) == hipErrorInvalidValue); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(devMem)); + free(hostMem); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Graph (adding memcpy 1D and kernel nodes) APIs from the hipGetProcAddress + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/graph/hipGetProcAddress_Graph_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_GraphAPIs_AddMemcpy1DKernelNodes") { + void* hipGraphAddMemcpyNode1D_ptr = nullptr; + void* hipGraphAddKernelNode_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipGraphAddMemcpyNode1D", + &hipGraphAddMemcpyNode1D_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphAddKernelNode", + &hipGraphAddKernelNode_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipGraphAddMemcpyNode1D_ptr)(hipGraphNode_t *, hipGraph_t, + const hipGraphNode_t *, size_t, void *, const void *, size_t, + hipMemcpyKind) = + reinterpret_cast + (hipGraphAddMemcpyNode1D_ptr); + + hipError_t (*dyn_hipGraphAddKernelNode_ptr)(hipGraphNode_t *, hipGraph_t, + const hipGraphNode_t *, size_t, const hipKernelNodeParams *) = + reinterpret_cast + (hipGraphAddKernelNode_ptr); + + int N = 40; + int Nbytes = N * sizeof(int); + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, 100); + + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + + hipGraphNode_t memcpyNodeH2D, kernelNode, memcpyNodeD2H; + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Validating hipGraphAddMemcpyNode1D API + // Prepare memcpyNodeH2D + HIP_CHECK(dyn_hipGraphAddMemcpyNode1D_ptr(&memcpyNodeH2D, graph, nullptr, 0, + devMem, hostMem, Nbytes, hipMemcpyHostToDevice)); + + // Validating hipGraphAddKernelNode API + // Prepare kernelNode with memcpyNodeH2D as a dependency + ::std::vector kernelNodeDependencies; + kernelNodeDependencies.push_back(memcpyNodeH2D); + + hipKernelNodeParams kernelNodeParams{}; + kernelNodeParams.func = reinterpret_cast(addOneKernel); + kernelNodeParams.gridDim = dim3(1, 1, 1); + kernelNodeParams.blockDim = dim3(1, 1, 1); + kernelNodeParams.sharedMemBytes = 0; + + void* kernelArgs[2] = { reinterpret_cast(&devMem), + reinterpret_cast(&N) }; + kernelNodeParams.kernelParams = kernelArgs; + kernelNodeParams.extra = nullptr; + + HIP_CHECK(dyn_hipGraphAddKernelNode_ptr(&kernelNode, graph, + kernelNodeDependencies.data(), kernelNodeDependencies.size(), + &kernelNodeParams)); + + // Prepare memcpyNodeD2H with kernelNode as a dependency + ::std::vector memcpyNodeD2HDependencies; + memcpyNodeD2HDependencies.push_back(kernelNode); + HIP_CHECK(dyn_hipGraphAddMemcpyNode1D_ptr(&memcpyNodeD2H, graph, + memcpyNodeD2HDependencies.data(), memcpyNodeD2HDependencies.size(), + hostMem, devMem, Nbytes, hipMemcpyDeviceToHost)); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + #ifdef _WIN32 + HIP_CHECK(hipStreamSynchronize(0)); + #endif + + REQUIRE(validateHostArray(hostMem, N, 101) == true); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(devMem)); + free(hostMem); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Graph (adding memset and memcpy nodes) APIs from the hipGetProcAddress + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/graph/hipGetProcAddress_Graph_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_GraphAPIs_AddMemsetMemcpyNodes") { + CHECK_IMAGE_SUPPORT + + void* hipGraphAddMemsetNode_ptr = nullptr; + void* hipGraphAddMemcpyNode_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipGraphAddMemsetNode", + &hipGraphAddMemsetNode_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphAddMemcpyNode", + &hipGraphAddMemcpyNode_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipGraphAddMemsetNode_ptr)(hipGraphNode_t *, hipGraph_t, + const hipGraphNode_t *, size_t, const hipMemsetParams *) = + reinterpret_cast + (hipGraphAddMemsetNode_ptr); + + hipError_t (*dyn_hipGraphAddMemcpyNode_ptr)(hipGraphNode_t *, hipGraph_t, + const hipGraphNode_t *, size_t, const hipMemcpy3DParms *) = + reinterpret_cast + (hipGraphAddMemcpyNode_ptr); + + size_t width = 1024; + size_t height = 1024; + int N = width * height; + int value = 120; + size_t pitch; + + char *devMemSrc = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMemSrc), + &pitch, width, height)); + REQUIRE(devMemSrc != nullptr); + + char* hostMemDst = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemDst != nullptr); + + hipGraphNode_t memsetNode, memcpyNode; + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Validating hipGraphAddMemsetNode API + // Prepare memsetNode + hipMemsetParams pMemsetParams{}; + pMemsetParams.dst = reinterpret_cast(devMemSrc); + pMemsetParams.value = value; + pMemsetParams.pitch = pitch; + pMemsetParams.elementSize = sizeof(char); + pMemsetParams.width = width; + pMemsetParams.height = height; + + HIP_CHECK(dyn_hipGraphAddMemsetNode_ptr(&memsetNode, graph, + nullptr, 0, &pMemsetParams)); + + // Validating hipGraphAddMemcpyNode API + // Prepare memcpyNode with memsetNode as a dependency + ::std::vector memcpyNodeDependencies; + memcpyNodeDependencies.push_back(memsetNode); + + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = make_hipPitchedPtr(devMemSrc, pitch, width, height); + myparms.dstPtr = make_hipPitchedPtr(hostMemDst, width, width, height); + myparms.extent = make_hipExtent(width, height, 1); + myparms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(dyn_hipGraphAddMemcpyNode_ptr(&memcpyNode, graph, + memcpyNodeDependencies.data(), + memcpyNodeDependencies.size(), + &myparms)); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + #ifdef _WIN32 + HIP_CHECK(hipStreamSynchronize(0)); + #endif + + REQUIRE(validateArrayT(hostMemDst, N, value) == true); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(devMemSrc)); + free(hostMemDst); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Graph (set/get params for memset and memcpy nodes) APIs from the + * - hipGetProcAddress and then validates the basic functionality of that + * - particular APIusing the funtion pointer. + * Test source + * ------------------------ + * - unit/graph/hipGetProcAddress_Graph_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_GraphAPIs_SetGetParamsMemsetMemcpy") { + CHECK_IMAGE_SUPPORT + + void* hipGraphMemsetNodeSetParams_ptr = nullptr; + void* hipGraphMemsetNodeGetParams_ptr = nullptr; + void* hipGraphMemcpyNodeSetParams_ptr = nullptr; + void* hipGraphMemcpyNodeGetParams_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipGraphMemsetNodeSetParams", + &hipGraphMemsetNodeSetParams_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphMemsetNodeGetParams", + &hipGraphMemsetNodeGetParams_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphMemcpyNodeSetParams", + &hipGraphMemcpyNodeSetParams_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphMemcpyNodeGetParams", + &hipGraphMemcpyNodeGetParams_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipGraphMemsetNodeSetParams_ptr)( + hipGraphNode_t, const hipMemsetParams *) = + reinterpret_cast + (hipGraphMemsetNodeSetParams_ptr); + + hipError_t (*dyn_hipGraphMemsetNodeGetParams_ptr)( + hipGraphNode_t, hipMemsetParams *) = + reinterpret_cast + (hipGraphMemsetNodeGetParams_ptr); + + hipError_t (*dyn_hipGraphMemcpyNodeSetParams_ptr)( + hipGraphNode_t, const hipMemcpy3DParms *) = + reinterpret_cast + (hipGraphMemcpyNodeSetParams_ptr); + + hipError_t (*dyn_hipGraphMemcpyNodeGetParams_ptr)( + hipGraphNode_t, hipMemcpy3DParms *) = + reinterpret_cast + (hipGraphMemcpyNodeGetParams_ptr); + + size_t width = 1024; + size_t height = 1024; + int N = width * height; + int value = 120; + size_t pitch; + + char *devMemSrc1 = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMemSrc1), + &pitch, width, height)); + REQUIRE(devMemSrc1 != nullptr); + + char* hostMemDst1 = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemDst1 != nullptr); + fillCharHostArray(hostMemDst1, N, 100); + + char *devMemSrc2 = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMemSrc2), + &pitch, width, height)); + REQUIRE(devMemSrc2 != nullptr); + + char* hostMemDst2 = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemDst2 != nullptr); + fillCharHostArray(hostMemDst2, N, 100); + + hipGraphNode_t memsetNode, memcpyNode; + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Prepare memset node + hipMemsetParams initialMemsetParams{}; + initialMemsetParams.dst = reinterpret_cast(devMemSrc1); + initialMemsetParams.value = value; + initialMemsetParams.pitch = pitch; + initialMemsetParams.elementSize = sizeof(char); + initialMemsetParams.width = width; + initialMemsetParams.height = height; + + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, + nullptr, 0, &initialMemsetParams)); + + hipMemsetParams receivedMemsetValues{}; + HIP_CHECK(dyn_hipGraphMemsetNodeGetParams_ptr(memsetNode, + &receivedMemsetValues)); + + REQUIRE(receivedMemsetValues.dst == devMemSrc1); + REQUIRE(receivedMemsetValues.value == value); + REQUIRE(receivedMemsetValues.pitch == pitch); + REQUIRE(receivedMemsetValues.elementSize == sizeof(char)); + REQUIRE(receivedMemsetValues.width == width); + REQUIRE(receivedMemsetValues.height == height); + + hipMemsetParams correctedMemsetParams{}; + correctedMemsetParams.dst = reinterpret_cast(devMemSrc2); + correctedMemsetParams.value = value; + correctedMemsetParams.pitch = pitch; + correctedMemsetParams.elementSize = sizeof(char); + correctedMemsetParams.width = width; + correctedMemsetParams.height = height; + + // Validating hipGraphMemsetNodeSetParams API + HIP_CHECK(dyn_hipGraphMemsetNodeSetParams_ptr(memsetNode, + &correctedMemsetParams)); + + // Validating hipGraphMemsetNodeGetParams API + HIP_CHECK(dyn_hipGraphMemsetNodeGetParams_ptr(memsetNode, + &receivedMemsetValues)); + + REQUIRE(receivedMemsetValues.dst == devMemSrc2); + REQUIRE(receivedMemsetValues.value == value); + REQUIRE(receivedMemsetValues.pitch == pitch); + REQUIRE(receivedMemsetValues.elementSize == sizeof(char)); + REQUIRE(receivedMemsetValues.width == width); + REQUIRE(receivedMemsetValues.height == height); + + // Prepare memcpyNode + ::std::vector memcpyNodeDependencies; + memcpyNodeDependencies.push_back(memsetNode); + + hipMemcpy3DParms initialParms{}; + initialParms.srcPos = make_hipPos(0, 0, 0); + initialParms.dstPos = make_hipPos(0, 0, 0); + initialParms.srcPtr = make_hipPitchedPtr(devMemSrc1, pitch, + width, height); + initialParms.dstPtr = make_hipPitchedPtr(hostMemDst1, width, + width, height); + initialParms.extent = make_hipExtent(width, height, 1); + initialParms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, + memcpyNodeDependencies.data(), + memcpyNodeDependencies.size(), &initialParms)); + + hipMemcpy3DParms receivedMemcpyValues{}; + HIP_CHECK(dyn_hipGraphMemcpyNodeGetParams_ptr(memcpyNode, + &receivedMemcpyValues)); + + REQUIRE(receivedMemcpyValues.srcArray == nullptr); + REQUIRE(receivedMemcpyValues.srcPos.x == 0); + REQUIRE(receivedMemcpyValues.srcPos.y == 0); + REQUIRE(receivedMemcpyValues.srcPos.z == 0); + REQUIRE(receivedMemcpyValues.srcPtr.ptr == devMemSrc1); + REQUIRE(receivedMemcpyValues.srcPtr.pitch == pitch); + REQUIRE(receivedMemcpyValues.srcPtr.xsize == width); + REQUIRE(receivedMemcpyValues.srcPtr.ysize == height); + REQUIRE(receivedMemcpyValues.dstArray == nullptr); + REQUIRE(receivedMemcpyValues.dstPos.x == 0); + REQUIRE(receivedMemcpyValues.dstPos.y == 0); + REQUIRE(receivedMemcpyValues.dstPos.z == 0); + REQUIRE(receivedMemcpyValues.dstPtr.ptr == hostMemDst1); + REQUIRE(receivedMemcpyValues.dstPtr.pitch == pitch); + REQUIRE(receivedMemcpyValues.dstPtr.xsize == width); + REQUIRE(receivedMemcpyValues.dstPtr.ysize == height); + REQUIRE(receivedMemcpyValues.extent.width == width); + REQUIRE(receivedMemcpyValues.extent.height == height); + REQUIRE(receivedMemcpyValues.extent.depth == 1); + REQUIRE(receivedMemcpyValues.kind == hipMemcpyDeviceToHost); + + hipMemcpy3DParms correctedParms{}; + correctedParms.srcPos = make_hipPos(0, 0, 0); + correctedParms.dstPos = make_hipPos(0, 0, 0); + correctedParms.srcPtr = make_hipPitchedPtr(devMemSrc2, pitch, width, height); + correctedParms.dstPtr = make_hipPitchedPtr(hostMemDst2, width, + width, height); + correctedParms.extent = make_hipExtent(width, height, 1); + correctedParms.kind = hipMemcpyDeviceToHost; + + // Validating hipGraphMemcpyNodeSetParams API + HIP_CHECK(dyn_hipGraphMemcpyNodeSetParams_ptr(memcpyNode, &correctedParms)); + + // Validating hipGraphMemcpyNodeGetParams API + HIP_CHECK(dyn_hipGraphMemcpyNodeGetParams_ptr(memcpyNode, + &receivedMemcpyValues)); + + REQUIRE(receivedMemcpyValues.srcArray == nullptr); + REQUIRE(receivedMemcpyValues.srcPos.x == 0); + REQUIRE(receivedMemcpyValues.srcPos.y == 0); + REQUIRE(receivedMemcpyValues.srcPos.z == 0); + REQUIRE(receivedMemcpyValues.srcPtr.ptr == devMemSrc2); + REQUIRE(receivedMemcpyValues.srcPtr.pitch == pitch); + REQUIRE(receivedMemcpyValues.srcPtr.xsize == width); + REQUIRE(receivedMemcpyValues.srcPtr.ysize == height); + REQUIRE(receivedMemcpyValues.dstArray == nullptr); + REQUIRE(receivedMemcpyValues.dstPos.x == 0); + REQUIRE(receivedMemcpyValues.dstPos.y == 0); + REQUIRE(receivedMemcpyValues.dstPos.z == 0); + REQUIRE(receivedMemcpyValues.dstPtr.ptr == hostMemDst2); + REQUIRE(receivedMemcpyValues.dstPtr.pitch == pitch); + REQUIRE(receivedMemcpyValues.dstPtr.xsize == width); + REQUIRE(receivedMemcpyValues.dstPtr.ysize == height); + REQUIRE(receivedMemcpyValues.extent.width == width); + REQUIRE(receivedMemcpyValues.extent.height == height); + REQUIRE(receivedMemcpyValues.extent.depth == 1); + REQUIRE(receivedMemcpyValues.kind == hipMemcpyDeviceToHost); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + #ifdef _WIN32 + HIP_CHECK(hipStreamSynchronize(0)); + #endif + + REQUIRE(validateArrayT(hostMemDst1, N, 100) == true); + REQUIRE(validateArrayT(hostMemDst2, N, 120) == true); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(devMemSrc1)); + HIP_CHECK(hipFree(devMemSrc2)); + free(hostMemDst1); + free(hostMemDst2); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Graph (set/get params for kernel) APIs from the + * - hipGetProcAddress and then validates the basic functionality of that + * - particular APIusing the funtion pointer. + * Test source + * ------------------------ + * - unit/graph/hipGetProcAddress_Graph_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_GraphAPIs_KernelNodeSetGetParams") { + void* hipGraphKernelNodeGetParams_ptr = nullptr; + void* hipGraphKernelNodeSetParams_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipGraphKernelNodeGetParams", + &hipGraphKernelNodeGetParams_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphKernelNodeSetParams", + &hipGraphKernelNodeSetParams_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipGraphKernelNodeGetParams_ptr)( + hipGraphNode_t, hipKernelNodeParams *) = + reinterpret_cast + (hipGraphKernelNodeGetParams_ptr); + + hipError_t (*dyn_hipGraphKernelNodeSetParams_ptr)( + hipGraphNode_t, const hipKernelNodeParams *) = + reinterpret_cast(hipGraphKernelNodeSetParams_ptr); + + int N = 40; + int Nbytes = N * sizeof(int); + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, 100); + + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + + hipGraphNode_t memcpyNodeH2D, kernelNode, memcpyNodeD2H; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Prepare memcpyNodeH2D + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNodeH2D, graph, nullptr, 0, + devMem, hostMem, Nbytes, hipMemcpyHostToDevice)); + + // Prepare kernelNode with memcpyNodeH2D as a dependency + ::std::vector kernelNodeDependencies; + kernelNodeDependencies.push_back(memcpyNodeH2D); + + void* kernelArgs[2] = { reinterpret_cast(&devMem), + reinterpret_cast(&N) }; + + hipKernelNodeParams kernelNodeParams{}; + kernelNodeParams.func = reinterpret_cast(addOneKernel); + kernelNodeParams.gridDim = dim3(1, 1, 1); + kernelNodeParams.blockDim = dim3(1, 1, 1); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = kernelArgs; + kernelNodeParams.extra = nullptr; + + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, + kernelNodeDependencies.data(), + kernelNodeDependencies.size(), &kernelNodeParams)); + + // Prepare memcpyNodeD2H with kernelNode as a dependency + ::std::vector memcpyNodeD2HDependencies; + memcpyNodeD2HDependencies.push_back(kernelNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNodeD2H, graph, + memcpyNodeD2HDependencies.data(), memcpyNodeD2HDependencies.size(), + hostMem, devMem, Nbytes, hipMemcpyDeviceToHost)); + + // Get and set Kernel node params + hipKernelNodeParams receivedKernelNodeParams; + + // Validating hipGraphKernelNodeGetParams API + HIP_CHECK(dyn_hipGraphKernelNodeGetParams_ptr(kernelNode, + &receivedKernelNodeParams)); + + REQUIRE(receivedKernelNodeParams.func == addOneKernel); + REQUIRE(receivedKernelNodeParams.gridDim.x == 1); + REQUIRE(receivedKernelNodeParams.gridDim.y == 1); + REQUIRE(receivedKernelNodeParams.gridDim.z == 1); + REQUIRE(receivedKernelNodeParams.blockDim.x == 1); + REQUIRE(receivedKernelNodeParams.blockDim.y == 1); + REQUIRE(receivedKernelNodeParams.blockDim.z == 1); + REQUIRE(*(reinterpret_cast(receivedKernelNodeParams.kernelParams[0])) + == *(reinterpret_cast(kernelArgs[0]))); + REQUIRE(*(reinterpret_cast(receivedKernelNodeParams.kernelParams[1])) + == *(reinterpret_cast(kernelArgs[1]))); + REQUIRE(receivedKernelNodeParams.extra == nullptr); + REQUIRE(receivedKernelNodeParams.sharedMemBytes == 0); + + hipKernelNodeParams correctedKernelNodeParams{}; + correctedKernelNodeParams.func = reinterpret_cast(addTwoKernel); + correctedKernelNodeParams.gridDim = dim3(2, 1, 1); + correctedKernelNodeParams.blockDim = dim3(2, 1, 1); + correctedKernelNodeParams.sharedMemBytes = 0; + correctedKernelNodeParams.kernelParams = kernelArgs; + correctedKernelNodeParams.extra = nullptr; + + // Validating hipGraphKernelNodeSetParams API + HIP_CHECK(dyn_hipGraphKernelNodeSetParams_ptr(kernelNode, + &correctedKernelNodeParams)); + + HIP_CHECK(dyn_hipGraphKernelNodeGetParams_ptr(kernelNode, + &receivedKernelNodeParams)); + + REQUIRE(receivedKernelNodeParams.func == addTwoKernel); + REQUIRE(receivedKernelNodeParams.gridDim.x == 2); + REQUIRE(receivedKernelNodeParams.gridDim.y == 1); + REQUIRE(receivedKernelNodeParams.gridDim.z == 1); + REQUIRE(receivedKernelNodeParams.blockDim.x == 2); + REQUIRE(receivedKernelNodeParams.blockDim.y == 1); + REQUIRE(receivedKernelNodeParams.blockDim.z == 1); + REQUIRE(*(reinterpret_cast(receivedKernelNodeParams.kernelParams[0])) + == *(reinterpret_cast(kernelArgs[0]))); + REQUIRE(*(reinterpret_cast(receivedKernelNodeParams.kernelParams[1])) + == *(reinterpret_cast(kernelArgs[1]))); + REQUIRE(receivedKernelNodeParams.extra == nullptr); + REQUIRE(receivedKernelNodeParams.sharedMemBytes == 0); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateHostArray(hostMem, N, 102) == true); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)) + HIP_CHECK(hipFree(devMem)); + free(hostMem); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Graph (set/get attributes for kernel) APIs from the + * - hipGetProcAddress and then validates the basic functionality of that + * - particular APIusing the funtion pointer. + * Test source + * ------------------------ + * - unit/graph/hipGetProcAddress_Graph_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_GraphAPIs_KernelNodeSetGetAttribute") { + void* hipGraphKernelNodeSetAttribute_ptr = nullptr; + void* hipGraphKernelNodeGetAttribute_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipGraphKernelNodeSetAttribute", + &hipGraphKernelNodeSetAttribute_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphKernelNodeGetAttribute", + &hipGraphKernelNodeGetAttribute_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipGraphKernelNodeSetAttribute_ptr)( + hipGraphNode_t, hipKernelNodeAttrID, const hipKernelNodeAttrValue *) = + reinterpret_cast(hipGraphKernelNodeSetAttribute_ptr); + + hipError_t (*dyn_hipGraphKernelNodeGetAttribute_ptr)(hipGraphNode_t, + hipKernelNodeAttrID, hipKernelNodeAttrValue *) = + reinterpret_cast(hipGraphKernelNodeGetAttribute_ptr); + + hipGraphNode_t kernelNode; + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipKernelNodeParams kernelNodeParams{}; + kernelNodeParams.func = reinterpret_cast(simpleKernel); + kernelNodeParams.gridDim = dim3(1, 1, 1); + kernelNodeParams.blockDim = dim3(1, 1, 1); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = nullptr; + kernelNodeParams.extra = nullptr; + + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, + nullptr, 0, &kernelNodeParams)); + + hipKernelNodeAttrValue attributeToSet, attributeToGet; + attributeToSet.cooperative = 1; + + // Validating hipGraphKernelNodeSetAttribute API + HIP_CHECK(dyn_hipGraphKernelNodeSetAttribute_ptr(kernelNode, + hipKernelNodeAttributeCooperative, &attributeToSet)); + + // Validating hipGraphKernelNodeGetAttribute API + HIP_CHECK(dyn_hipGraphKernelNodeGetAttribute_ptr(kernelNode, + hipKernelNodeAttributeCooperative, &attributeToGet)); + + REQUIRE(attributeToGet.cooperative == 1); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + #ifdef _WIN32 + HIP_CHECK(hipStreamSynchronize(0)); + #endif + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Graph (host) APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that + * - particular APIusing the funtion pointer. + * Test source + * ------------------------ + * - unit/graph/hipGetProcAddress_Graph_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_GraphAPIs_HostNode") { + void* hipGraphAddHostNode_ptr = nullptr; + void* hipGraphHostNodeGetParams_ptr = nullptr; + void* hipGraphHostNodeSetParams_ptr = nullptr; + void* hipGraphExecHostNodeSetParams_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipGraphAddHostNode", + &hipGraphAddHostNode_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphHostNodeGetParams", + &hipGraphHostNodeGetParams_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphHostNodeSetParams", + &hipGraphHostNodeSetParams_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphExecHostNodeSetParams", + &hipGraphExecHostNodeSetParams_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipGraphAddHostNode_ptr)(hipGraphNode_t *, hipGraph_t, + const hipGraphNode_t *, size_t, const hipHostNodeParams *) = + reinterpret_cast + (hipGraphAddHostNode_ptr); + hipError_t (*dyn_hipGraphHostNodeGetParams_ptr)(hipGraphNode_t, + hipHostNodeParams *) = + reinterpret_cast + (hipGraphHostNodeGetParams_ptr); + hipError_t (*dyn_hipGraphHostNodeSetParams_ptr)(hipGraphNode_t, + const hipHostNodeParams *) = + reinterpret_cast + (hipGraphHostNodeSetParams_ptr); + hipError_t (*dyn_hipGraphExecHostNodeSetParams_ptr)(hipGraphExec_t, + hipGraphNode_t, const hipHostNodeParams *) = + reinterpret_cast + (hipGraphExecHostNodeSetParams_ptr); + + // Validating hipGraphAddHostNode API + { + int hostInt = 10; + hipGraphNode_t hostNode; + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Prepare hostNode + hipHostNodeParams hostNodeParams; + hostNodeParams.fn = addTen; + hostNodeParams.userData = &hostInt; + + HIP_CHECK(dyn_hipGraphAddHostNode_ptr(&hostNode, graph, + nullptr, 0, &hostNodeParams)); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); + + REQUIRE(hostInt == 20); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + } + + // Validating hipGraphHostNodeGetParams, hipGraphHostNodeSetParams API's + { + int hostInt = 10; + int hostIntNew = 20; + + hipGraphNode_t hostNode; + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Prepare hostNode + hipHostNodeParams hostNodeParams; + hostNodeParams.fn = addTen; + hostNodeParams.userData = &hostInt; + + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, + nullptr, 0, &hostNodeParams)); + + hipHostNodeParams receivedHostNodeParams; + HIP_CHECK(dyn_hipGraphHostNodeGetParams_ptr(hostNode, + &receivedHostNodeParams)); + REQUIRE(receivedHostNodeParams.fn == addTen); + REQUIRE(*(reinterpret_cast(receivedHostNodeParams.userData)) == 10); + + hipHostNodeParams hostNodeParamsNew; + hostNodeParamsNew.fn = addTwenty; + hostNodeParamsNew.userData = &hostIntNew; + HIP_CHECK(dyn_hipGraphHostNodeSetParams_ptr(hostNode, &hostNodeParamsNew)); + + HIP_CHECK(dyn_hipGraphHostNodeGetParams_ptr(hostNode, + &receivedHostNodeParams)); + REQUIRE(receivedHostNodeParams.fn == addTwenty); + REQUIRE(*(reinterpret_cast(receivedHostNodeParams.userData)) == 20); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); + + REQUIRE(hostInt == 10); + REQUIRE(hostIntNew == 40); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + } + + // Validating hipGraphExecHostNodeSetParams API + { + int hostInt = 10; + int hostIntNew = 20; + + hipGraphNode_t hostNode; + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Prepare hostNode + hipHostNodeParams hostNodeParams; + hostNodeParams.fn = addTen; + hostNodeParams.userData = &hostInt; + + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, + nullptr, 0, &hostNodeParams)); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + // Update hostNode params + hipHostNodeParams hostNodeParamsNew; + hostNodeParamsNew.fn = addTwenty; + hostNodeParamsNew.userData = &hostIntNew; + HIP_CHECK(dyn_hipGraphExecHostNodeSetParams_ptr(graphExec, hostNode, + &hostNodeParamsNew)); + + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); + + REQUIRE(hostInt == 10); + REQUIRE(hostIntNew == 40); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + } +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of hipGraphExecUpdate API + * - from the hipGetProcAddress API + * - and then validates the basic functionality of that + * - particular APIusing the funtion pointer. + * Test source + * ------------------------ + * - unit/graph/hipGetProcAddress_Graph_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_GraphAPIs_ExecUpdate") { + void* hipGraphExecUpdate_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipGraphExecUpdate", + &hipGraphExecUpdate_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipGraphExecUpdate_ptr)(hipGraphExec_t, hipGraph_t, + hipGraphNode_t *, hipGraphExecUpdateResult *) = + reinterpret_cast + (hipGraphExecUpdate_ptr); + + int hostInt_1 = 10; + int hostInt_2 = 20; + + // Prepare graph1 with hostNode1 + hipGraph_t graph1 = nullptr; + HIP_CHECK(hipGraphCreate(&graph1, 0)); + hipGraphNode_t hostNode1; + hipHostNodeParams hostNodeParams1; + hostNodeParams1.fn = addTen; + hostNodeParams1.userData = &hostInt_1; + HIP_CHECK(hipGraphAddHostNode(&hostNode1, graph1, + nullptr, 0, &hostNodeParams1)); + + // Prepare graphExec with graph1 + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph1, nullptr, nullptr, 0)); + + // Prepare graph2 with hostNode2 + hipGraph_t graph2 = nullptr; + HIP_CHECK(hipGraphCreate(&graph2, 0)); + hipGraphNode_t hostNode2; + hipHostNodeParams hostNodeParams2; + hostNodeParams2.fn = addTwenty; + hostNodeParams2.userData = &hostInt_2; + HIP_CHECK(hipGraphAddHostNode(&hostNode2, graph2, + nullptr, 0, &hostNodeParams2)); + + // Update graphExec with graph2 + hipGraphNode_t hErrorNode_out = nullptr; + hipGraphExecUpdateResult updateResult_out; + + // Validating hipGraphExecUpdate API + HIP_CHECK(dyn_hipGraphExecUpdate_ptr(graphExec, graph2, + &hErrorNode_out, &updateResult_out)); + + REQUIRE(hErrorNode_out == nullptr); + REQUIRE(updateResult_out == hipGraphExecUpdateSuccess); + + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); + + REQUIRE(hostInt_1 == 10); + REQUIRE(hostInt_2 == 40); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Graph (for mem cpoy node set params) APIs from the + * - hipGetProcAddress and then validates the basic functionality of that + * - particular APIusing the funtion pointer. + * Test source + * ------------------------ + * - unit/graph/hipGetProcAddress_Graph_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_GraphAPIs_Memcpy1DSetParams") { + void* hipGraphMemcpyNodeSetParams1D_ptr = nullptr; + void* hipGraphExecMemcpyNodeSetParams1D_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipGraphMemcpyNodeSetParams1D", + &hipGraphMemcpyNodeSetParams1D_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphExecMemcpyNodeSetParams1D", + &hipGraphExecMemcpyNodeSetParams1D_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipGraphMemcpyNodeSetParams1D_ptr)( + hipGraphNode_t, void *, const void *, size_t, hipMemcpyKind) = + reinterpret_cast + (hipGraphMemcpyNodeSetParams1D_ptr); + + hipError_t (*dyn_hipGraphExecMemcpyNodeSetParams1D_ptr)( + hipGraphExec_t, hipGraphNode_t, void *, const void *, + size_t, hipMemcpyKind) = + reinterpret_cast + (hipGraphExecMemcpyNodeSetParams1D_ptr); + + int N = 40; + int Nbytes = N * sizeof(int); + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, 20); + + // Validating hipGraphMemcpyNodeSetParams1D API + { + int* devMem_1 = nullptr; + HIP_CHECK(hipMalloc(&devMem_1, Nbytes)); + REQUIRE(devMem_1 != nullptr); + fillDeviceArray(devMem_1, N, 10); + + int* devMem_2 = nullptr; + HIP_CHECK(hipMalloc(&devMem_2, Nbytes)); + REQUIRE(devMem_2 != nullptr); + fillDeviceArray(devMem_2, N, 10); + + hipGraphNode_t memcpyNodeH2D; + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNodeH2D, graph, nullptr, 0, + devMem_1, hostMem, Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(dyn_hipGraphMemcpyNodeSetParams1D_ptr(memcpyNodeH2D, devMem_2, + hostMem, Nbytes, hipMemcpyHostToDevice)); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + #ifdef _WIN32 + HIP_CHECK(hipStreamSynchronize(0)); + #endif + + REQUIRE(validateDeviceArray(devMem_1, N, 10) == true); + REQUIRE(validateDeviceArray(devMem_2, N, 20) == true); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(devMem_1)); + HIP_CHECK(hipFree(devMem_2)); + } + + // Validating hipGraphExecMemcpyNodeSetParams1D API + { + int* devMem_1 = nullptr; + HIP_CHECK(hipMalloc(&devMem_1, Nbytes)); + REQUIRE(devMem_1 != nullptr); + fillDeviceArray(devMem_1, N, 10); + + int* devMem_2 = nullptr; + HIP_CHECK(hipMalloc(&devMem_2, Nbytes)); + REQUIRE(devMem_2 != nullptr); + fillDeviceArray(devMem_2, N, 10); + + hipGraphNode_t memcpyNodeH2D; + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNodeH2D, graph, nullptr, 0, + devMem_1, hostMem, Nbytes, hipMemcpyHostToDevice)); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(dyn_hipGraphExecMemcpyNodeSetParams1D_ptr(graphExec, + memcpyNodeH2D, devMem_2, hostMem, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + #ifdef _WIN32 + HIP_CHECK(hipStreamSynchronize(0)); + #endif + + REQUIRE(validateDeviceArray(devMem_1, N, 10) == true); + REQUIRE(validateDeviceArray(devMem_2, N, 20) == true); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(devMem_1)); + HIP_CHECK(hipFree(devMem_2)); + } + free(hostMem); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Graph (memory allocation and free) APIs from the + * - hipGetProcAddress and then validates the basic functionality of that + * - particular APIusing the funtion pointer. + * Test source + * ------------------------ + * - unit/graph/hipGetProcAddress_Graph_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_GraphAPIs_MemAllocAndFree") { + void* hipGraphAddMemAllocNode_ptr = nullptr; + void* hipGraphAddMemFreeNode_ptr = nullptr; + void* hipGraphMemAllocNodeGetParams_ptr = nullptr; + void* hipGraphMemFreeNodeGetParams_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipGraphAddMemAllocNode", + &hipGraphAddMemAllocNode_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphAddMemFreeNode", + &hipGraphAddMemFreeNode_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphMemAllocNodeGetParams", + &hipGraphMemAllocNodeGetParams_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphMemFreeNodeGetParams", + &hipGraphMemFreeNodeGetParams_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipGraphAddMemAllocNode_ptr)(hipGraphNode_t *, hipGraph_t, + const hipGraphNode_t *, size_t, hipMemAllocNodeParams *) = + reinterpret_cast + (hipGraphAddMemAllocNode_ptr); + + hipError_t (*dyn_hipGraphAddMemFreeNode_ptr)(hipGraphNode_t *, hipGraph_t, + const hipGraphNode_t *, size_t, void *) = + reinterpret_cast + (hipGraphAddMemFreeNode_ptr); + + hipError_t (*dyn_hipGraphMemAllocNodeGetParams_ptr)(hipGraphNode_t, + hipMemAllocNodeParams *) = + reinterpret_cast + (hipGraphMemAllocNodeGetParams_ptr); + + hipError_t (*dyn_hipGraphMemFreeNodeGetParams_ptr)(hipGraphNode_t, void *) = + reinterpret_cast + (hipGraphMemFreeNodeGetParams_ptr); + + int N = 300; + int Nbytes = N * sizeof(int); + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, 10); + + int *devMem = nullptr; + + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipGraphNode_t memAllocNode, kernelNode, memcpyNodeD2H, memFreeNode; + + hipMemAllocNodeParams memAllocNodeParams{}; + memAllocNodeParams.poolProps.allocType = hipMemAllocationTypePinned; + memAllocNodeParams.poolProps.handleTypes = hipMemHandleTypeNone; + memAllocNodeParams.poolProps.location.type = hipMemLocationTypeDevice; + memAllocNodeParams.poolProps.location.id = 0; + memAllocNodeParams.poolProps.win32SecurityAttributes = nullptr; + memAllocNodeParams.poolProps.maxSize = 1024; + hipMemAccessDesc accessDescs; + accessDescs.location.id = 0; + accessDescs.location.type = hipMemLocationTypeDevice; + accessDescs.flags = hipMemAccessFlagsProtReadWrite; + memAllocNodeParams.accessDescs = &accessDescs; + memAllocNodeParams.accessDescCount = 1; + memAllocNodeParams.bytesize = Nbytes; + + // Validating hipGraphAddMemAllocNode API + HIP_CHECK(dyn_hipGraphAddMemAllocNode_ptr(&memAllocNode, graph, + nullptr, 0, &memAllocNodeParams)); + devMem = reinterpret_cast(memAllocNodeParams.dptr); + + ::std::vector kernelNodeDependencies; + kernelNodeDependencies.push_back(memAllocNode); + + hipKernelNodeParams kernelNodeParams{}; + kernelNodeParams.func = reinterpret_cast(fillArray); + kernelNodeParams.gridDim = dim3(1, 1, 1); + kernelNodeParams.blockDim = dim3(1, 1, 1); + kernelNodeParams.sharedMemBytes = 0; + int value = 20; + + void* kernelArgs[3] = { reinterpret_cast(&devMem), + reinterpret_cast(&N), + reinterpret_cast(&value) }; + kernelNodeParams.kernelParams = kernelArgs; + kernelNodeParams.extra = nullptr; + + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, + kernelNodeDependencies.data(), kernelNodeDependencies.size(), + &kernelNodeParams)); + + ::std::vector memcpyNodeD2HDependencies; + memcpyNodeD2HDependencies.push_back(kernelNode); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNodeD2H, graph, + memcpyNodeD2HDependencies.data(), memcpyNodeD2HDependencies.size(), + hostMem, devMem, Nbytes, hipMemcpyDeviceToHost)); + + ::std::vector memFreeNodeDependencies; + memFreeNodeDependencies.push_back(memcpyNodeD2H); + + // Validating hipGraphAddMemFreeNode API + HIP_CHECK(dyn_hipGraphAddMemFreeNode_ptr(&memFreeNode, graph, + memFreeNodeDependencies.data(), memFreeNodeDependencies.size(), + reinterpret_cast(devMem))); + + // Validating hipGraphMemAllocNodeGetParams API + hipMemAllocNodeParams recvdParams{}; + HIP_CHECK(dyn_hipGraphMemAllocNodeGetParams_ptr(memAllocNode, + &recvdParams)); + + REQUIRE(recvdParams.poolProps.allocType == hipMemAllocationTypePinned); + REQUIRE(recvdParams.poolProps.handleTypes == hipMemHandleTypeNone); + REQUIRE(recvdParams.poolProps.location.type == hipMemLocationTypeDevice); + REQUIRE(recvdParams.poolProps.location.id == 0); + REQUIRE(recvdParams.poolProps.win32SecurityAttributes == nullptr); + REQUIRE(recvdParams.poolProps.maxSize == 1024); + REQUIRE(recvdParams.accessDescs->location.id == 0); + REQUIRE(recvdParams.accessDescs->location.type == hipMemLocationTypeDevice); + REQUIRE(recvdParams.accessDescs->flags == hipMemAccessFlagsProtReadWrite); + REQUIRE(recvdParams.accessDescCount == 1); + REQUIRE(recvdParams.bytesize == Nbytes); + + // Validating hipGraphMemFreeNodeGetParams API + size_t dev_ptr = 0; + HIP_CHECK(dyn_hipGraphMemFreeNodeGetParams_ptr(memFreeNode, + reinterpret_cast(&dev_ptr))); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + #ifdef _WIN32 + HIP_CHECK(hipStreamSynchronize(0)); + #endif + + REQUIRE(validateHostArray(hostMem, N, 20) == true); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + free(hostMem); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Graph (for memset node set params) APIs from the + * - hipGetProcAddress and then validates the basic functionality of that + * - particular APIusing the funtion pointer. + * Test source + * ------------------------ + * - unit/graph/hipGetProcAddress_Graph_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_GraphAPIs_ExecMemsetMemcpySetParams") { + CHECK_IMAGE_SUPPORT + + void* hipGraphExecMemsetNodeSetParams_ptr = nullptr; + void* hipGraphExecMemcpyNodeSetParams_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipGraphExecMemsetNodeSetParams", + &hipGraphExecMemsetNodeSetParams_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphExecMemcpyNodeSetParams", + &hipGraphExecMemcpyNodeSetParams_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipGraphExecMemsetNodeSetParams_ptr)(hipGraphExec_t, + hipGraphNode_t, const hipMemsetParams *) = + reinterpret_cast(hipGraphExecMemsetNodeSetParams_ptr); + + hipError_t (*dyn_hipGraphExecMemcpyNodeSetParams_ptr)(hipGraphExec_t, + hipGraphNode_t, hipMemcpy3DParms *) = + reinterpret_cast(hipGraphExecMemcpyNodeSetParams_ptr); + + size_t width = 1024; + size_t height = 1024; + int N = width * height; + int value = 120; + size_t pitch; + + char *devMemSrc1 = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMemSrc1), + &pitch, width, height)); + REQUIRE(devMemSrc1 != nullptr); + + char* hostMemDst1 = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemDst1 != nullptr); + fillCharHostArray(hostMemDst1, N, 100); + + char *devMemSrc2 = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMemSrc2), + &pitch, width, height)); + REQUIRE(devMemSrc2 != nullptr); + + char* hostMemDst2 = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemDst2 != nullptr); + fillCharHostArray(hostMemDst2, N, 100); + + hipGraphNode_t memsetNode, memcpyNode; + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Prepare memset node + hipMemsetParams initialMemsetParams{}; + initialMemsetParams.dst = reinterpret_cast(devMemSrc1); + initialMemsetParams.value = value; + initialMemsetParams.pitch = pitch; + initialMemsetParams.elementSize = sizeof(char); + initialMemsetParams.width = width; + initialMemsetParams.height = height; + + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, + nullptr, 0, &initialMemsetParams)); + + // Prepare memcpyNode + ::std::vector memcpyNodeDependencies; + memcpyNodeDependencies.push_back(memsetNode); + + hipMemcpy3DParms initialParms{}; + initialParms.srcPos = make_hipPos(0, 0, 0); + initialParms.dstPos = make_hipPos(0, 0, 0); + initialParms.srcPtr = make_hipPitchedPtr(devMemSrc1, pitch, width, height); + initialParms.dstPtr = make_hipPitchedPtr(hostMemDst1, width, width, height); + initialParms.extent = make_hipExtent(width, height, 1); + initialParms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, + memcpyNodeDependencies.data(), + memcpyNodeDependencies.size(), &initialParms)); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + hipMemsetParams newMemsetParams{}; + newMemsetParams.dst = reinterpret_cast(devMemSrc2); + newMemsetParams.value = value; + newMemsetParams.pitch = pitch; + newMemsetParams.elementSize = sizeof(char); + newMemsetParams.width = width; + newMemsetParams.height = height; + + // Validating hipGraphExecMemsetNodeSetParams API + HIP_CHECK(dyn_hipGraphExecMemsetNodeSetParams_ptr(graphExec, memsetNode, + &newMemsetParams)); + + hipMemcpy3DParms newMemcpyParms{}; + newMemcpyParms.srcPos = make_hipPos(0, 0, 0); + newMemcpyParms.dstPos = make_hipPos(0, 0, 0); + newMemcpyParms.srcPtr = make_hipPitchedPtr(devMemSrc2, pitch, width, height); + newMemcpyParms.dstPtr = make_hipPitchedPtr(hostMemDst2, width, + width, height); + newMemcpyParms.extent = make_hipExtent(width, height, 1); + newMemcpyParms.kind = hipMemcpyDeviceToHost; + + // Validating hipGraphExecMemcpyNodeSetParams API + HIP_CHECK(dyn_hipGraphExecMemcpyNodeSetParams_ptr(graphExec, memcpyNode, + &newMemcpyParms)); + + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + #ifdef _WIN32 + HIP_CHECK(hipStreamSynchronize(0)); + #endif + + REQUIRE(validateArrayT(hostMemDst1, N, 100) == true); + REQUIRE(validateArrayT(hostMemDst2, N, 120) == true); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(devMemSrc1)); + HIP_CHECK(hipFree(devMemSrc2)); + free(hostMemDst1); + free(hostMemDst2); +} diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index 02a89bbb6e..71a221eed4 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -106,7 +106,7 @@ if(HIP_PLATFORM MATCHES "amd") hipMemVmm.cc hipArray.cc hipMemcpyDeviceToDeviceNoCU.cc - hipGetProcAddress_Memory_APIs.cc) + hipGetProcAddressMemoryApis.cc) if(UNIX) # Should be compiled for NVIDIA as well after EXSWHTEC-346 is addressed # For windows build error occurs undefined symbol: hipPointerSetAttribute diff --git a/projects/hip-tests/catch/unit/memory/hipGetProcAddress_Memory_APIs.cc b/projects/hip-tests/catch/unit/memory/hipGetProcAddressMemoryApis.cc similarity index 98% rename from projects/hip-tests/catch/unit/memory/hipGetProcAddress_Memory_APIs.cc rename to projects/hip-tests/catch/unit/memory/hipGetProcAddressMemoryApis.cc index 7570a1c523..cb29ea6b24 100644 --- a/projects/hip-tests/catch/unit/memory/hipGetProcAddress_Memory_APIs.cc +++ b/projects/hip-tests/catch/unit/memory/hipGetProcAddressMemoryApis.cc @@ -21,159 +21,7 @@ THE SOFTWARE. #include #include #include "hipMallocManagedCommon.hh" - -/** - * @addtogroup hipGetProcAddress hipGetProcAddress - * @{ - * @ingroup DeviceTest - * `hipGetProcAddress(const char* symbol, void** pfn, - * int hipVersion, uint64_t flags, - * hipDriverProcAddressQueryResult* symbolStatus);` - - * Gets the symbol's function address. - */ - -/** - * Local Function to fill the array with given value - */ -void fillHostArray(int *arr, int size, int value) { - for ( int i = 0; i < size; i++ ) { - arr[i] = value; - } -} - -/** - * Local Function to validate the array with given reference value - */ -bool validateHostArray(int *arr, int size, int refValue) { - for ( int i = 0; i < size; i++ ) { - if ( arr[i] != refValue ) { - return false; - } - } - return true; -} - -/** - * Local Function to fill the character array with given value - */ -void fillCharHostArray(char *arr, int size, int value) { - for ( int i = 0; i < size; i++ ) { - arr[i] = value; - } -} - -/** - * Local Function to validate the array with given reference value - */ -bool validateCharHostArray(char *arr, int size, int refValue) { - for ( int i = 0; i < size; i++ ) { - if ( arr[i] != refValue ) { - return false; - } - } - return true; -} - -/** - * Kernel to validate the array with given reference value - */ -__global__ void verifyArray(int *arr, int size, int refValue, int* status) { - for ( int i = 0; i < size; i++ ) { - if ( arr[i] != refValue ) { - *status = 0; - return; - } - } - *status = 1; - return; -} - -/** - * Local Function to validate the device array with given reference value - */ -bool validateDeviceArray(int *arr, int size, int refValue) { - int *devStatus = nullptr; - HIP_CHECK(hipMalloc(&devStatus, sizeof(int))); - REQUIRE(devStatus != nullptr); - - verifyArray<<<1, 1>>>(arr, size, refValue, devStatus); - int status; - HIP_CHECK(hipMemcpy(&status, devStatus, sizeof(int), hipMemcpyDeviceToHost)); - - HIP_CHECK(hipFree(devStatus)); - - if ( status == 1 ) { - return true; - } else { - return false; - } -} - -/** - * Kernel to fill the array with given value - */ -__global__ void fillArray(int *arr, int size, int value) { - for ( int i = 0; i < size; i++ ) { - arr[i] = value; - } -} - -/** - * Local Function to fill the device array with given value - */ -void fillDeviceArray(int *arr, int size, int value) { - fillArray<<<1, 1>>>(arr, size, value); -} - -/** - * Kernel to validate the character array with given reference value - */ -__global__ void verifyCharArray(char *arr, int size, - int refValue, int* status) { - for ( int i = 0; i < size; i++ ) { - if ( arr[i] != refValue ) { - *status = 0; - return; - } - } - *status = 1; - return; -} - -/** - * Local Function to validate the character device array with - * given reference value - */ -bool validateCharDeviceArray(char *arr, int size, int refValue) { - int *devStatus = nullptr; - HIP_CHECK(hipMalloc(&devStatus, sizeof(int))); - REQUIRE(devStatus != nullptr); - - verifyCharArray<<< 1, 1 >>>(arr, size, refValue, devStatus); - int status; - HIP_CHECK(hipMemcpy(&status, devStatus, sizeof(int), hipMemcpyDeviceToHost)); - - HIP_CHECK(hipFree(devStatus)); - - if ( status == 1 ) { - return true; - } else { - return false; - } -} - -/** - * Local Function to validate the array of different types - */ -template -bool validateArrayT(T *arr, int size, T value) { - for ( int i = 0; i < size; i++ ) { - if ( arr[i] != value ) { - return false; - } - } - return true; -} +#include "../device/hipGetProcAddressHelpers.hh" /** * Test Description diff --git a/projects/hip-tests/catch/unit/module/CMakeLists.txt b/projects/hip-tests/catch/unit/module/CMakeLists.txt index ddf94d4ae6..97a7fd59ec 100644 --- a/projects/hip-tests/catch/unit/module/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/module/CMakeLists.txt @@ -77,7 +77,7 @@ set(TEST_SRC ${TEST_SRC} hipExtModuleLaunchKernel.cc hipHccModuleLaunchKernel.cc - hipGetProcAddress_Module_APIs.cc) + hipGetProcAddressModuleApis.cc) add_custom_target(empty_module.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} diff --git a/projects/hip-tests/catch/unit/module/hipGetProcAddress_Module_APIs.cc b/projects/hip-tests/catch/unit/module/hipGetProcAddressModuleApis.cc similarity index 96% rename from projects/hip-tests/catch/unit/module/hipGetProcAddress_Module_APIs.cc rename to projects/hip-tests/catch/unit/module/hipGetProcAddressModuleApis.cc index 684ed3cfe2..ba95e2c8de 100644 --- a/projects/hip-tests/catch/unit/module/hipGetProcAddress_Module_APIs.cc +++ b/projects/hip-tests/catch/unit/module/hipGetProcAddressModuleApis.cc @@ -22,54 +22,7 @@ THE SOFTWARE. #include #include "hip/hip_ext.h" #include "hip_module_common.hh" - -/** - * @addtogroup hipGetProcAddress hipGetProcAddress - * @{ - * @ingroup DeviceTest - * `hipGetProcAddress(const char* symbol, void** pfn, - int hipVersion, uint64_t flags, - hipDriverProcAddressQueryResult* symbolStatus);` - - * Gets the symbol's function address. - */ - -/** - * Local Function to fill the array with given value - */ -void fillHostArray(int *arr, int size, int value) { - for ( int i = 0; i < size; i++ ) { - arr[i] = value; - } -} - -/** - * Local Function to validate the array with given reference value - */ -bool validateHostArray(int *arr, int size, int refValue) { - for ( int i = 0; i < size; i++ ) { - if ( arr[i] != refValue ) { - return false; - } - } - return true; -} - -/** - * Kernel to add one for each element in array - */ -__global__ void addOneKernel(int *a, int size) { - int offset = blockDim.x * blockIdx.x + threadIdx.x; - int stride = blockDim.x * gridDim.x; - for ( int i = offset; i < size; i+=stride ) { - a[i] += 1; - } -} - -/** - * A simple kernel to check some APIs functionality - */ -__global__ void sampleKernel() { -} +#include "../device/hipGetProcAddressHelpers.hh" /** * Test Description @@ -628,7 +581,7 @@ TEST_CASE("Unit_hipGetProcAddress_ModuleApisCooperativeKernels") { std::vector params(deviceCount); for (int i = 0; i < deviceCount; ++i) { - params[i].func = reinterpret_cast(sampleKernel); + params[i].func = reinterpret_cast(simpleKernel); params[i].gridDim = {1, 1, 1}; params[i].blockDim = {1, 1, 1}; params[i].args = nullptr; @@ -663,7 +616,7 @@ TEST_CASE("Unit_hipGetProcAddress_ModuleApisCooperativeKernels") { std::vector params(deviceCount); for (int i = 0; i < deviceCount; ++i) { - params[i].func = reinterpret_cast(sampleKernel); + params[i].func = reinterpret_cast(simpleKernel); params[i].gridDim = {1, 1, 1}; params[i].blockDim = {1, 1, 1}; params[i].args = nullptr; diff --git a/projects/hip-tests/catch/unit/streamperthread/CMakeLists.txt b/projects/hip-tests/catch/unit/streamperthread/CMakeLists.txt index e185942712..6082bf9191 100644 --- a/projects/hip-tests/catch/unit/streamperthread/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/streamperthread/CMakeLists.txt @@ -8,6 +8,11 @@ set(TEST_SRC hipStreamPerThrdCompilerOptn.cc ) +if(HIP_PLATFORM MATCHES "amd") + set(TEST_SRC ${TEST_SRC} + hipGetProcAddressSptApis.cc) +endif() + if(HIP_PLATFORM MATCHES "amd") set_source_files_properties(hipStreamPerThrdCompilerOptn.cc PROPERTIES COMPILE_OPTIONS "-fgpu-default-stream=per-thread") endif() diff --git a/projects/hip-tests/catch/unit/streamperthread/hipGetProcAddressSptApis.cc b/projects/hip-tests/catch/unit/streamperthread/hipGetProcAddressSptApis.cc new file mode 100644 index 0000000000..ceac3d0cf5 --- /dev/null +++ b/projects/hip-tests/catch/unit/streamperthread/hipGetProcAddressSptApis.cc @@ -0,0 +1,3208 @@ +/* +Copyright (c) 2024 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. +*/ + +#include +#include +#include +#include +#include "../device/hipGetProcAddressHelpers.hh" + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - spt (memory copy) APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/streamperthread/hipGetProcAddress_SPT_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_spt_MemCpy") { + void* hipMemcpy_spt_ptr = nullptr; + void* hipMemcpyAsync_spt_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress("hipMemcpy_spt", + &hipMemcpy_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemcpyAsync_spt", + &hipMemcpyAsync_spt_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipMemcpy_spt_ptr)(void *, const void *, + size_t, hipMemcpyKind) = + reinterpret_cast + (hipMemcpy_spt_ptr); + + hipError_t (*dyn_hipMemcpyAsync_spt_ptr)(void *, const void *, size_t, + hipMemcpyKind, hipStream_t) = + reinterpret_cast + (hipMemcpyAsync_spt_ptr); + + // Validating hipMemcpy_spt API + { + int N = 64; + int Nbytes = N * sizeof(int); + int value = 2; + // With flag hipMemcpyHostToHost + { + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, value); + + int* dstHostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(dstHostMem != nullptr); + HIP_CHECK(dyn_hipMemcpy_spt_ptr(dstHostMem, hostMem, Nbytes, + hipMemcpyHostToHost)); + REQUIRE(validateHostArray(dstHostMem, N, value) == true); + + free(hostMem); + free(dstHostMem); + } + // With flag hipMemcpyHostToDevice + { + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, value); + + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + HIP_CHECK(dyn_hipMemcpy_spt_ptr(devMem, hostMem, Nbytes, + hipMemcpyHostToDevice)); + REQUIRE(validateDeviceArray(devMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + // With flag hipMemcpyDeviceToHost + { + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + fillDeviceArray(devMem, N, value); + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + HIP_CHECK(dyn_hipMemcpy_spt_ptr(hostMem, devMem, Nbytes, + hipMemcpyDeviceToHost)); + REQUIRE(validateHostArray(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + // With flag hipMemcpyDeviceToDevice + { + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + fillDeviceArray(devMem, N, value); + + int* dstDevMem = nullptr; + HIP_CHECK(hipMalloc(&dstDevMem, Nbytes)); + REQUIRE(dstDevMem != nullptr); + HIP_CHECK(dyn_hipMemcpy_spt_ptr(dstDevMem, devMem, Nbytes, + hipMemcpyDeviceToDevice)); + REQUIRE(validateDeviceArray(dstDevMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + HIP_CHECK(hipFree(dstDevMem)); + } + // With flag hipMemcpyDeviceToDeviceNoCU + { + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + fillDeviceArray(devMem, N, value); + + int* dstDevMem = nullptr; + HIP_CHECK(hipMalloc(&dstDevMem, Nbytes)); + REQUIRE(dstDevMem != nullptr); + HIP_CHECK(dyn_hipMemcpy_spt_ptr(dstDevMem, devMem, Nbytes, + hipMemcpyDeviceToDeviceNoCU)); + REQUIRE(validateDeviceArray(dstDevMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + HIP_CHECK(hipFree(dstDevMem)); + } + // With flag hipMemcpyDefault - Host To Host + { + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, value); + + int* dstHostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(dstHostMem != nullptr); + HIP_CHECK(dyn_hipMemcpy_spt_ptr(dstHostMem, hostMem, Nbytes, + hipMemcpyDefault)); + REQUIRE(validateHostArray(dstHostMem, N, value) == true); + + free(hostMem); + free(dstHostMem); + } + // With flag hipMemcpyDefault - Host To Device + { + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, value); + + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + HIP_CHECK(dyn_hipMemcpy_spt_ptr(devMem, hostMem, Nbytes, + hipMemcpyDefault)); + REQUIRE(validateDeviceArray(devMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + // With flag hipMemcpyDefault - Device To Host + { + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + fillDeviceArray(devMem, N, value); + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + HIP_CHECK(dyn_hipMemcpy_spt_ptr(hostMem, devMem, Nbytes, + hipMemcpyDefault)); + REQUIRE(validateHostArray(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + // With flag hipMemcpyDefault - Device To Device + { + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + fillDeviceArray(devMem, N, value); + + int* dstDevMem = nullptr; + HIP_CHECK(hipMalloc(&dstDevMem, Nbytes)); + REQUIRE(dstDevMem != nullptr); + HIP_CHECK(dyn_hipMemcpy_spt_ptr(dstDevMem, devMem, Nbytes, + hipMemcpyDefault)); + REQUIRE(validateDeviceArray(dstDevMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + HIP_CHECK(hipFree(dstDevMem)); + } + } + + // Validating hipMemcpyAsync_spt API + { + int N = 4096; + const int Ns = 4; + int Nbytes = N * sizeof(int); + int value = 2; + + // With flag hipMemcpyHostToHost + { + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, value); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + int* dstHostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(dstHostMem != nullptr); + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + HIP_CHECK(dyn_hipMemcpyAsync_spt_ptr(dstHostMem + startIndex, + hostMem + startIndex, + (Nbytes/Ns), + hipMemcpyHostToHost, + stream[s])); + } + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamSynchronize(stream[s])); + } + + REQUIRE(validateHostArray(dstHostMem, N, value) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + + free(hostMem); + free(dstHostMem); + } + // With flag hipMemcpyHostToDevice + { + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, value); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + HIP_CHECK(dyn_hipMemcpyAsync_spt_ptr(devMem + startIndex, + hostMem + startIndex, + (Nbytes/Ns), + hipMemcpyHostToDevice, + stream[s])); + } + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamSynchronize(stream[s])); + } + + REQUIRE(validateDeviceArray(devMem, N, value) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + // With flag hipMemcpyDeviceToHost + { + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + fillDeviceArray(devMem, N, value); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + HIP_CHECK(dyn_hipMemcpyAsync_spt_ptr(hostMem + startIndex, + devMem + startIndex, + (Nbytes/Ns), + hipMemcpyDeviceToHost, + stream[s])); + } + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamSynchronize(stream[s])); + } + + REQUIRE(validateHostArray(hostMem, N, value) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + // With flag hipMemcpyDeviceToDevice + { + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + fillDeviceArray(devMem, N, value); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + int* dstDevMem = nullptr; + HIP_CHECK(hipMalloc(&dstDevMem, Nbytes)); + REQUIRE(dstDevMem != nullptr); + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + HIP_CHECK(dyn_hipMemcpyAsync_spt_ptr(dstDevMem + startIndex, + devMem + startIndex, + (Nbytes/Ns), + hipMemcpyDeviceToDevice, + stream[s])); + } + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamSynchronize(stream[s])); + } + + REQUIRE(validateDeviceArray(dstDevMem, N, value) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + + HIP_CHECK(hipFree(devMem)); + HIP_CHECK(hipFree(dstDevMem)); + } + // With flag hipMemcpyDeviceToDeviceNoCU + { + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + fillDeviceArray(devMem, N, value); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + int* dstDevMem = nullptr; + HIP_CHECK(hipMalloc(&dstDevMem, Nbytes)); + REQUIRE(dstDevMem != nullptr); + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + HIP_CHECK(dyn_hipMemcpyAsync_spt_ptr(dstDevMem + startIndex, + devMem + startIndex, + (Nbytes/Ns), + hipMemcpyDeviceToDeviceNoCU, + stream[s])); + } + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamSynchronize(stream[s])); + } + + REQUIRE(validateDeviceArray(dstDevMem, N, value) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + + HIP_CHECK(hipFree(devMem)); + HIP_CHECK(hipFree(dstDevMem)); + } + // With flag hipMemcpyDefault - Host To Host + { + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, value); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + int* dstHostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(dstHostMem != nullptr); + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + HIP_CHECK(dyn_hipMemcpyAsync_spt_ptr(dstHostMem + startIndex, + hostMem + startIndex, + (Nbytes/Ns), + hipMemcpyDefault, + stream[s])); + } + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamSynchronize(stream[s])); + } + + REQUIRE(validateHostArray(dstHostMem, N, value) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + + free(hostMem); + free(dstHostMem); + } + // With flag hipMemcpyDefault - Host To Device + { + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, value); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + HIP_CHECK(dyn_hipMemcpyAsync_spt_ptr(devMem + startIndex, + hostMem + startIndex, + (Nbytes/Ns), + hipMemcpyDefault, + stream[s])); + } + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamSynchronize(stream[s])); + } + + REQUIRE(validateDeviceArray(devMem, N, value) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + // With flag hipMemcpyDefault - Device To Host + { + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + fillDeviceArray(devMem, N, value); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + HIP_CHECK(dyn_hipMemcpyAsync_spt_ptr(hostMem + startIndex, + devMem + startIndex, + (Nbytes/Ns), + hipMemcpyDefault, + stream[s])); + } + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamSynchronize(stream[s])); + } + + REQUIRE(validateHostArray(hostMem, N, value) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + // With flag hipMemcpyDefault - Device To Device + { + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + fillDeviceArray(devMem, N, value); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + int* dstDevMem = nullptr; + HIP_CHECK(hipMalloc(&dstDevMem, Nbytes)); + REQUIRE(dstDevMem != nullptr); + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + HIP_CHECK(dyn_hipMemcpyAsync_spt_ptr(dstDevMem + startIndex, + devMem + startIndex, + (Nbytes/Ns), + hipMemcpyDefault, + stream[s])); + } + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamSynchronize(stream[s])); + } + + REQUIRE(validateDeviceArray(dstDevMem, N, value) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + + HIP_CHECK(hipFree(devMem)); + HIP_CHECK(hipFree(dstDevMem)); + } + } +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - spt (memset) APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/streamperthread/hipGetProcAddress_SPT_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_spt_Memset") { + void* hipMemset_spt_ptr = nullptr; + void* hipMemsetAsync_spt_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress("hipMemset_spt", + &hipMemset_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemsetAsync_spt", + &hipMemsetAsync_spt_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipMemset_spt_ptr)(void *, int, size_t) = + reinterpret_cast + (hipMemset_spt_ptr); + hipError_t (*dyn_hipMemsetAsync_spt_ptr)(void *, int, size_t, hipStream_t) = + reinterpret_cast + (hipMemsetAsync_spt_ptr); + + // Validating hipMemset_spt API + { + int N = 16; + int Nbytes = N * sizeof(char); + int value = 10; + + void* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + HIP_CHECK(dyn_hipMemset_spt_ptr(devMem, value, Nbytes)); + + char* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + HIP_CHECK(hipMemcpy(hostMem, devMem, Nbytes, hipMemcpyDeviceToHost)); + + REQUIRE(validateArrayT(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + + // Validating hipMemsetAsync_spt API + { + int N = 16; + int Nbytes = N * sizeof(char); + int value = 126; + const int Ns = 4; + + char* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + HIP_CHECK(dyn_hipMemsetAsync_spt_ptr(devMem + startIndex, + value, N/Ns, stream[s])); + } + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamSynchronize(stream[s])); + } + + char* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + HIP_CHECK(hipMemcpy(hostMem, devMem, Nbytes, hipMemcpyDeviceToHost)); + + REQUIRE(validateArrayT(hostMem, N, value) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - spt (memset 2D and 3D) APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/streamperthread/hipGetProcAddress_SPT_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_spt_Memset2D3D") { + CHECK_IMAGE_SUPPORT + + void* hipMemset2D_spt_ptr = nullptr; + void* hipMemset2DAsync_spt_ptr = nullptr; + void* hipMemset3D_spt_ptr = nullptr; + void* hipMemset3DAsync_spt_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress("hipMemset2D_spt", + &hipMemset2D_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemset2DAsync_spt", + &hipMemset2DAsync_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemset3D_spt", + &hipMemset3D_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemset3DAsync_spt", + &hipMemset3DAsync_spt_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipMemset2D_spt_ptr)(void *, size_t, int, size_t, size_t) = + reinterpret_cast + (hipMemset2D_spt_ptr); + hipError_t (*dyn_hipMemset2DAsync_spt_ptr)(void *, size_t, int, size_t, + size_t, hipStream_t) = + reinterpret_cast + (hipMemset2DAsync_spt_ptr); + hipError_t (*dyn_hipMemset3D_spt_ptr)(hipPitchedPtr, int, hipExtent) = + reinterpret_cast + (hipMemset3D_spt_ptr); + hipError_t (*dyn_hipMemset3DAsync_spt_ptr)(hipPitchedPtr, int, + hipExtent, hipStream_t) = + reinterpret_cast + (hipMemset3DAsync_spt_ptr); + size_t width = 1024; + size_t height = 1024; + size_t depth = 1024; + const int Ns = 4; + int value = 10; + size_t pitch; + + // Validating hipMemset2D_spt API + { + const int N = width * height; + + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + + HIP_CHECK(dyn_hipMemset2D_spt_ptr(devMem, pitch, value, width, height)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + HIP_CHECK(hipMemcpy2D(hostMem, width, devMem, pitch, + width, height, hipMemcpyDeviceToHost)); + + REQUIRE(validateArrayT(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + + // Validating hipMemset2DAsync_spt API + { + const int N = width * height; + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + HIP_CHECK(dyn_hipMemset2DAsync_spt_ptr(devMem + startIndex, pitch/Ns, + value, width/Ns, height/Ns, + stream[s])); + } + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamSynchronize(stream[s])); + } + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + HIP_CHECK(hipMemcpy2D(hostMem, width, devMem, pitch, + width, height, hipMemcpyDeviceToHost)); + + REQUIRE(validateArrayT(hostMem, N, value) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + + // Validating hipMemset3D_spt API + { + const int N = width * height * depth; + + hipPitchedPtr devMem; + hipExtent extent3d{width, height, depth}; + HIP_CHECK(hipMalloc3D(&devMem, extent3d)); + REQUIRE(devMem.ptr != nullptr); + + HIP_CHECK(dyn_hipMemset3D_spt_ptr(devMem, value, extent3d)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(hostMem, width , height, depth); + myparms.srcPtr = devMem; + myparms.extent = extent3d; + myparms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(hipMemcpy3D(&myparms)); + + REQUIRE(validateArrayT(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem.ptr)); + free(hostMem); + } + + // Validating hipMemset3DAsync_spt API + { + const int N = width * height * depth; + + hipPitchedPtr devMem; + hipExtent extent3d{width, height, depth}; + HIP_CHECK(hipMalloc3D(&devMem, extent3d)); + REQUIRE(devMem.ptr != nullptr); + + HIP_CHECK(dyn_hipMemset3DAsync_spt_ptr(devMem, value, extent3d, NULL)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(hostMem, width , height, depth); + myparms.srcPtr = devMem; + myparms.extent = extent3d; + myparms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(hipMemcpy3D(&myparms)); + + REQUIRE(validateArrayT(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem.ptr)); + free(hostMem); + } +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - spt (memory copy 2D) APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/streamperthread/hipGetProcAddress_SPT_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_spt_Memcpy2D") { + CHECK_IMAGE_SUPPORT + + void* hipMemcpy2D_spt_ptr = nullptr; + void* hipMemcpy2DAsync_spt_ptr = nullptr; + void* hipMemcpy2DToArray_spt_ptr = nullptr; + void* hipMemcpy2DToArrayAsync_spt_ptr = nullptr; + void* hipMemcpy2DFromArray_spt_ptr = nullptr; + void* hipMemcpy2DFromArrayAsync_spt_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress("hipMemcpy2D_spt", + &hipMemcpy2D_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemcpy2DAsync_spt", + &hipMemcpy2DAsync_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemcpy2DToArray_spt", + &hipMemcpy2DToArray_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemcpy2DToArrayAsync_spt", + &hipMemcpy2DToArrayAsync_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemcpy2DFromArray_spt", + &hipMemcpy2DFromArray_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemcpy2DFromArrayAsync_spt", + &hipMemcpy2DFromArrayAsync_spt_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipMemcpy2D_spt_ptr)(void *, size_t, const void *, size_t, + size_t, size_t, hipMemcpyKind) = + reinterpret_cast + (hipMemcpy2D_spt_ptr); + hipError_t (*dyn_hipMemcpy2DAsync_spt_ptr)(void *, size_t, const void *, + size_t, size_t, size_t, + hipMemcpyKind, hipStream_t) = + reinterpret_cast + (hipMemcpy2DAsync_spt_ptr); + hipError_t (*dyn_hipMemcpy2DToArray_spt_ptr)(hipArray_t, size_t, size_t, + const void *src, size_t, size_t, + size_t, hipMemcpyKind) = + reinterpret_cast + (hipMemcpy2DToArray_spt_ptr); + hipError_t (*dyn_hipMemcpy2DToArrayAsync_spt_ptr)(hipArray_t, size_t, size_t, + const void *src, size_t, + size_t, size_t, + hipMemcpyKind, + hipStream_t) = + reinterpret_cast + (hipMemcpy2DToArrayAsync_spt_ptr); + hipError_t (*dyn_hipMemcpy2DFromArray_spt_ptr)(void *, size_t, + hipArray_const_t, size_t, + size_t, size_t, size_t, + hipMemcpyKind) = + reinterpret_cast + (hipMemcpy2DFromArray_spt_ptr); + hipError_t (*dyn_hipMemcpy2DFromArrayAsync_spt_ptr)(void *, size_t, + hipArray_const_t, size_t, size_t, size_t, size_t, + hipMemcpyKind, hipStream_t) = + reinterpret_cast + (hipMemcpy2DFromArrayAsync_spt_ptr); + + // Validating hipMemcpy2D_spt API + { + size_t width = 256; + size_t height = 256; + const int N = width * height; + int value = 10; + size_t pitch; + + // With flag hipMemcpyHostToHost + { + char* sHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(sHostMem != nullptr); + fillCharHostArray(sHostMem, N, value); + + char* dHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(dHostMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2D_spt_ptr(dHostMem, width, sHostMem, width, + width, height, hipMemcpyHostToHost)); + + REQUIRE(validateCharHostArray(dHostMem, N, value) == true); + + free(sHostMem); + free(dHostMem); + } + + // With flag hipMemcpyHostToDevice + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2D_spt_ptr(devMem, pitch, hostMem, pitch, + width, height, hipMemcpyHostToDevice)); + + REQUIRE(validateCharDeviceArray(devMem, N, value) == true); + + free(hostMem); + HIP_CHECK(hipFree(devMem)); + } + + // With flag hipMemcpyDeviceToHost + { + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + HIP_CHECK(hipMemset2D(devMem, pitch, value, width, height)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2D_spt_ptr(hostMem, width, devMem, pitch, + width, height, hipMemcpyDeviceToHost)); + + REQUIRE(validateCharHostArray(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + + // With flag hipMemcpyDeviceToDevice + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2D_spt_ptr(dDevMem, dPitch, sDevMem, sPitch, + width, height, hipMemcpyDeviceToDevice)); + + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + } + + // With flag hipMemcpyDeviceToDeviceNoCU + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2D_spt_ptr(dDevMem, dPitch, sDevMem, sPitch, + width, height, hipMemcpyDeviceToDeviceNoCU)); + + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + } + + // With flag hipMemcpyDefault - Host To Host + { + char* sHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(sHostMem != nullptr); + fillCharHostArray(sHostMem, N, value); + + char* dHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(dHostMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2D_spt_ptr(dHostMem, width, sHostMem, width, + width, height, hipMemcpyDefault)); + + REQUIRE(validateCharHostArray(dHostMem, N, value) == true); + + free(sHostMem); + free(dHostMem); + } + + // With flag hipMemcpyDefault - Host To Device + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2D_spt_ptr(devMem, pitch, hostMem, pitch, + width, height, hipMemcpyDefault)); + + REQUIRE(validateCharDeviceArray(devMem, N, value) == true); + + free(hostMem); + HIP_CHECK(hipFree(devMem)); + } + + // With flag hipMemcpyDefault - Device To Host + { + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + HIP_CHECK(hipMemset2D(devMem, pitch, value, width, height)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2D_spt_ptr(hostMem, width, devMem, pitch, + width, height, hipMemcpyDefault)); + + REQUIRE(validateCharHostArray(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + + // With flag hipMemcpyDefault - Device To Device + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2D_spt_ptr(dDevMem, dPitch, sDevMem, sPitch, + width, height, hipMemcpyDefault)); + + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + } + } + + // Validating hipMemcpy2DAsync_spt API + { + size_t width = 256; + size_t height = 256; + const int N = width * height; + int value = 10; + size_t pitch; + + // With flag hipMemcpyHostToHost + { + char* sHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(sHostMem != nullptr); + fillCharHostArray(sHostMem, N, value); + + char* dHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(dHostMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DAsync_spt_ptr(dHostMem, width, sHostMem, width, + width, height, hipMemcpyHostToHost, NULL)); + + REQUIRE(validateCharHostArray(dHostMem, N, value) == true); + + free(sHostMem); + free(dHostMem); + } + + // With flag hipMemcpyHostToDevice + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DAsync_spt_ptr(devMem, pitch, hostMem, pitch, + width, height, hipMemcpyHostToDevice, NULL)); + + REQUIRE(validateCharDeviceArray(devMem, N, value) == true); + + free(hostMem); + HIP_CHECK(hipFree(devMem)); + } + + // With flag hipMemcpyDeviceToHost + { + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + HIP_CHECK(hipMemset2D(devMem, pitch, value, width, height)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DAsync_spt_ptr(hostMem, width, devMem, pitch, + width, height, hipMemcpyDeviceToHost, NULL)); + + REQUIRE(validateCharHostArray(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + + // With flag hipMemcpyDeviceToDevice + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DAsync_spt_ptr(dDevMem, dPitch, sDevMem, sPitch, + width, height, hipMemcpyDeviceToDevice, NULL)); + + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + } + + // With flag hipMemcpyDeviceToDeviceNoCU + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DAsync_spt_ptr(dDevMem, dPitch, sDevMem, sPitch, + width, height, hipMemcpyDeviceToDeviceNoCU, NULL)); + + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + } + + // With flag hipMemcpyDefault - Host To Host + { + char* sHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(sHostMem != nullptr); + fillCharHostArray(sHostMem, N, value); + + char* dHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(dHostMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DAsync_spt_ptr(dHostMem, width, sHostMem, width, + width, height, hipMemcpyDefault, NULL)); + + REQUIRE(validateCharHostArray(dHostMem, N, value) == true); + + free(sHostMem); + free(dHostMem); + } + + // With flag hipMemcpyDefault - Host To Device + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DAsync_spt_ptr(devMem, pitch, hostMem, pitch, + width, height, hipMemcpyDefault, NULL)); + + REQUIRE(validateCharDeviceArray(devMem, N, value) == true); + + free(hostMem); + HIP_CHECK(hipFree(devMem)); + } + + // With flag hipMemcpyDefault - Device To Host + { + char *devMem = nullptr; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMem), + &pitch, width, height)); + REQUIRE(devMem != nullptr); + HIP_CHECK(hipMemset2D(devMem, pitch, value, width, height)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DAsync_spt_ptr(hostMem, width, devMem, pitch, + width, height, hipMemcpyDefault, NULL)); + + REQUIRE(validateCharHostArray(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } + + // With flag hipMemcpyDefault - Device To Device + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DAsync_spt_ptr(dDevMem, dPitch, sDevMem, sPitch, + width, height, hipMemcpyDefault, NULL)); + + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + } + } + + // Validating hipMemcpy2DToArray_spt API + { + size_t width = 256; + size_t height = 256; + const int N = width * height; + int value = 10; + + // With flag hipMemcpyHostToDevice + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DToArray_spt_ptr(array, 0, 0, + hostMem, width, width, height, + hipMemcpyHostToDevice)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(hipMemcpy2DFromArray(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + free(hostMem); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDeviceToDevice + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DToArray_spt_ptr(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDeviceToDevice)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(hipMemcpy2DFromArray(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDeviceToDeviceNoCU + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DToArray_spt_ptr(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDeviceToDeviceNoCU)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(hipMemcpy2DFromArray(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDefault - Host To Device + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DToArray_spt_ptr(array, 0, 0, + hostMem, width, width, height, + hipMemcpyDefault)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(hipMemcpy2DFromArray(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + free(hostMem); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDefault - Device To Device + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DToArray_spt_ptr(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDefault)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(hipMemcpy2DFromArray(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + } + + // Validating hipMemcpy2DToArrayAsync_spt API + { + size_t width = 256; + size_t height = 256; + const int N = width * height; + int value = 10; + + // With flag hipMemcpyHostToDevice + { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DToArrayAsync_spt_ptr(array, 0, 0, + hostMem, width, width, height, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(hipMemcpy2DFromArray(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + free(hostMem); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDeviceToDevice + { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DToArrayAsync_spt_ptr(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDeviceToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(hipMemcpy2DFromArray(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(sDevMem)); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDeviceToDeviceNoCU + { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DToArrayAsync_spt_ptr(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDeviceToDeviceNoCU, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(hipMemcpy2DFromArray(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(sDevMem)); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDefault - Host To Device + { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DToArrayAsync_spt_ptr(array, 0, 0, + hostMem, width, width, height, + hipMemcpyDefault, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(hipMemcpy2DFromArray(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + free(hostMem); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDefault - Device To Device + { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DToArrayAsync_spt_ptr(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDefault, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(hipMemcpy2DFromArray(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(sDevMem)); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + } + + // Validating hipMemcpy2DFromArray API + { + size_t width = 256; + size_t height = 256; + const int N = width * height; + int value = 10; + + // With flag hipMemcpyDeviceToHost + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, + hostMem, width, width, height, + hipMemcpyHostToDevice)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DFromArray_spt_ptr(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + free(hostMem); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDeviceToDevice + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDeviceToDevice)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DFromArray_spt_ptr(dDevMem, width, + array, 0, 0, width, height, + hipMemcpyDeviceToDevice)); + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDeviceToDeviceNoCU + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDeviceToDevice)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DFromArray_spt_ptr(dDevMem, width, + array, 0, 0, width, height, + hipMemcpyDeviceToDeviceNoCU)); + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDefault - Device To Host + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, + hostMem, width, width, height, + hipMemcpyHostToDevice)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DFromArray_spt_ptr(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDefault)); + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + free(hostMem); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDefault - Device To Device + { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDeviceToDevice)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DFromArray_spt_ptr(dDevMem, width, + array, 0, 0, width, height, + hipMemcpyDefault)); + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + HIP_CHECK(hipFreeArray(array)); + } + } + + // Validating hipMemcpy2DFromArrayAsync API + { + size_t width = 256; + size_t height = 256; + const int N = width * height; + int value = 10; + + // With flag hipMemcpyDeviceToHost + { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, + hostMem, width, width, height, + hipMemcpyHostToDevice)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DFromArrayAsync_spt_ptr(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + free(hostMem); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDeviceToDevice + { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDeviceToDevice)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DFromArrayAsync_spt_ptr(dDevMem, width, + array, 0, 0, width, height, + hipMemcpyDeviceToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDeviceToDeviceNoCU + { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDeviceToDevice)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DFromArrayAsync_spt_ptr(dDevMem, width, + array, 0, 0, width, height, + hipMemcpyDeviceToDeviceNoCU, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDefault - Device To Host + { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, + hostMem, width, width, height, + hipMemcpyHostToDevice)); + + char* hostMemory = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMemory != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DFromArrayAsync_spt_ptr(hostMemory, width, + array, 0, 0, width, height, + hipMemcpyDefault, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharHostArray(hostMemory, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + free(hostMem); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); + } + // With flag hipMemcpyDefault - Device To Device + { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + &sPitch, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, + sDevMem, sPitch, width, height, + hipMemcpyDeviceToDevice)); + + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + &dPitch, width, height)); + REQUIRE(dDevMem != nullptr); + + HIP_CHECK(dyn_hipMemcpy2DFromArrayAsync_spt_ptr(dDevMem, width, + array, 0, 0, width, height, + hipMemcpyDefault, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + HIP_CHECK(hipFreeArray(array)); + } + } +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - spt (memory copy 3D) APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/streamperthread/hipGetProcAddress_SPT_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_spt_Memcpy3D") { + CHECK_IMAGE_SUPPORT + + void* hipMemcpy3D_spt_ptr = nullptr; + void* hipMemcpy3DAsync_spt_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress("hipMemcpy3D_spt", + &hipMemcpy3D_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemcpy3DAsync_spt", + &hipMemcpy3DAsync_spt_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipMemcpy3D_spt_ptr)(const struct hipMemcpy3DParms *) = + reinterpret_cast + (hipMemcpy3D_spt_ptr); + hipError_t (*dyn_hipMemcpy3DAsync_spt_ptr)(const struct hipMemcpy3DParms *, + hipStream_t) = + reinterpret_cast + (hipMemcpy3DAsync_spt_ptr); + + // Validating hipMemcpy3D_spt API + { + size_t width = 256; + size_t height = 256; + size_t depth = 256; + const int N = width * height * depth; + int value = 10; + hipExtent extent3d{width, height, depth}; + + // With flag hipMemcpyHostToHost + { + char* sHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(sHostMem != nullptr); + fillCharHostArray(sHostMem, N, value); + + char* dHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(dHostMem != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = make_hipPitchedPtr(sHostMem, width , height, depth); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(dHostMem, width , height, depth); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyHostToHost; + HIP_CHECK(dyn_hipMemcpy3D_spt_ptr(&myparms)); + + REQUIRE(validateCharHostArray(dHostMem, N, value) == true); + + free(sHostMem); + free(dHostMem); + } + + // With flag hipMemcpyHostToDevice + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipPitchedPtr devMem; + HIP_CHECK(hipMalloc3D(&devMem, extent3d)); + REQUIRE(devMem.ptr != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = make_hipPitchedPtr(hostMem, width , height, depth); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = devMem; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyHostToDevice; + HIP_CHECK(dyn_hipMemcpy3D_spt_ptr(&myparms)); + + REQUIRE(validateCharDeviceArray(reinterpret_cast(devMem.ptr), + N, value) == true); + + HIP_CHECK(hipFree(devMem.ptr)); + free(hostMem); + } + + // With flag hipMemcpyDeviceToHost + { + hipPitchedPtr devMem; + HIP_CHECK(hipMalloc3D(&devMem, extent3d)); + REQUIRE(devMem.ptr != nullptr); + HIP_CHECK(hipMemset3D(devMem, value, extent3d)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = devMem; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(hostMem, width , height, depth); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(dyn_hipMemcpy3D_spt_ptr(&myparms)); + + REQUIRE(validateCharHostArray(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem.ptr)); + free(hostMem); + } + + // With flag hipMemcpyDeviceToDevice + { + hipPitchedPtr sDevMem; + HIP_CHECK(hipMalloc3D(&sDevMem, extent3d)); + REQUIRE(sDevMem.ptr != nullptr); + HIP_CHECK(hipMemset3D(sDevMem, value, extent3d)); + + hipPitchedPtr dDevMem; + HIP_CHECK(hipMalloc3D(&dDevMem, extent3d)); + REQUIRE(dDevMem.ptr != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = sDevMem; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = dDevMem; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDeviceToDevice; + HIP_CHECK(dyn_hipMemcpy3D_spt_ptr(&myparms)); + + REQUIRE(validateCharDeviceArray(reinterpret_cast(dDevMem.ptr), + N, value) == true); + + HIP_CHECK(hipFree(sDevMem.ptr)); + HIP_CHECK(hipFree(dDevMem.ptr)); + } + + // With flag hipMemcpyDeviceToDeviceNoCU + { + hipPitchedPtr sDevMem; + HIP_CHECK(hipMalloc3D(&sDevMem, extent3d)); + REQUIRE(sDevMem.ptr != nullptr); + HIP_CHECK(hipMemset3D(sDevMem, value, extent3d)); + + hipPitchedPtr dDevMem; + HIP_CHECK(hipMalloc3D(&dDevMem, extent3d)); + REQUIRE(dDevMem.ptr != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = sDevMem; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = dDevMem; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDeviceToDeviceNoCU; + HIP_CHECK(dyn_hipMemcpy3D_spt_ptr(&myparms)); + + REQUIRE(validateCharDeviceArray(reinterpret_cast(dDevMem.ptr), + N, value) == true); + + HIP_CHECK(hipFree(sDevMem.ptr)); + HIP_CHECK(hipFree(dDevMem.ptr)); + } + + // With flag hipMemcpyDefault - Host To Host + { + char* sHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(sHostMem != nullptr); + fillCharHostArray(sHostMem, N, value); + + char* dHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(dHostMem != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = make_hipPitchedPtr(sHostMem, width , height, depth); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(dHostMem, width , height, depth); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDefault; + HIP_CHECK(dyn_hipMemcpy3D_spt_ptr(&myparms)); + + REQUIRE(validateCharHostArray(dHostMem, N, value) == true); + + free(sHostMem); + free(dHostMem); + } + + // With flag hipMemcpyDefault - Host To Device + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipPitchedPtr devMem; + HIP_CHECK(hipMalloc3D(&devMem, extent3d)); + REQUIRE(devMem.ptr != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = make_hipPitchedPtr(hostMem, width , height, depth); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = devMem; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDefault; + HIP_CHECK(dyn_hipMemcpy3D_spt_ptr(&myparms)); + + REQUIRE(validateCharDeviceArray(reinterpret_cast(devMem.ptr), + N, value) == true); + + HIP_CHECK(hipFree(devMem.ptr)); + free(hostMem); + } + + // With flag hipMemcpyDefault - Device To Host + { + hipPitchedPtr devMem; + HIP_CHECK(hipMalloc3D(&devMem, extent3d)); + REQUIRE(devMem.ptr != nullptr); + HIP_CHECK(hipMemset3D(devMem, value, extent3d)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = devMem; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(hostMem, width , height, depth); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDefault; + HIP_CHECK(dyn_hipMemcpy3D_spt_ptr(&myparms)); + + REQUIRE(validateCharHostArray(hostMem, N, value) == true); + + HIP_CHECK(hipFree(devMem.ptr)); + free(hostMem); + } + + // With flag hipMemcpyDefault - Device To Device + { + hipPitchedPtr sDevMem; + HIP_CHECK(hipMalloc3D(&sDevMem, extent3d)); + REQUIRE(sDevMem.ptr != nullptr); + HIP_CHECK(hipMemset3D(sDevMem, value, extent3d)); + + hipPitchedPtr dDevMem; + HIP_CHECK(hipMalloc3D(&dDevMem, extent3d)); + REQUIRE(dDevMem.ptr != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = sDevMem; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = dDevMem; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDefault; + HIP_CHECK(dyn_hipMemcpy3D_spt_ptr(&myparms)); + + REQUIRE(validateCharDeviceArray(reinterpret_cast(dDevMem.ptr), + N, value) == true); + + HIP_CHECK(hipFree(sDevMem.ptr)); + HIP_CHECK(hipFree(dDevMem.ptr)); + } + } + + // Validating hipMemcpy3DAsync_spt API + { + size_t width = 256; + size_t height = 256; + size_t depth = 256; + const int N = width * height * depth; + int value = 10; + hipExtent extent3d{width, height, depth}; + + // With flag hipMemcpyHostToHost + { + char* sHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(sHostMem != nullptr); + fillCharHostArray(sHostMem, N, value); + + char* dHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(dHostMem != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = make_hipPitchedPtr(sHostMem, width , height, depth); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(dHostMem, width , height, depth); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyHostToHost; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipMemcpy3DAsync_spt_ptr(&myparms, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharHostArray(dHostMem, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + free(sHostMem); + free(dHostMem); + } + + // With flag hipMemcpyHostToDevice + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipPitchedPtr devMem; + HIP_CHECK(hipMalloc3D(&devMem, extent3d)); + REQUIRE(devMem.ptr != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = make_hipPitchedPtr(hostMem, width , height, depth); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = devMem; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyHostToDevice; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipMemcpy3DAsync_spt_ptr(&myparms, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharDeviceArray(reinterpret_cast(devMem.ptr), + N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(devMem.ptr)); + free(hostMem); + } + + // With flag hipMemcpyDeviceToHost + { + hipPitchedPtr devMem; + HIP_CHECK(hipMalloc3D(&devMem, extent3d)); + REQUIRE(devMem.ptr != nullptr); + HIP_CHECK(hipMemset3D(devMem, value, extent3d)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = devMem; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(hostMem, width , height, depth); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDeviceToHost; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipMemcpy3DAsync_spt_ptr(&myparms, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharHostArray(hostMem, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(devMem.ptr)); + free(hostMem); + } + + // With flag hipMemcpyDeviceToDevice + { + hipPitchedPtr sDevMem; + HIP_CHECK(hipMalloc3D(&sDevMem, extent3d)); + REQUIRE(sDevMem.ptr != nullptr); + HIP_CHECK(hipMemset3D(sDevMem, value, extent3d)); + + hipPitchedPtr dDevMem; + HIP_CHECK(hipMalloc3D(&dDevMem, extent3d)); + REQUIRE(dDevMem.ptr != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = sDevMem; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = dDevMem; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDeviceToDevice; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipMemcpy3DAsync_spt_ptr(&myparms, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharDeviceArray(reinterpret_cast(dDevMem.ptr), + N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(sDevMem.ptr)); + HIP_CHECK(hipFree(dDevMem.ptr)); + } + + // With flag hipMemcpyDeviceToDeviceNoCU + { + hipPitchedPtr sDevMem; + HIP_CHECK(hipMalloc3D(&sDevMem, extent3d)); + REQUIRE(sDevMem.ptr != nullptr); + HIP_CHECK(hipMemset3D(sDevMem, value, extent3d)); + + hipPitchedPtr dDevMem; + HIP_CHECK(hipMalloc3D(&dDevMem, extent3d)); + REQUIRE(dDevMem.ptr != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = sDevMem; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = dDevMem; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDeviceToDeviceNoCU; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipMemcpy3DAsync_spt_ptr(&myparms, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharDeviceArray(reinterpret_cast(dDevMem.ptr), + N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(sDevMem.ptr)); + HIP_CHECK(hipFree(dDevMem.ptr)); + } + + // With flag hipMemcpyDefault - Host To Host + { + char* sHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(sHostMem != nullptr); + fillCharHostArray(sHostMem, N, value); + + char* dHostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(dHostMem != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = make_hipPitchedPtr(sHostMem, width , height, depth); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(dHostMem, width , height, depth); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDefault; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipMemcpy3DAsync_spt_ptr(&myparms, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharHostArray(dHostMem, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + free(sHostMem); + free(dHostMem); + } + + // With flag hipMemcpyDefault - Host To Device + { + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + fillCharHostArray(hostMem, N, value); + + hipPitchedPtr devMem; + HIP_CHECK(hipMalloc3D(&devMem, extent3d)); + REQUIRE(devMem.ptr != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = make_hipPitchedPtr(hostMem, width , height, depth); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = devMem; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDefault; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipMemcpy3DAsync_spt_ptr(&myparms, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharDeviceArray(reinterpret_cast(devMem.ptr), + N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(devMem.ptr)); + free(hostMem); + } + + // With flag hipMemcpyDefault - Device To Host + { + hipPitchedPtr devMem; + HIP_CHECK(hipMalloc3D(&devMem, extent3d)); + REQUIRE(devMem.ptr != nullptr); + HIP_CHECK(hipMemset3D(devMem, value, extent3d)); + + char* hostMem = reinterpret_cast(malloc( N * sizeof(char))); + REQUIRE(hostMem != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = devMem; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(hostMem, width , height, depth); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDefault; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipMemcpy3DAsync_spt_ptr(&myparms, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharHostArray(hostMem, N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(devMem.ptr)); + free(hostMem); + } + + // With flag hipMemcpyDefault - Device To Device + { + hipPitchedPtr sDevMem; + HIP_CHECK(hipMalloc3D(&sDevMem, extent3d)); + REQUIRE(sDevMem.ptr != nullptr); + HIP_CHECK(hipMemset3D(sDevMem, value, extent3d)); + + hipPitchedPtr dDevMem; + HIP_CHECK(hipMalloc3D(&dDevMem, extent3d)); + REQUIRE(dDevMem.ptr != nullptr); + + hipMemcpy3DParms myparms{}; + myparms.srcPtr = sDevMem; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPtr = dDevMem; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = extent3d; + myparms.kind = hipMemcpyDefault; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipMemcpy3DAsync_spt_ptr(&myparms, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(validateCharDeviceArray(reinterpret_cast(dDevMem.ptr), + N, value) == true); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(sDevMem.ptr)); + HIP_CHECK(hipFree(dDevMem.ptr)); + } + } +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - spt (kernel launch) APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/streamperthread/hipGetProcAddress_SPT_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_spt_LaunchKernel") { + void* hipLaunchKernel_spt_ptr = nullptr; + void* hipLaunchHostFunc_spt_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipLaunchKernel_spt", + &hipLaunchKernel_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipLaunchHostFunc_spt", + &hipLaunchHostFunc_spt_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipLaunchKernel_spt_ptr)(const void *, dim3, dim3, void **, + size_t, hipStream_t) = + reinterpret_cast + (hipLaunchKernel_spt_ptr); + hipError_t (*dyn_hipLaunchHostFunc_spt_ptr)(hipStream_t, + hipHostFn_t, void *) = + reinterpret_cast + (hipLaunchHostFunc_spt_ptr); + + const int N = 10; + const int Nbytes = 10 * sizeof(int); + + int *hostArr = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostArr != nullptr); + fillHostArray(hostArr, N, 10); + + int *devArr = nullptr; + HIP_CHECK(hipMalloc(&devArr, Nbytes)); + REQUIRE(devArr != nullptr); + HIP_CHECK(hipMemcpy(devArr, hostArr, Nbytes, hipMemcpyHostToDevice)); + + dim3 blocksPerGrid(1, 1, 1); + dim3 threadsPerBlock(1, 1, N); + + struct kernelArgs { + void *arr; + int size; + }; + kernelArgs kernelArg; + kernelArg.arr = devArr; + kernelArg.size = N; + void *kernel_args[] = {&kernelArg.arr, &kernelArg.size}; + + // Validating hipLaunchKernel_spt API + { + HIP_CHECK(dyn_hipLaunchKernel_spt_ptr( + reinterpret_cast(addOneKernel), + blocksPerGrid, threadsPerBlock, kernel_args, 0, nullptr)); + + HIP_CHECK(hipMemcpy(hostArr, devArr, Nbytes, hipMemcpyDeviceToHost)); + REQUIRE(validateHostArray(hostArr, N, 11) == true); + } + + // Validating hipLaunchHostFunc_spt API + { + int data = 30; + hipHostFn_t fn = addTen; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipLaunchHostFunc_spt_ptr(stream, fn, + reinterpret_cast(&data))); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(data == 40); + HIP_CHECK(hipStreamDestroy(stream)); + } + free(hostArr); + HIP_CHECK(hipFree(devArr)); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of + * - hipLaunchCooperativeKernel_spt from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/streamperthread/hipGetProcAddress_SPT_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_spt_LaunchCooperativeKernel") { + hipDeviceProp_t device_properties; + HIP_CHECK(hipGetDeviceProperties(&device_properties, 0)); + + if (!device_properties.cooperativeLaunch) { + HipTest::HIP_SKIP_TEST("Skipping since cooperative launch not supported"); + return; + } + + void* hipLaunchCooperativeKernel_spt_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipLaunchCooperativeKernel_spt", + &hipLaunchCooperativeKernel_spt_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipLaunchCooperativeKernel_spt_ptr)(const void *, dim3, + dim3, void **, unsigned int, hipStream_t) = + reinterpret_cast + (hipLaunchCooperativeKernel_spt_ptr); + + const int N = 10; + const int Nbytes = 10 * sizeof(int); + + int *hostArr = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostArr != nullptr); + fillHostArray(hostArr, N, 10); + + int *devArr = nullptr; + HIP_CHECK(hipMalloc(&devArr, Nbytes)); + REQUIRE(devArr != nullptr); + HIP_CHECK(hipMemcpy(devArr, hostArr, Nbytes, hipMemcpyHostToDevice)); + + dim3 blocksPerGrid(1, 1, 1); + dim3 threadsPerBlock(1, 1, N); + + struct kernelArgs { + void *arr; + int size; + }; + kernelArgs kernelArg; + kernelArg.arr = devArr; + kernelArg.size = N; + void *kernel_args[] = {&kernelArg.arr, &kernelArg.size}; + + // Validating hipLaunchCooperativeKernel_spt API + HIP_CHECK(dyn_hipLaunchCooperativeKernel_spt_ptr( + reinterpret_cast(addOneKernel), + blocksPerGrid, threadsPerBlock, + kernel_args, 0, 0)); + HIP_CHECK(hipMemcpy(hostArr, devArr, Nbytes, hipMemcpyDeviceToHost)); + REQUIRE(validateHostArray(hostArr, N, 11) == true); + + free(hostArr); + HIP_CHECK(hipFree(devArr)); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - spt (streams) APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/streamperthread/hipGetProcAddress_SPT_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_spt_Stream") { + void* hipStreamGetFlags_spt_ptr = nullptr; + void* hipStreamGetPriority_spt_ptr = nullptr; + void* hipStreamSynchronize_spt_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipStreamGetFlags_spt", + &hipStreamGetFlags_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamGetPriority_spt", + &hipStreamGetPriority_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamSynchronize_spt", + &hipStreamSynchronize_spt_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipStreamGetFlags_spt_ptr)(hipStream_t, unsigned int *) = + reinterpret_cast + (hipStreamGetFlags_spt_ptr); + hipError_t (*dyn_hipStreamGetPriority_spt_ptr)(hipStream_t, int *) = + reinterpret_cast + (hipStreamGetPriority_spt_ptr); + hipError_t (*dyn_hipStreamSynchronize_spt_ptr)(hipStream_t) = + reinterpret_cast + (hipStreamSynchronize_spt_ptr); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Validating hipStreamGetFlags_spt API + { + unsigned int flagsSpt, flagsSptWithPtr; + HIP_CHECK(hipStreamGetFlags_spt(stream, &flagsSpt)); + HIP_CHECK(dyn_hipStreamGetFlags_spt_ptr(stream, &flagsSptWithPtr)); + REQUIRE(flagsSptWithPtr == flagsSpt); + } + + // Validating hipStreamGetPriority_spt API + { + int prioritySpt, prioritySptWithPtr; + HIP_CHECK(hipStreamGetPriority_spt(stream, &prioritySpt)); + HIP_CHECK(dyn_hipStreamGetPriority_spt_ptr(stream, &prioritySptWithPtr)); + REQUIRE(prioritySptWithPtr == prioritySpt); + } + + HIP_CHECK(hipStreamDestroy(stream)); + + + // Validating hipStreamSynchronize_spt API + { + int N = 1024 * 1024; + const int Ns = 4; + int Nbytes = N * sizeof(int); + int value = 10; + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, value); + + hipStream_t stream[Ns]; + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamCreate(&stream[s])); + } + + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + HIP_CHECK(hipMemcpy(devMem, hostMem, Nbytes, hipMemcpyHostToDevice)); + + for ( int s = 0; s < Ns; s++ ) { + int startIndex = s * (N/Ns); + addOneKernel<<< 1, 1, 0, stream[s] >>>(devMem + startIndex , N/Ns); + } + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(dyn_hipStreamSynchronize_spt_ptr(stream[s])); + } + + HIP_CHECK(hipMemcpy(hostMem, devMem, Nbytes, hipMemcpyDeviceToHost)); + + REQUIRE(validateHostArray(hostMem, N, 11) == true); + + for ( int s = 0; s < Ns; s++ ) { + HIP_CHECK(hipStreamDestroy(stream[s])); + } + HIP_CHECK(hipFree(devMem)); + free(hostMem); + } +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - spt (Graph) APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/streamperthread/hipGetProcAddress_SPT_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_spt_Graph") { + void* hipStreamBeginCapture_spt_ptr = nullptr; + void* hipStreamIsCapturing_spt_ptr = nullptr; + void* hipStreamEndCapture_spt_ptr = nullptr; + void* hipStreamQuery_spt_ptr = nullptr; + void* hipStreamGetCaptureInfo_spt_ptr = nullptr; + void* hipStreamGetCaptureInfo_v2_spt_ptr = nullptr; + void* hipStreamAddCallback_spt_ptr = nullptr; + void* hipGraphLaunch_spt_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipStreamBeginCapture_spt", + &hipStreamBeginCapture_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamIsCapturing_spt", + &hipStreamIsCapturing_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamEndCapture_spt", + &hipStreamEndCapture_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamQuery_spt", + &hipStreamQuery_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamGetCaptureInfo_spt", + &hipStreamGetCaptureInfo_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamGetCaptureInfo_v2_spt", + &hipStreamGetCaptureInfo_v2_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamAddCallback_spt", + &hipStreamAddCallback_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipGraphLaunch_spt", + &hipGraphLaunch_spt_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipStreamBeginCapture_spt_ptr)( + hipStream_t, hipStreamCaptureMode) = + reinterpret_cast + (hipStreamBeginCapture_spt_ptr); + + hipError_t (*dyn_hipStreamIsCapturing_spt_ptr)( + hipStream_t, hipStreamCaptureStatus *) = + reinterpret_cast + (hipStreamIsCapturing_spt_ptr); + + hipError_t (*dyn_hipStreamEndCapture_spt_ptr)(hipStream_t, hipGraph_t *) = + reinterpret_cast + (hipStreamEndCapture_spt_ptr); + + hipError_t (*dyn_hipStreamQuery_spt_ptr)(hipStream_t) = + reinterpret_cast + (hipStreamQuery_spt_ptr); + + hipError_t (*dyn_hipStreamGetCaptureInfo_spt_ptr)( + hipStream_t, hipStreamCaptureStatus *, + unsigned long long *) = // NOLINT + reinterpret_cast // NOLINT + (hipStreamGetCaptureInfo_spt_ptr); + + hipError_t (*dyn_hipStreamGetCaptureInfo_v2_spt_ptr)( + hipStream_t, hipStreamCaptureStatus *, + unsigned long long *, // NOLINT + hipGraph_t *, const hipGraphNode_t **, size_t *) = + reinterpret_cast + (hipStreamGetCaptureInfo_v2_spt_ptr); + + hipError_t (*dyn_hipStreamAddCallback_spt_ptr)( + hipStream_t, hipStreamCallback_t, void *, unsigned int) = + reinterpret_cast + (hipStreamAddCallback_spt_ptr); + + hipError_t (*dyn_hipGraphLaunch_spt_ptr)(hipGraphExec_t, hipStream_t) = + reinterpret_cast + (hipGraphLaunch_spt_ptr); + + int N = 40; + int Nbytes = N * sizeof(int); + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, 10); + + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + + hipGraph_t graph = nullptr; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Validating hipStreamBeginCapture_spt API + HIP_CHECK(dyn_hipStreamBeginCapture_spt_ptr(stream, + hipStreamCaptureModeGlobal)); + + // Validating hipStreamIsCapturing_spt API + hipStreamCaptureStatus pCaptureStatus = hipStreamCaptureStatusNone; + HIP_CHECK(dyn_hipStreamIsCapturing_spt_ptr(stream, &pCaptureStatus)); + REQUIRE(pCaptureStatus == hipStreamCaptureStatusActive); + + HIP_CHECK(hipMemcpyAsync(devMem, hostMem, Nbytes, + hipMemcpyHostToDevice, stream)); + addOneKernel<<< 1, 1, 0, stream >>>(devMem , N); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, Nbytes, + hipMemcpyDeviceToHost, stream)); + + // Validating hipStreamGetCaptureInfo_spt API + hipStreamCaptureStatus + pCaptureStatusWithOrgApi = hipStreamCaptureStatusNone, + pCaptureStatusWithFuncPtr = hipStreamCaptureStatusNone; + + unsigned long long pIdWithOrgApi = 0, pIdWithFuncPtr = 0; // NOLINT + + HIP_CHECK(hipStreamGetCaptureInfo_spt(stream, + &pCaptureStatusWithOrgApi, &pIdWithOrgApi)); + HIP_CHECK(dyn_hipStreamGetCaptureInfo_spt_ptr(stream, + &pCaptureStatusWithFuncPtr, &pIdWithFuncPtr)); + + REQUIRE(pCaptureStatusWithFuncPtr == pCaptureStatusWithOrgApi); + REQUIRE(pIdWithFuncPtr == pIdWithOrgApi); + + // Validating hipStreamGetCaptureInfo_v2_spt API + hipStreamCaptureStatus captureStatus_out_org, captureStatus_out_ptr; + unsigned long long id_out_org, id_out_ptr; // NOLINT + hipGraph_t graph_out_org, graph_out_ptr; + const hipGraphNode_t *dependencies_out_org{}; + const hipGraphNode_t *dependencies_out_ptr{}; + size_t numDependencies_out_org, numDependencies_out_ptr; + + HIP_CHECK(hipStreamGetCaptureInfo_v2_spt(stream, &captureStatus_out_org, + &id_out_org, &graph_out_org, + &dependencies_out_org, &numDependencies_out_org)); + HIP_CHECK(dyn_hipStreamGetCaptureInfo_v2_spt_ptr(stream, + &captureStatus_out_ptr, &id_out_ptr, &graph_out_ptr, + &dependencies_out_ptr, &numDependencies_out_ptr)); + + REQUIRE(captureStatus_out_ptr == captureStatus_out_org); + REQUIRE(id_out_ptr == id_out_org); + REQUIRE(graph_out_ptr == graph_out_org); + + // Validating hipStreamEndCapture_spt API + HIP_CHECK(dyn_hipStreamEndCapture_spt_ptr(stream, &graph)); + + // Validating hipStreamAddCallback_spt API + int data = 200; + HIP_CHECK(dyn_hipStreamAddCallback_spt_ptr(stream, callBackFunction, + reinterpret_cast(&data), 0)); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + // Validating hipGraphLaunch_spt API + HIP_CHECK(dyn_hipGraphLaunch_spt_ptr(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Validating hipStreamQuery_spt API + REQUIRE(dyn_hipStreamQuery_spt_ptr(stream) == hipSuccess); + + REQUIRE(validateHostArray(hostMem, N, 11) == true); + REQUIRE(data == 300); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); + free(hostMem); + HIP_CHECK(hipFree(devMem)); +} + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - spt (Events) APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/streamperthread/hipGetProcAddress_SPT_APIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_spt_Event") { + void* hipEventRecord_spt_ptr = nullptr; + void* hipStreamWaitEvent_spt_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress( + "hipEventRecord_spt", + &hipEventRecord_spt_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress( + "hipStreamWaitEvent_spt", + &hipStreamWaitEvent_spt_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t (*dyn_hipEventRecord_spt_ptr)(hipEvent_t, hipStream_t) = + reinterpret_cast + (hipEventRecord_spt_ptr); + + hipError_t (*dyn_hipStreamWaitEvent_spt_ptr)(hipStream_t, hipEvent_t, + unsigned int) = + reinterpret_cast + (hipStreamWaitEvent_spt_ptr); + + int N = 40; + int Nbytes = N * sizeof(int); + + int* hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, 10); + + int* devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, Nbytes)); + REQUIRE(devMem != nullptr); + + // Validating hipEventRecord_spt API with null stream + { + hipEvent_t start = nullptr, stop = nullptr; + + HIP_CHECK(hipEventCreate(&start)); + REQUIRE(start != nullptr); + + HIP_CHECK(hipEventCreate(&stop)); + REQUIRE(stop != nullptr); + + HIP_CHECK(dyn_hipEventRecord_spt_ptr(start, NULL)); + + HIP_CHECK(hipMemcpy(devMem, hostMem, Nbytes, hipMemcpyHostToDevice)); + addOneKernel<<< 1, 1 >>>(devMem, N); + HIP_CHECK(hipMemcpy(hostMem, devMem, Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(dyn_hipEventRecord_spt_ptr(stop, NULL)); + HIP_CHECK(hipEventSynchronize(stop)); + + REQUIRE(validateHostArray(hostMem, N, 11) == true); + + float time = 0.0f; + HIP_CHECK(hipEventElapsedTime(&time, start, stop)); + REQUIRE(time > 0.0f); + + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(stop)); + } + + // Validating hipEventRecord_spt API with user defined stream + { + hipEvent_t start = nullptr, stop = nullptr; + + HIP_CHECK(hipEventCreate(&start)); + REQUIRE(start != nullptr); + + HIP_CHECK(hipEventCreate(&stop)); + REQUIRE(stop != nullptr); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(dyn_hipEventRecord_spt_ptr(start, stream)); + + HIP_CHECK(hipMemcpyAsync(devMem, hostMem, Nbytes, + hipMemcpyHostToDevice, stream)); + addOneKernel<<< 1, 1, 0, stream >>>(devMem, N); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, Nbytes, + hipMemcpyDeviceToHost, stream)); + + HIP_CHECK(dyn_hipEventRecord_spt_ptr(stop, stream)); + HIP_CHECK(hipEventSynchronize(stop)); + + REQUIRE(validateHostArray(hostMem, N, 12) == true); + + float time = 0.0f; + HIP_CHECK(hipEventElapsedTime(&time, start, stop)); + REQUIRE(time > 0.0f); + + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(stop)); + HIP_CHECK(hipStreamDestroy(stream)); + } + + // Validating hipStreamWaitEvent_spt API + { + hipEvent_t waitEvent = nullptr; + + HIP_CHECK(hipEventCreate(&waitEvent)); + REQUIRE(waitEvent != nullptr); + + hipStream_t stream1, stream2; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + + HIP_CHECK(hipMemcpyAsync(devMem, hostMem, Nbytes, + hipMemcpyHostToDevice, stream1)); + + HIP_CHECK(hipEventRecord(waitEvent, stream1)); + HIP_CHECK(dyn_hipStreamWaitEvent_spt_ptr(stream2, waitEvent, 0)); + + addOneKernel<<< 1, 1, 0, stream2 >>>(devMem , N); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, Nbytes, + hipMemcpyDeviceToHost, stream2)); + + HIP_CHECK(hipStreamSynchronize(stream2)); + + REQUIRE(validateHostArray(hostMem, N, 13) == true); + + HIP_CHECK(hipEventDestroy(waitEvent)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + } + free(hostMem); + HIP_CHECK(hipFree(devMem)); +} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt b/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt index 975308b5c4..3fd7375cfd 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt @@ -48,9 +48,10 @@ if(HIP_PLATFORM MATCHES "amd") hipMemMap.cc hipMemRelease.cc hipMemUnmap.cc - hipMemVmm_old.cc) + hipMemVmm_old.cc + hipGetProcAddressVmmApis.cc) endif() hip_add_exe_to_target(NAME VirtualMemoryManagementTest TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) \ No newline at end of file + TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipGetProcAddressVmmApis.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipGetProcAddressVmmApis.cc new file mode 100644 index 0000000000..04d6011c0e --- /dev/null +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipGetProcAddressVmmApis.cc @@ -0,0 +1,248 @@ +/* +Copyright (c) 2024 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. +*/ + +#include +#include +#include +#include "hip_vmm_common.hh" +#include "../device/hipGetProcAddressHelpers.hh" + +/** + * Test Description + * ------------------------ + * - This test will get the function pointer of different + * - Virtual Memory Management APIs from the hipGetProcAddress API + * - and then validates the basic functionality of that particular API + * - using the funtion pointer. + * Test source + * ------------------------ + * - unit/virtualMemoryManagement/hipGetProcAddress_VMMAPIs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipGetProcAddress_VMM") { + int value = 0; + HIP_CHECK(hipDeviceGetAttribute(&value, + hipDeviceAttributeVirtualMemoryManagementSupported, 0)); + if (value == 0) { + HipTest::HIP_SKIP_TEST("Machine does not support VMM. Skipping Test.."); + return; + } + + void* hipMemGetAllocationGranularity_ptr = nullptr; + void* hipMemCreate_ptr = nullptr; + void* hipMemAddressReserve_ptr = nullptr; + void* hipMemMap_ptr = nullptr; + void* hipMemSetAccess_ptr = nullptr; + void* hipMemUnmap_ptr = nullptr; + void* hipMemAddressFree_ptr = nullptr; + void* hipMemRelease_ptr = nullptr; + void* hipMemGetAccess_ptr = nullptr; + void* hipMemGetAllocationPropertiesFromHandle_ptr = nullptr; + void* hipMemRetainAllocationHandle_ptr = nullptr; + + int currentHipVersion = 0; + HIP_CHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIP_CHECK(hipGetProcAddress("hipMemGetAllocationGranularity", + &hipMemGetAllocationGranularity_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemCreate", + &hipMemCreate_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemAddressReserve", + &hipMemAddressReserve_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemMap", + &hipMemMap_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemSetAccess", + &hipMemSetAccess_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemUnmap", + &hipMemUnmap_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemAddressFree", + &hipMemAddressFree_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemRelease", + &hipMemRelease_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemGetAccess", + &hipMemGetAccess_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemGetAllocationPropertiesFromHandle", + &hipMemGetAllocationPropertiesFromHandle_ptr, + currentHipVersion, 0, nullptr)); + HIP_CHECK(hipGetProcAddress("hipMemRetainAllocationHandle", + &hipMemRetainAllocationHandle_ptr, + currentHipVersion, 0, nullptr)); + + hipError_t(*dyn_hipMemGetAllocationGranularity_ptr)( + size_t *, const hipMemAllocationProp *, + hipMemAllocationGranularity_flags) = + reinterpret_cast + (hipMemGetAllocationGranularity_ptr); + + hipError_t(*dyn_hipMemCreate_ptr)( + hipMemGenericAllocationHandle_t *, + size_t, const hipMemAllocationProp *, + uint64_t) = + reinterpret_cast + (hipMemCreate_ptr); + + hipError_t(*dyn_hipMemAddressReserve_ptr)( + void **, size_t, size_t, void *, uint64_t) = + reinterpret_cast + (hipMemAddressReserve_ptr); + + hipError_t(*dyn_hipMemMap_ptr)( + void *, size_t, size_t, + hipMemGenericAllocationHandle_t, uint64_t) = + reinterpret_cast + (hipMemMap_ptr); + + hipError_t(*dyn_hipMemSetAccess_ptr)( + void *, size_t, const hipMemAccessDesc *, size_t) = + reinterpret_cast + (hipMemSetAccess_ptr); + + hipError_t(*dyn_hipMemUnmap_ptr)( + void *, size_t) = + reinterpret_cast + (hipMemUnmap_ptr); + + hipError_t(*dyn_hipMemAddressFree_ptr)( + void *, size_t) = + reinterpret_cast + (hipMemAddressFree_ptr); + + hipError_t(*dyn_hipMemRelease_ptr)(hipMemGenericAllocationHandle_t) = + reinterpret_cast + (hipMemRelease_ptr); + + hipError_t(*dyn_hipMemGetAccess_ptr)( + uint64_t *, const hipMemLocation *, void *) = + reinterpret_cast + (hipMemGetAccess_ptr); + + hipError_t(*dyn_hipMemGetAllocationPropertiesFromHandle_ptr)( + hipMemAllocationProp *, hipMemGenericAllocationHandle_t) = + reinterpret_cast + (hipMemGetAllocationPropertiesFromHandle_ptr); + + hipError_t(*dyn_hipMemRetainAllocationHandle_ptr)( + hipMemGenericAllocationHandle_t *, void *) = + reinterpret_cast + (hipMemRetainAllocationHandle_ptr); + + const int N = 10; + const int Nbytes = 10 * sizeof(int); + int *hostMem = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hostMem != nullptr); + fillHostArray(hostMem, N, 10); + + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.requestedHandleType = hipMemHandleTypeNone; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = 0; + + // Validating hipMemGetAllocationGranularity API + size_t granularity = 0; + size_t granularityWithFuncPtr = 0; + + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, + &prop, hipMemAllocationGranularityMinimum)); + HIP_CHECK(dyn_hipMemGetAllocationGranularity_ptr(&granularityWithFuncPtr, + &prop, hipMemAllocationGranularityMinimum)); + + REQUIRE(granularity > 0); + REQUIRE(granularityWithFuncPtr > 0); + REQUIRE(granularityWithFuncPtr == granularity); + + // Validating hipMemCreate API + hipMemGenericAllocationHandle_t handle = nullptr; + HIP_CHECK(dyn_hipMemCreate_ptr(&handle, granularity, &prop, 0)); + REQUIRE(handle != nullptr); + + // Validating hipMemAddressReserve API + void* ptr = nullptr; + HIP_CHECK(dyn_hipMemAddressReserve_ptr(&ptr, granularity, 0, 0, 0)); + REQUIRE(ptr != nullptr); + + // Validating hipMemMap API + HIP_CHECK(dyn_hipMemMap_ptr(ptr, granularity, 0, handle, 0)); + REQUIRE(ptr != nullptr); + + // Validating hipMemSetAccess API + hipMemAccessDesc desc; + desc.location.type = hipMemLocationTypeDevice; + desc.location.id = 0; + desc.flags = hipMemAccessFlagsProtReadWrite; + HIP_CHECK(dyn_hipMemSetAccess_ptr(ptr, granularity, &desc, 1)); + REQUIRE(ptr != nullptr); + + // Performing some operations on ptr, to validate it + HIP_CHECK(hipMemcpy(ptr, hostMem, Nbytes, hipMemcpyHostToDevice)); + addOneKernel<<< 1, 1 >>>(reinterpret_cast(ptr), N); + HIP_CHECK(hipMemcpy(hostMem, ptr, Nbytes, hipMemcpyDeviceToHost)); + validateHostArray(hostMem, N, 11); + + // Validating hipMemGetAccess API + uint64_t flags; + hipMemLocation location = {hipMemLocationTypeDevice, 0}; + HIP_CHECK(dyn_hipMemGetAccess_ptr(&flags, &location, ptr)); + REQUIRE(flags == hipMemAccessFlagsProtReadWrite); + + // Validating hipMemGetAllocationPropertiesFromHandle API + hipMemAllocationProp requiredProp; + HIP_CHECK(dyn_hipMemGetAllocationPropertiesFromHandle_ptr( + &requiredProp, handle)); + REQUIRE(requiredProp.type == hipMemAllocationTypePinned); + REQUIRE(requiredProp.requestedHandleType == hipMemHandleTypeNone); + REQUIRE(requiredProp.location.type == hipMemLocationTypeDevice); + REQUIRE(requiredProp.location.id == 0); + + // Validating hipMemRetainAllocationHandle API + hipMemGenericAllocationHandle_t requiredHandle = nullptr; + HIP_CHECK(dyn_hipMemRetainAllocationHandle_ptr(&requiredHandle, ptr)); + REQUIRE(requiredHandle == handle); + + // Validating hipMemUnmap, hipMemAddressFree, hipMemRelease API's + HIP_CHECK(dyn_hipMemUnmap_ptr(ptr, granularity)); + HIP_CHECK(dyn_hipMemAddressFree_ptr(ptr, granularity)); + HIP_CHECK(dyn_hipMemRelease_ptr(handle)); + + // Performing operation on ptr, to check it is invalidated or not + REQUIRE(hipMemcpy(ptr, hostMem, Nbytes, hipMemcpyHostToDevice) + == hipErrorInvalidValue); + free(hostMem); +}