From 942fc5de3f62e0feca06fca2999fda86cbb8f07c Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 12 Jun 2017 11:53:25 +0530 Subject: [PATCH 1/5] Update P2P test for using memcpy and kernel tests Change-Id: Ib0f8fc9425e6e85fd11d7d02395c52bc713dcb37 [ROCm/clr commit: ce6e45567790704c2a3918155371bdeee2908425] --- .../runtimeApi/memory/p2p_copy_coherency.cpp | 90 +++++++++++++------ 1 file changed, 61 insertions(+), 29 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp index 6bc6235454..a5d79464d0 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp @@ -33,10 +33,14 @@ THE SOFTWARE. #ifdef __HIP_PLATFORM_HCC__ #include +#define USE_HCC_MEMTRACKER 0 #endif #define USE_HSA_COPY 1 +int elementSizes[] = {16, 1024,524288}; +int nSizes = sizeof(elementSizes) / sizeof(int); + int enablePeers(int dev0, int dev1) { int canAccessPeer01, canAccessPeer10; @@ -54,16 +58,25 @@ int enablePeers(int dev0, int dev1) return 0; }; - __global__ void -memsetIntKernel(int * ptr, int val, size_t numElements) +memsetIntKernel(/*hipLaunchParm lp,*/ int * ptr, const int val, size_t numElements) { int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); - if (gid < numElements) { - ptr[gid] = val; + int stride = hipBlockDim_x * hipGridDim_x ; + for (size_t i= gid; i< numElements; i+=stride){ + ptr[i] = val; } }; +__global__ void +memcpyIntKernel(/*hipLaunchParm lp, */const int * src, int* dst, size_t numElements) +{ + int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); + int stride = hipBlockDim_x * hipGridDim_x ; + for (size_t i= gid; i< numElements; i+=stride){ + dst[i] = src[i]; + } +}; void checkReverse(const int *ptr, int numElements, int expected) { for (int i=numElements-1; i>=0; i--) { @@ -76,52 +89,66 @@ void checkReverse(const int *ptr, int numElements, int expected) { printf ("test: OK\n"); } - -void runTest(bool stepAIsCopy, hipStream_t gpu0Stream, hipStream_t gpu1Stream, int numElements, - int * dataGpu0, int *dataGpu1, int *dataHost, int expected) +void runTest(bool stepAIsCopy, bool hostSync, hipStream_t gpu0Stream, hipStream_t gpu1Stream, int numElements, + int * dataGpu0_0, int * dataGpu0_1, int *dataGpu1, int *dataHost, int expected) { hipEvent_t e; - HIPCHECK(hipEventCreateWithFlags(&e,0)); - - printf ("test: runTest with %s\n", stepAIsCopy ? "copy" : "kernel"); + if(!hostSync) { + HIPCHECK(hipEventCreateWithFlags(&e,0)); + } const size_t sizeElements = numElements * sizeof(int); + printf ("test: runTest with %zu bytes %s with hostSync %s\n", sizeElements, stepAIsCopy ? "copy" : "kernel", hostSync ? "enabled" : "disabled"); hipStream_t stepAStream = gpu0Stream; if (stepAIsCopy) { #ifdef USE_HSA_COPY - HIPCHECK(hipMemcpyAsync(dataGpu1, dataGpu0, sizeElements, hipMemcpyDeviceToDevice, stepAStream)); + HIPCHECK(hipMemcpyAsync(dataGpu1, dataGpu0_0, sizeElements, hipMemcpyDeviceToDevice, stepAStream)); #endif } else { - assert(0); // not yet supported. + //assert(0); // not yet supported. + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements); + hipLaunchKernelGGL(memcpyIntKernel, dim3(blocks), dim3(threadsPerBlock), 0, gpu0Stream, + dataGpu0_0, dataGpu1, numElements); } - HIPCHECK(hipEventRecord(e, stepAStream)); - HIPCHECK(hipStreamWaitEvent(gpu1Stream, e, 0)); + if(!hostSync) { + HIPCHECK(hipEventRecord(e, stepAStream)); + HIPCHECK(hipStreamWaitEvent(gpu1Stream, e, 0)); + } else { + HIPCHECK(hipStreamSynchronize(stepAStream)); + } - HIPCHECK(hipMemcpyAsync(dataHost, dataGpu1, sizeElements, hipMemcpyDeviceToHost, gpu1Stream)); + HIPCHECK(hipMemcpyAsync(dataGpu0_1, dataGpu1, sizeElements, hipMemcpyDeviceToDevice, gpu1Stream)); - HIPCHECK(hipStreamSynchronize(gpu1Stream)); + if(!hostSync) { + HIPCHECK(hipEventRecord(e, gpu1Stream)); + } else { + HIPCHECK(hipStreamSynchronize(gpu1Stream)); + } + + HIPCHECK(hipMemcpyAsync(dataHost, dataGpu0_1, sizeElements, hipMemcpyDeviceToHost, gpu0Stream)); + HIPCHECK(hipStreamSynchronize(gpu0Stream)); checkReverse(dataHost, numElements, expected); } - -void testMultiGpu0(int dev0, int dev1, int numElements) +void testMultiGpu(int dev0, int dev1, int numElements, bool hostSync, bool useMemcpy) { const size_t sizeElements = numElements * sizeof(int); - int * dataGpu0, *dataGpu1, *dataHost; + int * dataGpu0_0, * dataGpu0_1, *dataGpu1, *dataHost; hipStream_t gpu0Stream, gpu1Stream; const int expected = 42; unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements); HIPCHECK(hipSetDevice(dev0)); - HIPCHECK(hipMalloc(&dataGpu0, sizeElements)); + HIPCHECK(hipMalloc(&dataGpu0_0, sizeElements)); + HIPCHECK(hipMalloc(&dataGpu0_1, sizeElements)); HIPCHECK(hipStreamCreate(&gpu0Stream)); hipLaunchKernelGGL(memsetIntKernel, dim3(blocks), dim3(threadsPerBlock), 0, gpu0Stream, - dataGpu0, expected, numElements); + dataGpu0_0, expected, numElements); HIPCHECK(hipDeviceSynchronize()); @@ -135,18 +162,19 @@ void testMultiGpu0(int dev0, int dev1, int numElements) HIPCHECK(hipHostMalloc(&dataHost, sizeElements)); memset(dataHost, 13, sizeElements); -#ifdef __HIP_PLATFORM_HCC__ +#if USE_HCC_MEMTRACKER hc::am_memtracker_print(0x0); #endif - + printf (" test: init complete\n"); + runTest(useMemcpy , hostSync, gpu0Stream, gpu1Stream, numElements, dataGpu0_0,dataGpu0_1, dataGpu1, dataHost, expected); - runTest(true/*stepAIsCopy*/, gpu0Stream, gpu1Stream, numElements, dataGpu0, dataGpu1, dataHost, expected); - + HIPCHECK(hipFree(dataGpu0_0)); + HIPCHECK(hipFree(dataGpu0_1)); + HIPCHECK(hipFree(dataGpu1)); + HIPCHECK(hipHostFree(dataHost)); }; - - int main(int argc, char *argv[]) { HipTest::parseStandardArguments(argc, argv, true); @@ -168,8 +196,12 @@ int main(int argc, char *argv[]) return -1; }; - //testMultiGpu0(dev0, dev1, numElements); - + for(int index = 1;index < nSizes;index++) { + testMultiGpu(dev0, dev1, elementSizes[index] , false /* GPU Synchronization*/, true); + testMultiGpu(dev0, dev1, elementSizes[index] , true /*Host Synchronization*/, true); + testMultiGpu(dev0, dev1, elementSizes[index] , true /*Host Synchronization*/, false); + testMultiGpu(dev0, dev1, elementSizes[index] , false /*Host Synchronization*/, false); + } passed(); From f48cbf564a0bc34dff00969a11534746f4e109c1 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 13 Jun 2017 13:35:50 +0530 Subject: [PATCH 2/5] Input args NULL check in hipChooseDevice Change-Id: I1a7b8cded2f81d739645bbf3dab2f04bb9c3c796 [ROCm/clr commit: 0efd737767a56cf7d28b6a0cfbe2774cbe934e08] --- projects/clr/hipamd/src/hip_device.cpp | 110 +++++++++++++------------ 1 file changed, 58 insertions(+), 52 deletions(-) diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp index 93c1c20484..05db4c2b30 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/projects/clr/hipamd/src/hip_device.cpp @@ -415,72 +415,78 @@ hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop ) int inPropCount = 0; int matchedPropCount = 0; hipError_t e = hipSuccess; - ihipGetDeviceCount( &deviceCount ); - *device = 0; - for (int i = 0; i < deviceCount; i++) { - ihipGetDeviceProperties( &tempProp, i ); - if(prop->major != 0) { - inPropCount++; - if(tempProp.major >= prop->major) { - matchedPropCount++; - } - if(prop->minor != 0) { + if((device == NULL) || (prop == NULL)) { + e = hipErrorInvalidValue; + } + if(e == hipSuccess) { + ihipGetDeviceCount( &deviceCount ); + *device = 0; + for (int i = 0; i < deviceCount; i++) { + ihipGetDeviceProperties( &tempProp, i ); + if(prop->major != 0) { inPropCount++; - if(tempProp.minor >= prop->minor) { - matchedPropCount++; - } + if(tempProp.major >= prop->major) { + matchedPropCount++; + } + if(prop->minor != 0) { + inPropCount++; + if(tempProp.minor >= prop->minor) { + matchedPropCount++; + } + } } - } - if(prop->totalGlobalMem != 0) { - inPropCount++; - if(tempProp.totalGlobalMem >= prop->totalGlobalMem) { - matchedPropCount++; + if(prop->totalGlobalMem != 0) { + inPropCount++; + if(tempProp.totalGlobalMem >= prop->totalGlobalMem) { + matchedPropCount++; + } } - } - if(prop->sharedMemPerBlock != 0) { - inPropCount++; - if(tempProp.sharedMemPerBlock >= prop->sharedMemPerBlock) { - matchedPropCount++; + if(prop->sharedMemPerBlock != 0) { + inPropCount++; + if(tempProp.sharedMemPerBlock >= prop->sharedMemPerBlock) { + matchedPropCount++; + } } - } - if(prop->maxThreadsPerBlock != 0) { - inPropCount++; - if(tempProp.maxThreadsPerBlock >= prop->maxThreadsPerBlock ) { - matchedPropCount++; + if(prop->maxThreadsPerBlock != 0) { + inPropCount++; + if(tempProp.maxThreadsPerBlock >= prop->maxThreadsPerBlock ) { + matchedPropCount++; + } } - } - if(prop->totalConstMem != 0) { - inPropCount++; - if(tempProp.totalConstMem >= prop->totalConstMem ) { - matchedPropCount++; + if(prop->totalConstMem != 0) { + inPropCount++; + if(tempProp.totalConstMem >= prop->totalConstMem ) { + matchedPropCount++; + } } - } - if(prop->multiProcessorCount != 0) { - inPropCount++; - if(tempProp.multiProcessorCount >= prop->multiProcessorCount ) { - matchedPropCount++; + if(prop->multiProcessorCount != 0) { + inPropCount++; + if(tempProp.multiProcessorCount >= prop->multiProcessorCount ) { + matchedPropCount++; + } } - } - if(prop->maxThreadsPerMultiProcessor != 0) { - inPropCount++; - if(tempProp.maxThreadsPerMultiProcessor >= prop->maxThreadsPerMultiProcessor ) { - matchedPropCount++; + if(prop->maxThreadsPerMultiProcessor != 0) { + inPropCount++; + if(tempProp.maxThreadsPerMultiProcessor >= prop->maxThreadsPerMultiProcessor ) { + matchedPropCount++; + } } - } - if(prop->memoryClockRate != 0) { - inPropCount++; - if(tempProp.memoryClockRate >= prop->memoryClockRate ) { - matchedPropCount++; + if(prop->memoryClockRate != 0) { + inPropCount++; + if(tempProp.memoryClockRate >= prop->memoryClockRate ) { + matchedPropCount++; + } + } + if(inPropCount == matchedPropCount) { + *device = i; } - } - if(inPropCount == matchedPropCount) { - *device = i; - } #if 0 else{ e= hipErrorInvalidValue; } #endif + } } return ihipLogStatus(e); } + From cfe2dc9fc090c2ed071b3be5c6aec39f47356c7a Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 13 Jun 2017 20:25:11 +0300 Subject: [PATCH 3/5] [HIPIFY] Initial sync HIPIFY with HIP by CUDA Driver API functions. + CUDA_Driver_API_functions_supported_by_HIP.md update. + Initial update of HIPIFY with CUDA driver API functions: 1.Error Handling, 2.Initialization, 3.Version Management, 5-6.Device Management, 7.Primary Context, 8-9.Context, 10.Module Management, 11.Memory Management. + Sync HIP functions against CUDA Driver and Runtime API functions. + Typo fixes. ToDo: 12-30 modules of CUDA Driver API. [ROCm/clr commit: 82b37fe48116aa6f122b0d2a2fc087fcb25f9a32] --- ...A_Driver_API_functions_supported_by_HIP.md | 123 +++++++++++-- .../clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 169 +++++++++++------- 2 files changed, 221 insertions(+), 71 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index ad9d791a6d..d4b54438bb 100644 --- a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -386,65 +386,168 @@ | define |`CUDA_ARRAY3D_TEXTURE_GATHER` | | This flag must be set in order to perform texture gather operations on a CUDA array. | | define |`CUDA_VERSION` | | CUDA API version number. | - ## **2. Error Handling** | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| - +| `cuGetErrorName` | | Gets the string representation of an error code enum name. | +| `cuGetErrorString` | | Gets the string description of an error code. | ## **3. Initialization** | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| - +| `cuInit` | `hipInit` | Initialize the CUDA driver API. | ## **4. Version Management** | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| - +| `cuDriverGetVersion` | `hipDriverGetVersion` | Returns the CUDA driver version. | ## **5. Device Management** | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cuDriverGetVersion` | `hipGetDevice` | Returns a handle to a compute device. | +| `cuDeviceGetAttribute` | `hipDeviceGetAttribute` | Returns information about the device. | +| `cuDeviceGetCount` | `hipGetDeviceCount` | Returns the number of compute-capable devices. | +| `cuDeviceGetName` | `hipDeviceGetName` | Returns an identifer string for the device. | +| `cuDeviceTotalMem` | `hipDeviceTotalMem` | Returns the total amount of memory on the device. | ## **6. Device Management [DEPRECATED]** | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| - +| `cuDeviceComputeCapability` | `hipDeviceComputeCapability` | Returns the compute capability of the device. | +| `cuDeviceGetProperties` | `hipGetDeviceProperties` | Returns properties for a selected device. | ## **7. Primary Context Management** | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| - +| `cuDevicePrimaryCtxGetState` | `hipDevicePrimaryCtxGetState` | Get the state of the primary context. | +| `cuDevicePrimaryCtxRelease` | `hipDevicePrimaryCtxRelease` | Release the primary context on the GPU. | +| `cuDevicePrimaryCtxReset` | `hipDevicePrimaryCtxReset` | Destroy all allocations and reset all state on the primary context. | +| `cuDevicePrimaryCtxRetain` | `hipDevicePrimaryCtxRetain` | Retain the primary context on the GPU. | +| `cuDevicePrimaryCtxSetFlags` | `hipDevicePrimaryCtxSetFlags` | Set flags for the primary context. | ## **8. Context Management** | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| - +| `cuCtxCreate` | `hipCtxCreate` | Create a CUDA context. | +| `cuCtxDestroy` | `hipCtxDestroy` | Destroy a CUDA context. | +| `cuCtxGetApiVersion` | `hipCtxGetApiVersion` | Gets the context's API version. | +| `cuCtxGetCacheConfig` | `hipCtxGetCacheConfig` | Returns the preferred cache configuration for the current context. | +| `cuCtxGetCurrent` | `hipCtxGetCurrent` | Returns the CUDA context bound to the calling CPU thread. | +| `cuCtxGetDevice` | `hipCtxGetDevice` | Returns the device ID for the current context. | +| `cuCtxGetFlags` | `hipCtxGetFlags` | Returns the flags for the current context. | +| `cuCtxGetLimit` | | Returns resource limits. | +| `cuCtxGetSharedMemConfig` | `hipCtxGetSharedMemConfig` | Returns the current shared memory configuration for the current context. | +| `cuCtxGetStreamPriorityRange` | | Returns numerical values that correspond to the least and greatest stream priorities. | +| `cuCtxPopCurrent` | `hipCtxPopCurrent` | Pops the current CUDA context from the current CPU thread. | +| `cuCtxPushCurrent` | `hipCtxPushCurrent` | Pushes a context on the current CPU thread. | +| `cuCtxSetCacheConfig` | `hipCtxSetCacheConfig` | Sets the preferred cache configuration for the current context. | +| `cuCtxSetCurrent` | `hipCtxSetCurrent` | Binds the specified CUDA context to the calling CPU thread. | +| `cuCtxSetLimit` | | Set resource limits. | +| `cuCtxSetSharedMemConfig` | `hipCtxSetSharedMemConfig` | Sets the shared memory configuration for the current context. | +| `cuCtxSynchronize` | `hipCtxSynchronize` | Block for a context's tasks to complete. | ## **9. Context Management [DEPRECATED]** | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| - +| `cuCtxAttach` | | Increment a context's usage-count. | +| `cuCtxDetach` | | Decrement a context's usage-count. | ## **10. Module Management** | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| - +| `cuLinkAddData` | | Add an input to a pending linker invocation. | +| `cuLinkAddFile` | | Add a file input to a pending linker invocation. | +| `cuLinkComplete` | | Complete a pending linker invocation. | +| `cuLinkCreate` | | Creates a pending JIT linker invocation. | +| `cuLinkDestroy` | | Destroys state for a JIT linker invocation. | +| `cuModuleGetFunction` | `hipModuleGetFunction` | Returns a function handle. | +| `cuModuleGetGlobal` | `hipModuleGetGlobal` | Returns a global pointer from a module. | +| `cuModuleGetSurfRef` | | Returns a handle to a surface reference. | +| `cuModuleGetTexRef` | | Returns a handle to a texture reference. | +| `cuModuleLoad` | `hipModuleLoad` | Loads a compute module. | +| `cuModuleLoadData` | `hipModuleLoadData` | Load a module's data. | +| `cuModuleLoadDataEx` | `hipModuleLoadDataEx` | Load a module's data with options. | +| `cuModuleLoadFatBinary` | | Load a module's data. | +| `cuModuleUnload` | `hipModuleUnload` | Unloads a module. | ## **11. Memory Management** | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| - +| `cuArray3DCreate` | | Creates a 3D CUDA array. | +| `cuArray3DGetDescriptor` | | Get a 3D CUDA array descriptor. | +| `cuArrayCreate` | | Creates a 1D or 2D CUDA array. | +| `cuArrayDestroy` | | Destroys a CUDA array. | +| `cuArrayGetDescriptor` | | Get a 1D or 2D CUDA array descriptor. | +| `cuDeviceGetByPCIBusId` | `hipDeviceGetByPCIBusId` | Returns a handle to a compute device. | +| `cuDeviceGetPCIBusId` | `hipDeviceGetPCIBusId` | Returns a PCI Bus Id string for the device. | +| `cuIpcCloseMemHandle` | | Close memory mapped with cuIpcOpenMemHandle. | +| `cuIpcGetEventHandle` | | Gets an interprocess handle for a previously allocated event. | +| `cuIpcGetMemHandle` | | Gets an interprocess memory handle for an existing device memory allocation. | +| `cuIpcOpenEventHandle` | | Opens an interprocess event handle for use in the current process. | +| `cuIpcOpenMemHandle` | | Opens an interprocess memory handle exported from another process and returns a device pointer usable in the local process. | +| `cuMemAlloc` | `hipMalloc` | Allocates device memory. | +| `cuMemAllocHost` | | Allocates page-locked host memory. | +| `cuMemAllocManaged` | | Allocates memory that will be automatically managed by the Unified Memory system. | +| `cuMemAllocPitch` | | Allocates pitched device memory. | +| `cuMemcpy` | | Copies memory. | +| `cuMemcpy2D` | | Copies memory for 2D arrays. | +| `cuMemcpy2DAsync` | | Copies memory for 2D arrays. | +| `cuMemcpy2DUnaligned` | | Copies memory for 2D arrays. | +| `cuMemcpy3D` | | Copies memory for 3D arrays. | +| `cuMemcpy3DAsync` | | Copies memory for 3D arrays. | +| `cuMemcpy3DPeer` | | Copies memory between contexts. | +| `cuMemcpy3DPeerAsync` | | Copies memory between contexts asynchronously. | +| `cuMemcpyAsync` | | Copies memory asynchronously. | +| `cuMemcpyAtoA` | | Copies memory from Array to Array. | +| `cuMemcpyAtoD` | | Copies memory from Array to Device. | +| `cuMemcpyAtoH` | | Copies memory from Array to Host. | +| `cuMemcpyAtoHAsync` | | Copies memory from Array to Host. | +| `cuMemcpyDtoA` | | Copies memory from Device to Array. | +| `cuMemcpyDtoD` | `hipMemcpyDtoD` | Copies memory from Device to Device. | +| `cuMemcpyDtoDAsync` | `hipMemcpyDtoDAsync` | Copies memory from Device to Device. | +| `cuMemcpyDtoH` | `hipMemcpyDtoH` | Copies memory from Device to Host. | +| `cuMemcpyDtoHAsync` | `hipMemcpyDtoHAsync` | Copies memory from Device to Host. | +| `cuMemcpyHtoA` | | Copies memory from Host to Array. | +| `cuMemcpyHtoAAsync` | | Copies memory from Host to Array. | +| `cuMemcpyHtoD` | `hipMemcpyHtoD` | Copies memory from Host to Device. | +| `cuMemcpyHtoDAsync` | `hipMemcpyHtoDAsync` | Copies memory from Host to Device. | +| `cuMemcpyPeer` | | Copies device memory between two contexts. | +| `cuMemcpyPeerAsync` | | Copies device memory between two contexts asynchronously. | +| `cuMemFree` | `hipFree` | Frees device memory. | +| `cuMemFreeHost` | `hipFreeHost` | Frees page-locked host memory. | +| `cuMemGetAddressRange` | | Get information on memory allocations. | +| `cuMemGetInfo` | `hipMemGetInfo` | Gets free and total memory. | +| `cuMemHostAlloc` | `hipHostMalloc` | Allocates page-locked host memory. | +| `cuMemHostGetDevicePointer` | | Passes back device pointer of mapped pinned memory. | +| `cuMemHostGetFlags` | | Passes back flags that were used for a pinned allocation. | +| `cuMemHostRegister` | `hipHostRegister` | Registers an existing host memory range for use by CUDA. | +| `cuMemHostUnregister` | `hipHostUnregister` | Unregisters a memory range that was registered with cuMemHostRegister. | +| `cuMemsetD16` | | Initializes device memory. | +| `cuMemsetD16Async` | | Sets device memory. | +| `cuMemsetD2D16` | | Initializes device memory. | +| `cuMemsetD2D16Async` | | Sets device memory. | +| `cuMemsetD2D32` | | Initializes device memory. | +| `cuMemsetD2D32Async` | | Sets device memory. | +| `cuMemsetD2D8` | | Initializes device memory. | +| `cuMemsetD2D8Async` | | Sets device memory. | +| `cuMemsetD32` | `hipMemset` | Initializes device memory. | +| `cuMemsetD32Async` | `hipMemsetAsync` | Sets device memory. | +| `cuMemsetD2D8` | | Initializes device memory. | +| `cuMemsetD2D8Async` | | Sets device memory. | +| `cuMipmappedArrayCreate` | | Creates a CUDA mipmapped array. | +| `cuMipmappedArrayDestroy` | | Destroys a CUDA mipmapped array. | +| `cuMipmappedArrayGetLevel` | | Gets a mipmap level of a CUDA mipmapped array. | ## **12. Unified Addressing** diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index 0825285b51..de4da78451 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -391,7 +391,7 @@ struct cuda2hipMap { cuda2hipRename["cudaErrorUnknown"] = {"hipErrorUnknown", CONV_ERR, API_RUNTIME}; // 30 ///////////////////////////// CUDA DRIVER API ///////////////////////////// - // enums + // structs cuda2hipRename["CUDA_ARRAY3D_DESCRIPTOR"] = {"HIP_ARRAY3D_DESCRIPTOR", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; cuda2hipRename["CUDA_ARRAY_DESCRIPTOR"] = {"HIP_ARRAY_DESCRIPTOR", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; cuda2hipRename["CUDA_MEMCPY2D"] = {"HIP_MEMCPY2D", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; @@ -453,16 +453,16 @@ struct cuda2hipMap { cuda2hipRename["CU_MEM_RANGE_ATTRIBUTE_LAST_PREFETCH_LOCATION"] = {"hipMemRangeAttributeLastPrefetchLocation", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 4 // API_RUNTIME ANALOGUE (cudaMemRangeAttributeLastPrefetchLocation = 4) // Context flags - cuda2hipRename["CUctx_flags"] = {"hipCctx_flags", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; - cuda2hipRename["CU_CTX_SCHED_AUTO"] = {"HIP_CTX_SCHED_AUTO", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x00 - cuda2hipRename["CU_CTX_SCHED_SPIN"] = {"HIP_CTX_SCHED_SPIN", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x01 - cuda2hipRename["CU_CTX_SCHED_YIELD"] = {"HIP_CTX_SCHED_YIELD", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x02 - cuda2hipRename["CU_CTX_SCHED_BLOCKING_SYNC"] = {"HIP_CTX_SCHED_BLOCKING_SYNC", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x04 - cuda2hipRename["CU_CTX_BLOCKING_SYNC"] = {"HIP_CTX_BLOCKING_SYNC", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x04 - cuda2hipRename["CU_CTX_SCHED_MASK"] = {"HIP_CTX_SCHED_MASK", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x07 - cuda2hipRename["CU_CTX_MAP_HOST"] = {"HIP_CTX_MAP_HOST", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x08 - cuda2hipRename["CU_CTX_LMEM_RESIZE_TO_MAX"] = {"HIP_CTX_LMEM_RESIZE_TO_MAX", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x10 - cuda2hipRename["CU_CTX_FLAGS_MASK"] = {"HIP_CTX_FLAGS_MASK", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x1f + cuda2hipRename["CUctx_flags"] = {"hipCctx_flags", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["CU_CTX_SCHED_AUTO"] = {"HIP_CTX_SCHED_AUTO", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x00 + cuda2hipRename["CU_CTX_SCHED_SPIN"] = {"HIP_CTX_SCHED_SPIN", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x01 + cuda2hipRename["CU_CTX_SCHED_YIELD"] = {"HIP_CTX_SCHED_YIELD", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x02 + cuda2hipRename["CU_CTX_SCHED_BLOCKING_SYNC"] = {"HIP_CTX_SCHED_BLOCKING_SYNC", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x04 + cuda2hipRename["CU_CTX_BLOCKING_SYNC"] = {"HIP_CTX_BLOCKING_SYNC", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x04 + cuda2hipRename["CU_CTX_SCHED_MASK"] = {"HIP_CTX_SCHED_MASK", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x07 + cuda2hipRename["CU_CTX_MAP_HOST"] = {"HIP_CTX_MAP_HOST", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x08 + cuda2hipRename["CU_CTX_LMEM_RESIZE_TO_MAX"] = {"HIP_CTX_LMEM_RESIZE_TO_MAX", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x10 + cuda2hipRename["CU_CTX_FLAGS_MASK"] = {"HIP_CTX_FLAGS_MASK", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}; // 0x1f // Defines cuda2hipRename["CU_LAUNCH_PARAM_BUFFER_POINTER"] = {"HIP_LAUNCH_PARAM_BUFFER_POINTER", CONV_DEV, API_DRIVER}; // ((void*)0x01) @@ -882,52 +882,79 @@ struct cuda2hipMap { cuda2hipRename["CU_STREAM_MEM_OP_WRITE_VALUE_32"] = {"hipStreamBatchMemOpWriteValue32", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; // 2 cuda2hipRename["CU_STREAM_MEM_OP_FLUSH_REMOTE_WRITES"] = {"hipStreamBatchMemOpFlushRemoteWrites", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; // 3 + // Error Handling + cuda2hipRename["cuGetErrorName"] = {"hipGetErrorName___", CONV_ERR, API_DRIVER, HIP_UNSUPPORTED}; // cudaGetErrorName (hipGetErrorName) has different signature + cuda2hipRename["cuGetErrorString"] = {"hipGetErrorString___", CONV_ERR, API_DRIVER, HIP_UNSUPPORTED}; // cudaGetErrorString (hipGetErrorString) has different signature + // Init cuda2hipRename["cuInit"] = {"hipInit", CONV_DRIVER, API_DRIVER}; // Driver - cuda2hipRename["cuDriverGetVersion"] = {"hipDriverGetVersion", CONV_DRIVER, API_DRIVER}; + cuda2hipRename["cuDriverGetVersion"] = {"hipDriverGetVersion", CONV_DRIVER, API_DRIVER}; - // Context + // Context Management cuda2hipRename["cuCtxCreate_v2"] = {"hipCtxCreate", CONV_CONTEXT, API_DRIVER}; cuda2hipRename["cuCtxDestroy_v2"] = {"hipCtxDestroy", CONV_CONTEXT, API_DRIVER}; - cuda2hipRename["cuCtxPopCurrent_v2"] = {"hipCtxPopCurrent", CONV_CONTEXT, API_DRIVER}; - cuda2hipRename["cuCtxPushCurrent_v2"] = {"hipCtxPushCurrent", CONV_CONTEXT, API_DRIVER}; - cuda2hipRename["cuCtxSetCurrent"] = {"hipCtxSetCurrent", CONV_CONTEXT, API_DRIVER}; - cuda2hipRename["cuCtxGetCurrent"] = {"hipCtxGetCurrent", CONV_CONTEXT, API_DRIVER}; - cuda2hipRename["cuCtxGetDevice"] = {"hipCtxGetDevice", CONV_CONTEXT, API_DRIVER}; cuda2hipRename["cuCtxGetApiVersion"] = {"hipCtxGetApiVersion", CONV_CONTEXT, API_DRIVER}; cuda2hipRename["cuCtxGetCacheConfig"] = {"hipCtxGetCacheConfig", CONV_CONTEXT, API_DRIVER}; - cuda2hipRename["cuCtxSetCacheConfig"] = {"hipCtxSetCacheConfig", CONV_CONTEXT, API_DRIVER}; - cuda2hipRename["cuCtxSetSharedMemConfig"] = {"hipCtxSetSharedMemConfig", CONV_CONTEXT, API_DRIVER}; - cuda2hipRename["cuCtxGetSharedMemConfig"] = {"hipCtxGetSharedMemConfig", CONV_CONTEXT, API_DRIVER}; - cuda2hipRename["cuCtxSynchronize"] = {"hipCtxSynchronize", CONV_CONTEXT, API_DRIVER}; + cuda2hipRename["cuCtxGetCurrent"] = {"hipCtxGetCurrent", CONV_CONTEXT, API_DRIVER}; + cuda2hipRename["cuCtxGetDevice"] = {"hipCtxGetDevice", CONV_CONTEXT, API_DRIVER}; cuda2hipRename["cuCtxGetFlags"] = {"hipCtxGetFlags", CONV_CONTEXT, API_DRIVER}; + cuda2hipRename["cuCtxGetLimit"] = {"hipCtxGetLimit", CONV_CONTEXT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuCtxGetSharedMemConfig"] = {"hipCtxGetSharedMemConfig", CONV_CONTEXT, API_DRIVER}; + cuda2hipRename["cuCtxGetStreamPriorityRange"] = {"hipCtxGetStreamPriorityRange", CONV_CONTEXT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuCtxPopCurrent_v2"] = {"hipCtxPopCurrent", CONV_CONTEXT, API_DRIVER}; + cuda2hipRename["cuCtxPushCurrent_v2"] = {"hipCtxPushCurrent", CONV_CONTEXT, API_DRIVER}; + cuda2hipRename["cuCtxSetCacheConfig"] = {"hipCtxSetCacheConfig", CONV_CONTEXT, API_DRIVER}; + cuda2hipRename["cuCtxSetCurrent"] = {"hipCtxSetCurrent", CONV_CONTEXT, API_DRIVER}; + cuda2hipRename["cuCtxSetLimit"] = {"hipCtxSetLimit", CONV_CONTEXT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuCtxSetSharedMemConfig"] = {"hipCtxSetSharedMemConfig", CONV_CONTEXT, API_DRIVER}; + cuda2hipRename["cuCtxSynchronize"] = {"hipCtxSynchronize", CONV_CONTEXT, API_DRIVER}; + // Context Management [DEPRECATED] + cuda2hipRename["cuCtxAttach"] = {"hipCtxAttach", CONV_CONTEXT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuCtxDetach"] = {"hipCtxDetach", CONV_CONTEXT, API_DRIVER, HIP_UNSUPPORTED}; + + // Peer Context Memory Access cuda2hipRename["cuCtxEnablePeerAccess"] = {"hipCtxEnablePeerAccess", CONV_CONTEXT, API_DRIVER}; cuda2hipRename["cuCtxDisablePeerAccess"] = {"hipCtxDisablePeerAccess", CONV_CONTEXT, API_DRIVER}; - // unsupported yet by HIP - cuda2hipRename["cuCtxSetLimit"] = {"hipCtxSetLimit", CONV_CONTEXT, API_DRIVER, HIP_UNSUPPORTED}; - cuda2hipRename["cuCtxGetLimit"] = {"hipCtxGetLimit", CONV_CONTEXT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuDeviceCanAccessPeer"] = {"hipDeviceCanAccessPeer", CONV_DEV, API_DRIVER}; + // Primary Context Management cuda2hipRename["cuDevicePrimaryCtxGetState"] = {"hipDevicePrimaryCtxGetState", CONV_CONTEXT, API_DRIVER}; cuda2hipRename["cuDevicePrimaryCtxRelease"] = {"hipDevicePrimaryCtxRelease", CONV_CONTEXT, API_DRIVER}; - cuda2hipRename["cuDevicePrimaryCtxRetain"] = {"hipDevicePrimaryCtxRetain", CONV_CONTEXT, API_DRIVER}; cuda2hipRename["cuDevicePrimaryCtxReset"] = {"hipDevicePrimaryCtxReset", CONV_CONTEXT, API_DRIVER}; + cuda2hipRename["cuDevicePrimaryCtxRetain"] = {"hipDevicePrimaryCtxRetain", CONV_CONTEXT, API_DRIVER}; cuda2hipRename["cuDevicePrimaryCtxSetFlags"] = {"hipDevicePrimaryCtxSetFlags", CONV_CONTEXT, API_DRIVER}; - // Device + // Device Management cuda2hipRename["cuDeviceGet"] = {"hipGetDevice", CONV_DEV, API_DRIVER}; cuda2hipRename["cuDeviceGetName"] = {"hipDeviceGetName", CONV_DEV, API_DRIVER}; cuda2hipRename["cuDeviceGetCount"] = {"hipGetDeviceCount", CONV_DEV, API_DRIVER}; cuda2hipRename["cuDeviceGetAttribute"] = {"hipDeviceGetAttribute", CONV_DEV, API_DRIVER}; - cuda2hipRename["cuDeviceGetProperties"] = {"hipGetDeviceProperties", CONV_DEV, API_DRIVER}; cuda2hipRename["cuDeviceGetPCIBusId"] = {"hipDeviceGetPCIBusId", CONV_DEV, API_DRIVER}; cuda2hipRename["cuDeviceGetByPCIBusId"] = {"hipDeviceGetByPCIBusId", CONV_DEV, API_DRIVER}; - cuda2hipRename["cuDeviceTotalMem_v2"] = {"hipDeviceTotalMem", CONV_DEV, API_DRIVER}; + + // Device Management [DEPRECATED] cuda2hipRename["cuDeviceComputeCapability"] = {"hipDeviceComputeCapability", CONV_DEV, API_DRIVER}; - cuda2hipRename["cuDeviceCanAccessPeer"] = {"hipDeviceCanAccessPeer", CONV_DEV, API_DRIVER}; + cuda2hipRename["cuDeviceGetProperties"] = {"hipGetDeviceProperties", CONV_DEV, API_DRIVER}; + + // Module Management + cuda2hipRename["cuLinkAddData"] = {"hipLinkAddData", CONV_EVENT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuLinkAddFile"] = {"hipLinkAddFile", CONV_EVENT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuLinkComplete"] = {"hipLinkComplete", CONV_EVENT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuLinkCreate"] = {"hipLinkCreate", CONV_EVENT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuLinkDestroy"] = {"hipLinkDestroy", CONV_EVENT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuModuleGetFunction"] = {"hipModuleGetFunction", CONV_MODULE, API_DRIVER}; + cuda2hipRename["cuModuleGetGlobal_v2"] = {"hipModuleGetGlobal", CONV_MODULE, API_DRIVER}; + cuda2hipRename["cuModuleGetSurfRef"] = {"hipModuleGetSurfRef", CONV_EVENT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuModuleGetTexRef"] = {"hipModuleGetTexRef", CONV_EVENT, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuModuleLoad"] = {"hipModuleLoad", CONV_MODULE, API_DRIVER}; + cuda2hipRename["cuModuleLoadData"] = {"hipModuleLoadData", CONV_MODULE, API_DRIVER}; + cuda2hipRename["cuModuleLoadDataEx"] = {"hipModuleLoadDataEx", CONV_MODULE, API_DRIVER}; + cuda2hipRename["cuModuleLoadFatBinary"] = {"hipModuleLoadFatBinary", CONV_MODULE, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuModuleUnload"] = {"hipModuleUnload", CONV_MODULE, API_DRIVER}; // unsupported yet by HIP [CUDA 8.0.44] // P2P Attributes @@ -960,16 +987,6 @@ struct cuda2hipMap { cuda2hipRename["cuEventRecord"] = {"hipEventRecord", CONV_EVENT, API_DRIVER}; cuda2hipRename["cuEventSynchronize"] = {"hipEventSynchronize", CONV_EVENT, API_DRIVER}; - // Module - cuda2hipRename["cuModuleGetFunction"] = {"hipModuleGetFunction", CONV_MODULE, API_DRIVER}; - cuda2hipRename["cuModuleGetGlobal_v2"] = {"hipModuleGetGlobal", CONV_MODULE, API_DRIVER}; - cuda2hipRename["cuModuleLoad"] = {"hipModuleLoad", CONV_MODULE, API_DRIVER}; - cuda2hipRename["cuModuleLoadData"] = {"hipModuleLoadData", CONV_MODULE, API_DRIVER}; - // unsupported yet by HIP - cuda2hipRename["cuModuleLoadDataEx"] = {"hipModuleLoadDataEx", CONV_MODULE, API_DRIVER}; - cuda2hipRename["cuModuleLoadFatBinary"] = {"hipModuleLoadFatBinary", CONV_MODULE, API_DRIVER, HIP_UNSUPPORTED}; - - cuda2hipRename["cuModuleUnload"] = {"hipModuleUnload", CONV_MODULE, API_DRIVER}; cuda2hipRename["cuLaunchKernel"] = {"hipModuleLaunchKernel", CONV_MODULE, API_DRIVER}; // Streams @@ -986,39 +1003,69 @@ struct cuda2hipMap { cuda2hipRename["cuStreamWaitEvent"] = {"hipStreamWaitEvent", CONV_STREAM, API_DRIVER}; // Memory management + cuda2hipRename["cuArray3DCreate"] = {"hipArray3DCreate", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuArray3DGetDescriptor"] = {"hipArray3DGetDescriptor", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuArrayCreate"] = {"hipArrayCreate", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuArrayDestroy"] = {"hipArrayDestroy", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuArrayGetDescriptor"] = {"hipArrayGetDescriptor", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuIpcCloseMemHandle"] = {"hipIpcCloseMemHandle", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuIpcGetEventHandle"] = {"hipIpcGetEventHandle", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuIpcGetMemHandle"] = {"hipIpcGetMemHandle", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuIpcOpenEventHandle"] = {"hipIpcOpenEventHandle", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuIpcOpenMemHandle"] = {"hipIpcOpenMemHandle", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; cuda2hipRename["cuMemAlloc_v2"] = {"hipMalloc", CONV_MEM, API_DRIVER}; - cuda2hipRename["cuMemFree_v2"] = {"hipFree", CONV_MEM, API_DRIVER}; - - cuda2hipRename["cuMemHostAlloc"] = {"hipHostMalloc", CONV_MEM, API_DRIVER}; - cuda2hipRename["cuMemFreeHost"] = {"hipHostFree", CONV_MEM, API_DRIVER}; - + cuda2hipRename["cuMemAllocHost"] = {"hipMemAllocHost", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemAllocManaged"] = {"hipMemAllocManaged", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemAllocPitch"] = {"hipMemAllocPitch__", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; // Not equal to cudaMemAllocPitch due to different signatures + cuda2hipRename["cuMemcpy"] = {"hipMemcpy__", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; // Not equal to cudaMemcpy due to different signatures + cuda2hipRename["cuMemcpy2D"] = {"hipMemcpy2D__", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; // Not equal to cudaMemcpy2D due to different signatures + cuda2hipRename["cuMemcpy2DAsync"] = {"hipMemcpy2DAsync__", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; // Not equal to cudaMemcpy2DAsync due to different signatures + cuda2hipRename["cuMemcpy2DUnaligned"] = {"hipMemcpy2DUnaligned", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemcpy3D"] = {"hipMemcpy3D__", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; // Not equal to cudaMemcpy3D due to different signatures + cuda2hipRename["cuMemcpy3DAsync"] = {"hipMemcpy3DAsync__", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}; // Not equal to cudaMemcpy3DAsync due to different signatures + cuda2hipRename["cuMemcpy3DPeer"] = {"hipMemcpy3DPeer__", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}; // Not equal to cudaMemcpy3DPeer due to different signatures + cuda2hipRename["cuMemcpy3DPeerAsync"] = {"hipMemcpy3DPeerAsync__", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}; // Not equal to cudaMemcpy3DPeerAsync due to different signatures + cuda2hipRename["cuMemcpyAsync"] = {"hipMemcpyAsync__", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}; // Not equal to cudaMemcpyAsync due to different signatures + cuda2hipRename["cuMemcpyAtoA"] = {"hipMemcpyAtoA", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemcpyAtoD"] = {"hipMemcpyAtoD", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemcpyAtoH"] = {"hipMemcpyAtoH", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemcpyAtoHAsync"] = {"hipMemcpyAtoHAsync", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemcpyDtoA"] = {"hipMemcpyDtoA", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; cuda2hipRename["cuMemcpyDtoD_v2"] = {"hipMemcpyDtoD", CONV_MEM, API_DRIVER}; cuda2hipRename["cuMemcpyDtoDAsync_v2"] = {"hipMemcpyDtoDAsync", CONV_MEM, API_DRIVER}; cuda2hipRename["cuMemcpyDtoH_v2"] = {"hipMemcpyDtoH", CONV_MEM, API_DRIVER}; cuda2hipRename["cuMemcpyDtoHAsync_v2"] = {"hipMemcpyDtoHAsync", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemcpyHtoA"] = {"hipMemcpyHtoA", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemcpyHtoAAsync"] = {"hipMemcpyHtoAAsync", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; cuda2hipRename["cuMemcpyHtoD_v2"] = {"hipMemcpyHtoD", CONV_MEM, API_DRIVER}; cuda2hipRename["cuMemcpyHtoDAsync_v2"] = {"hipMemcpyHtoDAsync", CONV_MEM, API_DRIVER}; - - // unsupported yet by HIP - cuda2hipRename["cuMemsetD8_v2"] = {"hipMemsetD8", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; - cuda2hipRename["cuMemsetD8Async"] = {"hipMemsetD8Async", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; - cuda2hipRename["cuMemsetD2D8_v2"] = {"hipMemsetD2D8", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; - cuda2hipRename["cuMemsetD2D8Async"] = {"hipMemsetD2D8Async", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemcpyPeerAsync"] = {"hipMemcpyPeerAsync__", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; // Not equal to cudaMemcpyPeerAsync due to different signatures + cuda2hipRename["cuMemcpyPeer"] = {"hipMemcpyPeer__", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; // Not equal to cudaMemcpyPeer due to different signatures + cuda2hipRename["cuMemFree_v2"] = {"hipFree", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemFreeHost"] = {"hipHostFree", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemGetAddressRange"] = {"hipMemGetAddressRange", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemGetInfo_v2"] = {"hipMemGetInfo", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemHostAlloc"] = {"hipHostMalloc", CONV_MEM, API_DRIVER}; // API_Runtime ANALOGUE (cudaHostAlloc) + cuda2hipRename["cuMemHostGetDevicePointer"] = {"hipMemHostGetDevicePointer", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemHostGetFlags"] = {"hipMemHostGetFlags", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemHostRegister_v2"] = {"hipHostRegister", CONV_MEM, API_DRIVER}; // API_Runtime ANALOGUE (cudaHostAlloc) + cuda2hipRename["cuMemHostUnregister"] = {"hipHostUnregister", CONV_MEM, API_DRIVER}; // API_Runtime ANALOGUE (cudaHostUnregister) cuda2hipRename["cuMemsetD16_v2"] = {"hipMemsetD16", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; cuda2hipRename["cuMemsetD16Async"] = {"hipMemsetD16Async", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; cuda2hipRename["cuMemsetD2D16_v2"] = {"hipMemsetD2D16", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; cuda2hipRename["cuMemsetD2D16Async"] = {"hipMemsetD2D16Async", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; - - cuda2hipRename["cuMemsetD32_v2"] = {"hipMemset", CONV_MEM, API_DRIVER}; - cuda2hipRename["cuMemsetD32Async"] = {"hipMemsetAsync", CONV_MEM, API_DRIVER}; - // unsupported yet by HIP cuda2hipRename["cuMemsetD2D32_v2"] = {"hipMemsetD2D32", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; cuda2hipRename["cuMemsetD2D32Async"] = {"hipMemsetD2D32Async", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; - - cuda2hipRename["cuMemGetInfo_v2"] = {"hipMemGetInfo", CONV_MEM, API_DRIVER}; - cuda2hipRename["cuMemHostRegister_v2"] = {"hipHostRegister", CONV_MEM, API_DRIVER}; - cuda2hipRename["cuMemHostUnregister"] = {"hipHostUnregister", CONV_MEM, API_DRIVER}; - + cuda2hipRename["cuMemsetD2D8_v2"] = {"hipMemsetD2D8", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemsetD2D8Async"] = {"hipMemsetD2D8Async", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemsetD32_v2"] = {"hipMemset", CONV_MEM, API_DRIVER}; // API_Runtime ANALOGUE (cudaMemset) + cuda2hipRename["cuMemsetD32Async"] = {"hipMemsetAsync", CONV_MEM, API_DRIVER}; // API_Runtime ANALOGUE (cudaMemsetAsync) + cuda2hipRename["cuMemsetD8_v2"] = {"hipMemsetD8", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMemsetD8Async"] = {"hipMemsetD8Async", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMipmappedArrayCreate"] = {"hipMipmappedArrayCreate", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMipmappedArrayDestroy"] = {"hipMipmappedArrayDestroy", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; + cuda2hipRename["cuMipmappedArrayGetLevel"] = {"hipMipmappedArrayGetLevel", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}; + // unsupported yet by HIP [CUDA 8.0.44] cuda2hipRename["cuMemPrefetchAsync"] = {"hipMemPrefetchAsync___", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}; // [CUDA 8.0.44] // no API_Runtime ANALOGUE (cudaMemPrefetchAsync has different signature) cuda2hipRename["cuMemAdvise"] = {"hipMemAdvise", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}; // [CUDA 8.0.44] // API_Runtime ANALOGUE (cudaMemAdvise) cuda2hipRename["cuMemRangeGetAttribute"] = {"hipMemRangeGetAttribute", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}; // [CUDA 8.0.44] // API_Runtime ANALOGUE (cudaMemRangeGetAttribute) @@ -1298,7 +1345,7 @@ struct cuda2hipMap { // Attributes cuda2hipRename["cudaDeviceGetAttribute"] = {"hipDeviceGetAttribute", CONV_DEV, API_RUNTIME}; - cuda2hipRename["cudaDeviceAttr"] = {"hipDeviceAttribute_t", CONV_TYPE, API_RUNTIME}; // API_DRIVER ANALOGUE (CUdevice_attribute) + cuda2hipRename["cudaDeviceAttr"] = {"hipDeviceAttribute_t", CONV_TYPE, API_RUNTIME}; // API_DRIVER ANALOGUE (CUdevice_attribute) cuda2hipRename["cudaDevAttrMaxThreadsPerBlock"] = {"hipDeviceAttributeMaxThreadsPerBlock", CONV_DEV, API_RUNTIME}; // 1 // API_DRIVER ANALOGUE (CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK = 1) cuda2hipRename["cudaDevAttrMaxBlockDimX"] = {"hipDeviceAttributeMaxBlockDimX", CONV_DEV, API_RUNTIME}; // 2 // API_DRIVER ANALOGUE (CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X = 2) cuda2hipRename["cudaDevAttrMaxBlockDimY"] = {"hipDeviceAttributeMaxBlockDimY", CONV_DEV, API_RUNTIME}; // 3 // API_DRIVER ANALOGUE (CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y = 3) From b237b831014f67424ea82be6311506b4dbcce89c Mon Sep 17 00:00:00 2001 From: Sandeep Kumar Date: Mon, 12 Jun 2017 17:14:12 +0530 Subject: [PATCH 4/5] Add peer2peer bandwidth and latency test Change-Id: I6d88e4aa9f6e64096af16579eebef4740734203e [ROCm/clr commit: 01842faa1cc15d6ae1f4bd88a91b2b98b6d6a750] --- .../hipBusBandwidth/hipBusBandwidth.cpp | 395 +++++++++++++++++- 1 file changed, 372 insertions(+), 23 deletions(-) diff --git a/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 09f78543c9..b3b0b3e4a6 100644 --- a/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -16,13 +16,15 @@ int p_iterations = 10; int p_beatsperiteration=1; int p_device = 0; int p_detailed = 0; -bool p_async = 0; +bool p_async = 0; int p_alignedhost = 0; // align host allocs to this granularity, in bytes. 64 or 4096 are good values to try. -int p_onesize = 0; +int p_onesize = 0; bool p_h2d = true; bool p_d2h = true; bool p_bidir = true; +bool p_p2p = false; + //#define NO_CHECK @@ -70,7 +72,7 @@ std::string sizeToString(int size) // **************************************************************************** -hipError_t memcopy(void * dst, const void *src, size_t sizeBytes, enum hipMemcpyKind kind) +hipError_t memcopy(void * dst, const void *src, size_t sizeBytes, enum hipMemcpyKind kind ) { if (p_async) { return hipMemcpyAsync(dst, src, sizeBytes, kind, NULL); @@ -632,6 +634,9 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB) } + + + #define failed(...) \ printf ("error: ");\ printf (__VA_ARGS__);\ @@ -646,6 +651,326 @@ int parseInt(const char *str, int *output) } +void checkPeer2PeerSupport() +{ + int deviceCnt; + hipGetDeviceCount(&deviceCnt); + std::cout << "Total no. of available gpu #" << deviceCnt << "\n" << std::endl; + + for(int deviceId=0; deviceIdhost then host-->GPU2)\n\n" << std::endl; +} + +void enablePeer2Peer(int currentGpu, int peerGpu) +{ + int canAccessPeer; + + hipSetDevice(currentGpu); + hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu); + + if(canAccessPeer==1){ + hipDeviceEnablePeerAccess(peerGpu, 0); + } +} + +void disablePeer2Peer(int currentGpu, int peerGpu) +{ + int canAccessPeer; + + hipSetDevice(currentGpu); + hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu); + + if(canAccessPeer==1){ + hipDeviceDisablePeerAccess(peerGpu); + } +} + +std::string gpuIDToString(int gpuID) +{ + using namespace std; + stringstream ss; + ss << gpuID; + return ss.str(); +} + +void RunBenchmark_P2P_Unidir(ResultDatabase &resultDB) +{ + int gpuCount; + hipGetDeviceCount(&gpuCount); + + int currentGpu, peerGpu; + + long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + + for (currentGpu=0; currentGpu1) { + sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), p_beatsperiteration); + } else { + sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str()); + } + + string cGpu, pGpu; + cGpu = gpuIDToString(currentGpu); + pGpu = gpuIDToString(peerGpu); + + resultDB.AddResult(std::string("p2p_uni") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "GB/sec", speed); + resultDB.AddResult(std::string("P2P_uni") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "ms", t); + + if (p_onesize) { + break; + } + } + + } + + if (p_onesize) { + numMaxFloats = sizeToBytes(p_onesize) / sizeof(float); + } + + disablePeer2Peer(currentGpu, peerGpu); + + hipEventDestroy(start); + hipEventDestroy(stop); + + // Cleanup + hipFree((void*)currentGpuMem); + hipFree((void*)peerGpuMem); + CHECK_HIP_ERROR(); + + hipSetDevice(peerGpu); + hipDeviceReset(); + + hipSetDevice(currentGpu); + hipDeviceReset(); + } + + } + +} + +void RunBenchmark_P2P_Bidir(ResultDatabase &resultDB) { + + int gpuCount; + hipGetDeviceCount(&gpuCount); + + hipStream_t stream[2]; + + int currentGpu, peerGpu; + + long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + + for (currentGpu=0; currentGpu1) { + sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), p_beatsperiteration); + } else { + sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str()); + } + + string cGpu, pGpu; + cGpu = gpuIDToString(currentGpu); + pGpu = gpuIDToString(peerGpu); + + resultDB.AddResult(std::string("p2p_bi") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "GB/sec", speed); + resultDB.AddResult(std::string("P2P_bi") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "ms", t); + + if (p_onesize) { + break; + } + } + + } + + if (p_onesize) { + numMaxFloats = sizeToBytes(p_onesize) / sizeof(float); + } + + disablePeer2Peer(currentGpu, peerGpu); + + hipEventDestroy(start); + hipEventDestroy(stop); + + for (int i=0; i<2; i++) { + hipStreamDestroy(stream[i]); + + hipFree((void*)currentGpuMem[i]); + hipFree((void*)peerGpuMem[i]); + CHECK_HIP_ERROR(); + } + + hipSetDevice(peerGpu); + hipDeviceReset(); + + hipSetDevice(currentGpu); + hipDeviceReset(); + } + } +} + + void printConfig() { hipDeviceProp_t props; hipGetDeviceProperties(&props, p_device); @@ -662,9 +987,9 @@ void help() { printf (" --d2h : Run only device-to-host test.\n"); printf (" --h2d : Run only host-to-device test.\n"); printf (" --bidir : Run only bidir copy test.\n"); + printf (" --p2p : Run only peer2peer unidir and bidir copy tests.\n"); printf (" --verbose : Print verbose status messages as test is run.\n"); printf (" --detailed : Print detailed report (including all trials).\n"); - printf (" --async : Use hipMemcpyAsync(with NULL stream) for H2D/D2H. Default uses hipMemcpy.\n"); printf (" --onesize, -o : Only run one measurement, at specified size (in KB, or if negative in bytes)\n"); @@ -712,6 +1037,12 @@ int parseStandardArguments(int argc, char *argv[]) p_d2h = false; p_bidir = true; + } else if (!strcmp(arg, "--p2p")) { + p_h2d = false; + p_d2h = false; + p_bidir = false; + p_p2p = true; + } else if (!strcmp(arg, "--help") || (!strcmp(arg, "-h"))) { help(); exit(EXIT_SUCCESS); @@ -737,39 +1068,57 @@ int main(int argc, char *argv[]) { parseStandardArguments(argc, argv); - printConfig(); + if (p_p2p) { + checkPeer2PeerSupport(); - if (p_h2d) { - ResultDatabase resultDB; - RunBenchmark_H2D(resultDB); + ResultDatabase resultDB_Unidir, resultDB_Bidir; - resultDB.DumpSummary(std::cout); + RunBenchmark_P2P_Unidir(resultDB_Unidir); + RunBenchmark_P2P_Bidir(resultDB_Bidir); + + resultDB_Unidir.DumpSummary(std::cout); + resultDB_Bidir.DumpSummary(std::cout); if (p_detailed) { - resultDB.DumpDetailed(std::cout); + resultDB_Unidir.DumpDetailed(std::cout); + resultDB_Bidir.DumpDetailed(std::cout); } } + else { + printConfig(); - if (p_d2h) { - ResultDatabase resultDB; - RunBenchmark_D2H(resultDB); + if (p_h2d) { + ResultDatabase resultDB; + RunBenchmark_H2D(resultDB); - resultDB.DumpSummary(std::cout); + resultDB.DumpSummary(std::cout); - if (p_detailed) { - resultDB.DumpDetailed(std::cout); + if (p_detailed) { + resultDB.DumpDetailed(std::cout); + } + } + + if (p_d2h) { + ResultDatabase resultDB; + RunBenchmark_D2H(resultDB); + + resultDB.DumpSummary(std::cout); + + if (p_detailed) { + resultDB.DumpDetailed(std::cout); + } } - } - if (p_bidir) { - ResultDatabase resultDB; - RunBenchmark_Bidir(resultDB); + if (p_bidir) { + ResultDatabase resultDB; + RunBenchmark_Bidir(resultDB); - resultDB.DumpSummary(std::cout); + resultDB.DumpSummary(std::cout); - if (p_detailed) { - resultDB.DumpDetailed(std::cout); + if (p_detailed) { + resultDB.DumpDetailed(std::cout); + } } } } From 86f2c1e395c0ed35e978c6b531edfd2904e17a5f Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 14 Jun 2017 09:45:46 +0530 Subject: [PATCH 5/5] Bump HIP base version to 1.2 Change-Id: I8ecc164afed4383f78579ed86a5c8c11a73b0780 [ROCm/clr commit: 4fa2090b78dde2de636ffb89ab7cebdc58849f9a] --- projects/clr/hipamd/bin/hipconfig | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/bin/hipconfig b/projects/clr/hipamd/bin/hipconfig index 663a1e14cd..39fdab5a99 100755 --- a/projects/clr/hipamd/bin/hipconfig +++ b/projects/clr/hipamd/bin/hipconfig @@ -1,7 +1,7 @@ #!/usr/bin/perl -w $HIP_BASE_VERSION_MAJOR = "1"; -$HIP_BASE_VERSION_MINOR = "0"; +$HIP_BASE_VERSION_MINOR = "2"; # Need perl > 5.10 to use logic-defined or use 5.006; use v5.10.1;