From cbdc8c277c60e6c73ba985a75e70bc61300b6bfe Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 15 Mar 2016 06:30:16 -0500 Subject: [PATCH 01/14] Deprecating hipMallocHost to hipHostAlloc --- tests/src/hipMemcpy.cpp | 6 ++-- tests/src/hipMemcpyAsync.cpp | 8 ++--- tests/src/hipMemcpy_simple.cpp | 4 +-- tests/src/hipMultiThreadStreams2.cpp | 54 ++++++++++++++-------------- tests/src/hipPointerAttrib.cpp | 8 ++--- tests/src/hipStream.h | 4 +-- tests/src/test_common.h | 8 ++--- 7 files changed, 46 insertions(+), 46 deletions(-) diff --git a/tests/src/hipMemcpy.cpp b/tests/src/hipMemcpy.cpp index 7dc8de8d2e..a27135907b 100644 --- a/tests/src/hipMemcpy.cpp +++ b/tests/src/hipMemcpy.cpp @@ -36,7 +36,7 @@ void printSep() // The subroutine allocates memory , copies to device, runs a vector add kernel, copies back, and checks the result. // // IN: numElements controls the number of elements used for allocations. -// IN: usePinnedHost : If true, allocate host with hipMallocHost and is pinned ; else allocate host memory with malloc. +// IN: usePinnedHost : If true, allocate host with hipHostAlloc and is pinned ; else allocate host memory with malloc. // IN: useHostToHost : If true, add an extra host-to-host copy. // IN: useDeviceToDevice : If true, add an extra deviceto-device copy after result is produced. // IN: useMemkindDefault : If true, use memkinddefault (runtime figures out direction). if false, use explicit memcpy direction. @@ -67,8 +67,8 @@ void memcpytest2(size_t numElements, bool usePinnedHost, bool useHostToHost, boo if (useHostToHost) { if (usePinnedHost) { - HIPCHECK ( hipMallocHost(&A_hh, sizeElements) ); - HIPCHECK ( hipMallocHost(&B_hh, sizeElements) ); + HIPCHECK ( hipHostAlloc((void**)&A_hh, sizeElements, hipHostAllocDefault) ); + HIPCHECK ( hipHostAlloc((void**)&B_hh, sizeElements, hipHostAllocDefault) ); } else { A_hh = (T*)malloc(sizeElements); B_hh = (T*)malloc(sizeElements); diff --git a/tests/src/hipMemcpyAsync.cpp b/tests/src/hipMemcpyAsync.cpp index 2254f9bdae..bf50f85bbb 100644 --- a/tests/src/hipMemcpyAsync.cpp +++ b/tests/src/hipMemcpyAsync.cpp @@ -33,7 +33,7 @@ void simpleNegTest() size_t Nbytes = N*sizeof(float); A_malloc = (float*)malloc(Nbytes); - HIPCHECK(hipMallocHost(&A_pinned, Nbytes)); + HIPCHECK(hipHostAlloc((void**)&A_pinned, Nbytes, hipHostAllocDefault)); HIPCHECK(hipMalloc(&A_d, Nbytes)); @@ -61,7 +61,7 @@ struct HostTraits static void *Alloc(size_t sizeBytes) { void *p; - HIPCHECK(hipMallocHost(&p, sizeBytes)); + HIPCHECK(hipHostAlloc((void**)&p, sizeBytes, hipHostAllocDefault)); return p; }; }; @@ -181,8 +181,8 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies, T *A_d; T *A_h1, *A_h2; - HIPCHECK(hipMallocHost(&A_h1, Nbytes)); - HIPCHECK(hipMallocHost(&A_h2, Nbytes)); + HIPCHECK(hipHostAlloc((void**)&A_h1, Nbytes, hipHostAllocDefault)); + HIPCHECK(hipHostAlloc((void**)&A_h2, Nbytes, hipHostAllocDefault)); HIPCHECK(hipMalloc(&A_d, Nbytes)); for (int i=0; i Date: Tue, 15 Mar 2016 13:39:15 -0500 Subject: [PATCH 02/14] v2 deprecating hipMallocHost with hipHostAlloc --- samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 88d547e2c5..3d2aa659ae 100644 --- a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -103,7 +103,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) float *hostMem = NULL; if (p_pinned) { - hipMallocHost((void**)&hostMem, sizeof(float) * numMaxFloats); + hipHostAlloc((void**)&hostMem, sizeof(float) * numMaxFloats, hipHostAllocDefault); while (hipGetLastError() != hipSuccess) { // drop the size and try again @@ -115,7 +115,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) return; } numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; - hipMallocHost((void**)&hostMem, sizeof(float) * numMaxFloats); + hipHostAlloc((void**)&hostMem, sizeof(float) * numMaxFloats, hipHostAllocDefault); } } else @@ -217,9 +217,9 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) float *hostMem2; if (p_pinned) { - hipMallocHost((void**)&hostMem1, sizeof(float)*numMaxFloats); + hipHostAlloc((void**)&hostMem1, sizeof(float)*numMaxFloats, hipHostAllocDefault); hipError_t err1 = hipGetLastError(); - hipMallocHost((void**)&hostMem2, sizeof(float)*numMaxFloats); + hipHostAlloc((void**)&hostMem2, sizeof(float)*numMaxFloats, hipHostAllocDefault); hipError_t err2 = hipGetLastError(); while (err1 != hipSuccess || err2 != hipSuccess) { From 31d8f60e5695b39eecdeca20d7132ae019f6eb3e Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 15 Mar 2016 12:37:24 -0500 Subject: [PATCH 03/14] added performance metrics for kernel dispatch --- samples/1_Utils/hipDispatchLatency/Makefile | 15 +++ .../hipDispatchLatency/hipDispatchLatency.cpp | 120 ++++++++++++++++++ 2 files changed, 135 insertions(+) create mode 100644 samples/1_Utils/hipDispatchLatency/Makefile create mode 100644 samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp diff --git a/samples/1_Utils/hipDispatchLatency/Makefile b/samples/1_Utils/hipDispatchLatency/Makefile new file mode 100644 index 0000000000..b4543e6f99 --- /dev/null +++ b/samples/1_Utils/hipDispatchLatency/Makefile @@ -0,0 +1,15 @@ +HIP_PATH?=$(shell hipconfig -p) +HIPCC=$(HIP_PATH)/bin/hipcc + +EXE=hipDispatchLatency + +all: install + +$(EXE): hipDispatchLatency.cpp + $(HIPCC) hipDispatchLatency.cpp -o $@ + +install: $(EXE) + cp $(EXE) $(HIP_PATH)/bin + +clean: + rm -f *.o $(EXE) diff --git a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp new file mode 100644 index 0000000000..fbdbced063 --- /dev/null +++ b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -0,0 +1,120 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include"hip_runtime.h" +#include +#include + +#define check(msg, status) \ +if(status != hipSuccess){ \ + printf("%s failed.\n",#msg); \ + exit(1); \ +} + +#define LEN 1024*1024 +#define SIZE LEN * sizeof(float) +#define ITER 10000 + +__global__ void One(hipLaunchParm lp, float* Ad){ +} + +int main(){ + +hipError_t err; +float *A, *Ad; + +A = new float[LEN]; + +for(int i=0;i Date: Tue, 15 Mar 2016 21:05:15 -0500 Subject: [PATCH 04/14] Added single kernel launch to sample --- .../hipDispatchLatency/hipDispatchLatency.cpp | 169 ++++++++++-------- 1 file changed, 90 insertions(+), 79 deletions(-) diff --git a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index fbdbced063..a47931958f 100644 --- a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -36,85 +36,96 @@ __global__ void One(hipLaunchParm lp, float* Ad){ int main(){ -hipError_t err; -float *A, *Ad; + hipError_t err; + float *A, *Ad; -A = new float[LEN]; - -for(int i=0;i Date: Tue, 15 Mar 2016 14:22:00 -0500 Subject: [PATCH 05/14] corrected first and second kernel dispatch --- samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index a47931958f..9827b94229 100644 --- a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -64,13 +64,13 @@ int main(){ hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad); hipEventRecord(stop); hipEventElapsedTime(&mS, start, stop); - std::cout<<"First Kernel Launch: \t\t"< Date: Wed, 16 Mar 2016 05:15:03 -0500 Subject: [PATCH 06/14] added cudaHostRegister test --- tests/src/cudaRegister.cu | 80 +++++++++++++++++++++++++++++++++++++++ 1 file changed, 80 insertions(+) create mode 100644 tests/src/cudaRegister.cu diff --git a/tests/src/cudaRegister.cu b/tests/src/cudaRegister.cu new file mode 100644 index 0000000000..1fc456f5fa --- /dev/null +++ b/tests/src/cudaRegister.cu @@ -0,0 +1,80 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include +#include +#include +#include + +#define LEN 1024*1024 +#define SIZE LEN * sizeof(float) +#define ITER 1024*128 + +#define check(msg, status){ \ +if(status != cudaSuccess) { \ + printf("%s failed. \n", #msg); \ +} \ +} + +__global__ void Inc1(float *Ad){ + int tx = threadIdx.x + blockIdx.x * blockDim.x; + if(tx < 1 ){ + for(int i=0;i>>(Ad); + A[1] = -(ITER*1.0f); + cudaDeviceSynchronize(); + std::cout<>>(Ad); + A[1] = -(ITER*1.0f); + cudaDeviceSynchronize(); + std::cout< Date: Wed, 16 Mar 2016 05:24:08 -0500 Subject: [PATCH 07/14] increased iteration size --- tests/src/cudaRegister.cu | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/tests/src/cudaRegister.cu b/tests/src/cudaRegister.cu index 1fc456f5fa..6ea038a71a 100644 --- a/tests/src/cudaRegister.cu +++ b/tests/src/cudaRegister.cu @@ -26,7 +26,7 @@ THE SOFTWARE. #define LEN 1024*1024 #define SIZE LEN * sizeof(float) -#define ITER 1024*128 +#define ITER 1024*1024 #define check(msg, status){ \ if(status != cudaSuccess) { \ @@ -66,15 +66,16 @@ int main(){ dim3 dimGrid(LEN/512,1,1); dim3 dimBlock(512,1,1); Inc1<<>>(Ad); - A[1] = -(ITER*1.0f); + A[0] = -(ITER*1.0f); + std::cout<>>(Ad); - A[1] = -(ITER*1.0f); + A[0] = -(ITER*1.0f); cudaDeviceSynchronize(); - std::cout< Date: Wed, 16 Mar 2016 05:30:29 -0500 Subject: [PATCH 08/14] test/src [v3] clean up --- tests/src/cudaRegister.cu | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/tests/src/cudaRegister.cu b/tests/src/cudaRegister.cu index 6ea038a71a..8d97c7ae59 100644 --- a/tests/src/cudaRegister.cu +++ b/tests/src/cudaRegister.cu @@ -24,7 +24,7 @@ THE SOFTWARE. #include #include -#define LEN 1024*1024 +#define LEN 1024 #define SIZE LEN * sizeof(float) #define ITER 1024*1024 @@ -67,15 +67,16 @@ int main(){ dim3 dimBlock(512,1,1); Inc1<<>>(Ad); A[0] = -(ITER*1.0f); - std::cout<>>(Ad); A[0] = -(ITER*1.0f); + std::cout<<"Diff cache line before completion: \t"< Date: Wed, 16 Mar 2016 07:04:40 -0500 Subject: [PATCH 09/14] tests/src [v4] Added feature for partial writes on CPU --- tests/src/cudaRegister.cu | 20 ++++++++++++++------ 1 file changed, 14 insertions(+), 6 deletions(-) diff --git a/tests/src/cudaRegister.cu b/tests/src/cudaRegister.cu index 8d97c7ae59..b9953a5a3f 100644 --- a/tests/src/cudaRegister.cu +++ b/tests/src/cudaRegister.cu @@ -34,26 +34,32 @@ if(status != cudaSuccess) { \ } \ } -__global__ void Inc1(float *Ad){ +__global__ void Inc1(float *Ad, float *Bd){ int tx = threadIdx.x + blockIdx.x * blockDim.x; if(tx < 1 ){ for(int i=0;i>>(Ad); + Inc1<<>>(Ad, Bd); + sleep(3); A[0] = -(ITER*1.0f); std::cout<<"Same cache line before completion: \t"<< A[0]<>>(Ad); + Inc2<<>>(Ad, Bd); + sleep(3); A[0] = -(ITER*1.0f); std::cout<<"Diff cache line before completion: \t"< Date: Wed, 16 Mar 2016 07:16:51 -0500 Subject: [PATCH 10/14] Added performance test for memcpy --- tests/src/hipPerfMemcpy.cpp | 41 +++++++++++++++++++++++++++++++++++++ 1 file changed, 41 insertions(+) create mode 100644 tests/src/hipPerfMemcpy.cpp diff --git a/tests/src/hipPerfMemcpy.cpp b/tests/src/hipPerfMemcpy.cpp new file mode 100644 index 0000000000..d351dbc868 --- /dev/null +++ b/tests/src/hipPerfMemcpy.cpp @@ -0,0 +1,41 @@ +#include"test_common.h" +#include +#include + +#define NUM_SIZE 8 +#define NUM_ITER 12 +static size_t size[NUM_SIZE]; + +void setup(){ + for(int i=0;i Date: Wed, 16 Mar 2016 07:32:54 -0500 Subject: [PATCH 11/14] src/ fixed hipHostAllocDefault flags --- src/hip_hcc.cpp | 34 +++++++++++++++++----------------- tests/src/hipHostAlloc.cpp | 9 ++++++--- 2 files changed, 23 insertions(+), 20 deletions(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 41aea51c0c..af03c49dcd 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -2054,31 +2054,31 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){ auto device = ihipGetTlsDefaultDevice(); if(device){ - if(flags & hipHostAllocDefault){ - const unsigned am_flags = amHostPinned; + if(flags == hipHostAllocDefault){ + const unsigned am_flags = amHostPinned; - *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); - if(sizeBytes && (*ptr == NULL)){ - hip_status = hipErrorMemoryAllocation; - }else{ - hc::am_memtracker_update(*ptr, device->_device_index, 0); - } - tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); + *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); + if(sizeBytes && (*ptr == NULL)){ + hip_status = hipErrorMemoryAllocation; + }else{ + hc::am_memtracker_update(*ptr, device->_device_index, 0); + } + tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); } if(flags & hipHostAllocMapped){ - const unsigned am_flags = amHostPinned; + const unsigned am_flags = amHostPinned; - *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); - if(sizeBytes && (*ptr == NULL)){ - hip_status = hipErrorMemoryAllocation; - }else{ - hc::am_memtracker_update(*ptr, device->_device_index, flags); + *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); + if(sizeBytes && (*ptr == NULL)){ + hip_status = hipErrorMemoryAllocation; + }else{ + hc::am_memtracker_update(*ptr, device->_device_index, flags); // void *srcPtr; // hsa_status_t hsa_status = hsa_amd_memory_lock((*ptr), sizeBytes, &device->_hsa_agent, 1, &srcPtr); // assert(hsa_status == HSA_STATUS_SUCCESS); // hc::am_memtracker_add(srcPtr, sizeBytes, device->_acc, false); - } - tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); + } + tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); } } return ihipLogStatus(hip_status); diff --git a/tests/src/hipHostAlloc.cpp b/tests/src/hipHostAlloc.cpp index 072582d85f..fe82a0c78f 100644 --- a/tests/src/hipHostAlloc.cpp +++ b/tests/src/hipHostAlloc.cpp @@ -39,14 +39,14 @@ int device; HIPCHECK(hipGetDevice(&device)); HIPCHECK(hipGetDeviceProperties(&prop, device)); if(prop.canMapHostMemory != 1){ -std::cout<<"Exiting..."< Date: Wed, 16 Mar 2016 08:01:53 -0500 Subject: [PATCH 12/14] changed flag in hipHostRegister --- src/hip_hcc.cpp | 16 ++++------------ 1 file changed, 4 insertions(+), 12 deletions(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index af03c49dcd..c839b9cf8e 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -2055,9 +2055,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){ if(device){ if(flags == hipHostAllocDefault){ - const unsigned am_flags = amHostPinned; - - *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); + *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned); if(sizeBytes && (*ptr == NULL)){ hip_status = hipErrorMemoryAllocation; }else{ @@ -2066,17 +2064,11 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){ tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); } if(flags & hipHostAllocMapped){ - const unsigned am_flags = amHostPinned; - - *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); + *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned); if(sizeBytes && (*ptr == NULL)){ hip_status = hipErrorMemoryAllocation; }else{ hc::am_memtracker_update(*ptr, device->_device_index, flags); -// void *srcPtr; -// hsa_status_t hsa_status = hsa_amd_memory_lock((*ptr), sizeBytes, &device->_hsa_agent, 1, &srcPtr); -// assert(hsa_status == HSA_STATUS_SUCCESS); -// hc::am_memtracker_add(srcPtr, sizeBytes, device->_acc, false); } tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); } @@ -2142,10 +2134,10 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) return ihipLogStatus(hipErrorInvalidValue); } if(device){ - if(flags == hipHostAllocDefault){ + if(flags == hipHostRegisterDefault){ hsa_status_t hsa_status = hsa_amd_memory_lock(hostPtr, sizeBytes, &device->_hsa_agent, 1, &srcPtr); if(hsa_status == HSA_STATUS_SUCCESS){ - hip_status = hipSuccess; + hip_status = hipSuccess; }else{ hip_status = hipErrorMemoryAllocation; } From 456f8b845584b961de5c7ca20ce55308dae0b671 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 18 Mar 2016 11:23:44 -0500 Subject: [PATCH 13/14] Create CUDA_Runtime_API_functions_supported_by_HIP.md --- ..._Runtime_API_functions_supported_by_HIP.md | 269 ++++++++++++++++++ 1 file changed, 269 insertions(+) create mode 100644 CUDA_Runtime_API_functions_supported_by_HIP.md diff --git a/CUDA_Runtime_API_functions_supported_by_HIP.md b/CUDA_Runtime_API_functions_supported_by_HIP.md new file mode 100644 index 0000000000..b22856ee58 --- /dev/null +++ b/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -0,0 +1,269 @@ +**1. Device Management** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaChooseDevice` | | Select compute-device which best matches criteria. | +| `cudaDeviceGetAttribute` | | Returns information about the device. | +| `cudaDeviceGetByPCIBusId` | | Returns a handle to a compute device. | +| `cudaDeviceGetCacheConfig` | `hipDeviceGetCacheConfig` | Returns the preferred cache configuration for the current device. | +| `cudaDeviceGetLimit` | | Returns resource limits. | +| `cudaDeviceGetPCIBusId` | | Returns a PCI Bus Id string for the device. | +| `cudaDeviceGetSharedMemConfig` | `hipDeviceGetSharedMemConfig` | Returns the shared memory configuration for the current device. | +| `cudaDeviceGetStreamPriorityRange` | | Returns numerical values that correspond to the least and greatest stream priorities. | +| `cudaDeviceReset` | `hipDeviceReset` | Destroy all allocations and reset all state on the current device in the current process. | +| `cudaDeviceSetCacheConfig` | `hipDeviceSetCacheConfig` | Sets the preferred cache configuration for the current device. | +| `cudaDeviceSetLimit` | | Set resource limits. | +| `cudaDeviceSetSharedMemConfig` | `hipDeviceSetSharedMemConfig` | Sets the shared memory configuration for the current device. | +| `cudaDeviceSynchronize` | `hipDeviceSynchronize` | Wait for compute device to finish. | +| `cudaGetDevice` | `hipGetDevice` | Returns which device is currently being used. | +| `cudaGetDeviceCount` | `hipGetDeviceCount` | Returns the number of compute-capable devices. | +| `cudaGetDeviceFlags` | | Gets the flags for the current device. | +| `cudaGetDeviceProperties` | `hipDeviceGetProperties` | Returns information about the compute-device. | +| `cudaIpcCloseMemHandle` | | Close memory mapped with cudaIpcOpenMemHandle. | +| `cudaIpcGetEventHandle` | | Gets an interprocess handle for a previously allocated event. | +| `cudaIpcGetMemHandle` | | Gets an interprocess memory handle for an existing device memory allocation. | +| `cudaIpcOpenEventHandle` | | Opens an interprocess event handle for use in the current process. | +| `cudaIpcOpenMemHandle` | | Opens an interprocess memory handle exported from another process and returns a device pointer usable in the local process. | +| `cudaSetDevice` | `hipSetDevice` | Set device to be used for GPU executions. | +| `cudaSetDeviceFlags` | | Sets flags to be used for device executions. | +| `cudaSetValidDevices` | | Set a list of devices that can be used for CUDA. |` + +**2. Error Handling** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaGetErrorName` | `hipGetErrorName` | Returns the string representation of an error code enum name. | +| `cudaGetErrorString` | `hipGetErrorString` | Returns the description string for an error code. | +| `cudaGetLastError` | `hipGetLastError` | Returns the last error from a runtime call. | +| `cudaPeekAtLastError` | `hipPeekAtLastError` | Returns the last error from a runtime call. | + +**3. Stream Management** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaStreamAddCallback` | | Add a callback to a compute stream. | +| `cudaStreamAttachMemAsync` | | Attach memory to a stream asynchronously. | +| `cudaStreamCreate` | `hipStreamCreate` | Create an asynchronous stream. | +| `cudaStreamCreateWithFlags` | `hipStreamCreateWithFlags` | Create an asynchronous stream. | +| `cudaStreamCreateWithPriority` | | Create an asynchronous stream with the specified priority. | +| `cudaStreamDestroy` | `hipStreamDestroy` | Destroys and cleans up an asynchronous stream. | +| `cudaStreamGetFlags` | | Query the flags of a stream. | +| `cudaStreamGetPriority` | | Query the priority of a stream. | +| `cudaStreamQuery` | | Queries an asynchronous stream for completion status. | +| `cudaStreamSynchronize` | `hipStreamSynchronize` | Waits for stream tasks to complete. | +| `cudaStreamWaitEvent` | `hipStreamWaitEvent` | Make a compute stream wait on an event. | + +**4. Event Management** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaEventCreate` | `hipEventCreate` | Creates an event object. | +| `cudaEventCreateWithFlags` | `hipEventCreateWithFlags` | Creates an event object with the specified flags. | +| `cudaEventDestroy` | `hipEventDestroy` | Destroys an event object. | +| `cudaEventElapsedTime` | `hipEventElapsedTime` | Computes the elapsed time between events. | +| `cudaEventQuery` | `hipEventQuery` | Queries an event's status. | +| `cudaEventRecord` | `hipEventRecord` | Records an event. | +| `cudaEventSynchronize` | `hipEventSynchronize` | Waits for an event to complete. | + +**5. Execution Control** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaFuncGetAttributes` | | Find out attributes for a given function. | +| `cudaFuncSetCacheConfig` | `hipFuncSetCacheConfig` | Sets the preferred cache configuration for a device function. | +| `cudaFuncSetSharedMemConfig` | | Sets the shared memory configuration for a device function. | +| `cudaGetParameterBuffer` | | Obtains a parameter buffer. | +| `cudaGetParameterBufferV2` | | Launches a specified kernel. | +| `cudaLaunchKernel` | `hipLaunchKernel` | Launches a device function. | +| `cudaSetDoubleForDevice` | | Converts a double argument to be executed on a device. | +| `cudaSetDoubleForHost` | | Converts a double argument after execution on a device. | + +**6. Occupancy** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaOccupancyMaxActiveBlocksPerMultiprocessor` | | Returns occupancy for a device function. | +| `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | | Returns occupancy for a device function with the specified flags. | + +**7. Memory Management** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaArrayGetInfo` | | Gets info about the specified cudaArray. | +| `cudaFree` | `hipFree` | Frees memory on the device. | +| `cudaFreeArray` | | Frees an array on the device. | +| `cudaFreeHost` | `hipHostFree` | Frees page-locked memory. | +| `cudaFreeMipmappedArray` | | Frees a mipmapped array on the device. | +| `cudaGetMipmappedArrayLevel` | | Gets a mipmap level of a CUDA mipmapped array. | +| `cudaGetSymbolAddress` | | Finds the address associated with a CUDA symbol. | +| `cudaGetSymbolSize` | | Finds the size of the object associated with a CUDA symbol. | +| `cudaHostAlloc` | | Allocates page-locked memory on the host. | +| `cudaHostGetDevicePointer` | | Passes back device pointer of mapped host memory allocated by cudaHostAlloc or registered by cudaHostRegister. | +| `cudaHostGetFlags` | | Passes back flags used to allocate pinned host memory allocated by cudaHostAlloc. | +| `cudaHostRegister` | | Registers an existing host memory range for use by CUDA. | +| `cudaHostUnregister` | | Unregisters a memory range that was registered with cudaHostRegister. | +| `cudaMalloc` | `hipMalloc` | Allocate memory on the device. | +| `cudaMalloc3D` | | Allocates logical 1D, 2D, or 3D memory objects on the device. | +| `cudaMalloc3DArray` | | Allocate an array on the device. | +| `cudaMallocArray` | | Allocate an array on the device. | +| `cudaMallocHost` | `hipHostAlloc` | Allocates page-locked memory on the host. | +| `cudaMallocManaged` | | Allocates memory that will be automatically managed by the Unified Memory system. | +| `cudaMallocMipmappedArray` | | Allocate a mipmapped array on the device. | +| `cudaMallocPitch` | | Allocates pitched memory on the device. | +| `cudaMemGetInfo` | | Gets free and total device memory. | +| `cudaMemcpy` | `hipMemcpy` | Copies data between host and device. | +| `cudaMemcpy2D` | | Copies data between host and device. | +| `cudaMemcpy2DArrayToArray` | | Copies data between host and device. | +| `cudaMemcpy2DAsync` | | Copies data between host and device. | +| `cudaMemcpy2DFromArray` | | Copies data between host and device. | +| `cudaMemcpy2DFromArrayAsync` | | Copies data between host and device. | +| `cudaMemcpy2DToArray` | | Copies data between host and device. | +| `cudaMemcpy2DToArrayAsync` | | Copies data between host and device. | +| `cudaMemcpy3D` | | Copies data between 3D objects. | +| `cudaMemcpy3DAsync` | | Copies data between 3D objects. | +| `cudaMemcpy3DPeer` | | Copies memory between devices. | +| `cudaMemcpy3DPeerAsync` | | Copies memory between devices asynchronously. | +| `cudaMemcpyArrayToArray` | | Copies data between host and device. | +| `cudaMemcpyAsync` | `hipMemcpyAsync` | Copies data between host and device. | +| `cudaMemcpyFromArray` | `MemcpyFromArray` | Copies data between host and device. | +| `cudaMemcpyFromArrayAsync` | | Copies data between host and device. | +| `cudaMemcpyFromSymbol` | `hipMemcpyFromSymbol` | Copies data from the given symbol on the device. | +| `cudaMemcpyFromSymbolAsync` | | Copies data from the given symbol on the device. | +| `cudaMemcpyPeer` | `hipMemcpyPeer` | Copies memory between two devices. | +| `cudaMemcpyPeerAsync` | `hipMemcpyPeerAsync` | Copies memory between two devices asynchronously. | +| `cudaMemcpyToArray` | | Copies data between host and device. | +| `cudaMemcpyToArrayAsync` | | Copies data between host and device. | +| `cudaMemcpyToSymbol` | `hipMemcpyToSymbol` | Copies data to the given symbol on the device. | +| `cudaMemcpyToSymbolAsync` | | Copies data to the given symbol on the device. | +| `cudaMemset` | `hipMemset` | Initializes or sets device memory to a value. | +| `cudaMemset2D` | | Initializes or sets device memory to a value. | +| `cudaMemset2DAsync` | | Initializes or sets device memory to a value. | +| `cudaMemset3D` | | Initializes or sets device memory to a value. | +| `cudaMemset3DAsync` | | Initializes or sets device memory to a value. | +| `cudaMemsetAsync` | `hipMemsetAsync` | Initializes or sets device memory to a value. | +| `make\_cudaExtent` | | Returns a cudaExtent based on input parameters. | +| `make\_cudaPitchedPtr` | | Returns a cudaPitchedPtr based on input parameters. | +| `make\_cudaPos` | | Returns a cudaPos based on input parameters. | + +**8. Unified Addressing** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaPointerGetAttributes` | | Returns attributes about a specified pointer. | + +**9. Peer Device Memory Access** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaDeviceCanAccessPeer` | `hipDeviceCanAccessPeer` | Queries if a device may directly access a peer device's memory. | +| `cudaDeviceDisablePeerAccess` | `hipDeviceDisablePeerAccess` | Disables direct access to memory allocations on a peer device. | +| `cudaDeviceEnablePeerAccess` | `hipDeviceEnablePeerAccess` | Enables direct access to memory allocations on a peer device. | + +**10. OpenGL Interoperability** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaGLGetDevices` | | Gets the CUDA devices associated with the current OpenGL context. | +| `cudaGraphicsGLRegisterBuffer` | | Registers an OpenGL buffer object. | +| `cudaGraphicsGLRegisterImage` | | Register an OpenGL texture or renderbuffer object. | +| `cudaWGLGetDevice` | | Gets the CUDA device associated with hGpu. | + +**11. Graphics Interoperability** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaGraphicsMapResources` | | Map graphics resources for access by CUDA. | +| `cudaGraphicsResourceGetMappedMipmappedArray` | | Get a mipmapped array through which to access a mapped graphics resource. | +| `cudaGraphicsResourceGetMappedPointer` | | Get a device pointer through which to access a mapped graphics resource. | +| `cudaGraphicsResourceSetMapFlags` | | Set usage flags for mapping a graphics resource. | +| `cudaGraphicsSubResourceGetMappedArray` | | Get an array through which to access a subresource of a mapped graphics resource. | +| `cudaGraphicsUnmapResources` | | Unmap graphics resources. | +| `cudaGraphicsUnregisterResource` | | Unregisters a graphics resource for access by CUDA. | + +**12. Texture Reference Management** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaBindTexture` | | Binds a memory area to a texture. | +| `cudaBindTexture2D` | | Binds a 2D memory area to a texture. | +| `cudaBindTextureToArray` | | Binds an array to a texture. | +| `cudaBindTextureToMipmappedArray` | | Binds a mipmapped array to a texture. | +| `cudaCreateChannelDesc` | | Returns a channel descriptor using the specified format. | +| `cudaGetChannelDesc` | | Get the channel descriptor of an array. | +| `cudaGetTextureAlignmentOffset` | | Get the alignment offset of a texture. | +| `cudaGetTextureReference` | | Get the texture reference associated with a symbol. | +| `cudaUnbindTexture` | | Unbinds a texture. | + +**13. Surface Reference Management** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaBindSurfaceToArray` | | Binds an array to a surface. | +| `cudaGetSurfaceReference` | | Get the surface reference associated with a symbol. | + +**14. Texture Object Management** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaCreateTextureObject` | | Creates a texture object. | +| `cudaDestroyTextureObject` | | Destroys a texture object. | +| `cudaGetTextureObjectResourceDesc` | | Returns a texture object's resource descriptor. | +| `cudaGetTextureObjectResourceViewDesc` | | Returns a texture object's resource view descriptor. | +| `cudaGetTextureObjectTextureDesc` | | Returns a texture object's texture descriptor. | + +**15. Surface Object Management** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaCreateSurfaceObject` | | Creates a surface object. | +| `cudaDestroySurfaceObject` | | Destroys a surface object. | +| `cudaGetSurfaceObjectResourceDesc` | | Returns a surface object's resource descriptor Returns the resource descriptor for the surface object specified by surfObject. | + +**16. Version Management** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaDriverGetVersion` | `hipDriverGetVersion` | Returns the CUDA driver version. | +| `cudaRuntimeGetVersion` | | Returns the CUDA Runtime version. | + +**17. C++ API Routines (7.0 contains, 7.5 doesn’t)** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaBindSurfaceToArra`y | | Binds an array to a surface. | +| `cudaBindTexture` | | Binds a memory area to a texture. | +| `cudaBindTexture2D` | | Binds a 2D memory area to a texture. | +| `cudaBindTextureToArray` | | Binds an array to a texture. | +| `cudaBindTextureToMipmappedArray` | | Binds a mipmapped array to a texture. | +| `cudaCreateChannelDesc` | | Returns a channel descriptor using the specified format. | +| `cudaEventCreate` | | Creates an event object with the specified flags. | +| `cudaFuncGetAttributes` | | Find out attributes for a given function. | +| `cudaFuncSetCacheConfig` | | Sets the preferred cache configuration for a device function. | +| `cudaGetSymbolAddress` | | Finds the address associated with a CUDA symbol | +| `cudaGetSymbolSize` | | Finds the size of the object associated with a CUDA symbol. | +| `cudaGetTextureAlignmentOffset` | | Get the alignment offset of a texture. | +| `cudaLaunch` | | Launches a device function. | +| `cudaLaunchKernel` | | Launches a device function. | +| `cudaMallocHost` | | Allocates page-locked memory on the host | +| `cudaMallocManaged` | | Allocates memory that will be automatically managed by the Unified Memory system. | +| `cudaMemcpyFromSymbol` | | Copies data from the given symbol on the device. | +| `cudaMemcpyFromSymbolAsync` | | Copies data from the given symbol on the device. | +| `cudaMemcpyToSymbol` | | Copies data to the given symbol on the device. | +| `cudaMemcpyToSymbolAsync` | | Async copies data to the given symbol on the device. | +| `cudaOccupancyMaxActiveBlocksPerMultiprocessor` | | Returns occupancy for a device function. | +| `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | | Returns occupancy for a device function with the specified flags. | +| `cudaOccupancyMaxPotentialBlockSize` | | Returns grid and block size that achieves maximum potential occupancy for a device function. | +| `cudaOccupancyMaxPotentialBlockSizeVariableSMem` | | Returns grid and block size that achieves maximum potential occupancy for a device function. | +| `cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags` | | Returns grid and block size that achieves maximum potential occupancy for a device function. | +| `cudaOccupancyMaxPotentialBlockSizeWithFlags` | | Returns grid and block size that achived maximum potential occupancy for a device function with the specified flags. | +| `cudaSetupArgument` | | Configure a device launch. | +| `cudaStreamAttachMemAsync` | | Attach memory to a stream asynchronously. | +| `cudaUnbindTexture` | | Unbinds a texture. | + +**18. Profiler Control** + +| **CUDA** | **HIP** | **CUDA description** | +|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| +| `cudaProfilerInitialize` | | Initialize the CUDA profiler. | +| `cudaProfilerStart` | `hipProfilerStart` | Enable profiling. | +| `cudaProfilerStop` | `hipProfilerStop` | Disable profiling. | From 27b6a620d9d7f5301e6978726f3f39dcd58d87f3 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 18 Mar 2016 11:28:06 -0500 Subject: [PATCH 14/14] Update CUDA_Runtime_API_functions_supported_by_HIP.md --- CUDA_Runtime_API_functions_supported_by_HIP.md | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/CUDA_Runtime_API_functions_supported_by_HIP.md b/CUDA_Runtime_API_functions_supported_by_HIP.md index b22856ee58..f1f0599626 100644 --- a/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -97,9 +97,9 @@ | `cudaGetMipmappedArrayLevel` | | Gets a mipmap level of a CUDA mipmapped array. | | `cudaGetSymbolAddress` | | Finds the address associated with a CUDA symbol. | | `cudaGetSymbolSize` | | Finds the size of the object associated with a CUDA symbol. | -| `cudaHostAlloc` | | Allocates page-locked memory on the host. | -| `cudaHostGetDevicePointer` | | Passes back device pointer of mapped host memory allocated by cudaHostAlloc or registered by cudaHostRegister. | -| `cudaHostGetFlags` | | Passes back flags used to allocate pinned host memory allocated by cudaHostAlloc. | +| `cudaHostAlloc` | `hipHostAlloc` | Allocates page-locked memory on the host. | +| `cudaHostGetDevicePointer` | `hipHostGetDevicePointer` | Passes back device pointer of mapped host memory allocated by cudaHostAlloc or registered by cudaHostRegister. | +| `cudaHostGetFlags` | `hipHostGetFlags` | Passes back flags used to allocate pinned host memory allocated by cudaHostAlloc. | | `cudaHostRegister` | | Registers an existing host memory range for use by CUDA. | | `cudaHostUnregister` | | Unregisters a memory range that was registered with cudaHostRegister. | | `cudaMalloc` | `hipMalloc` | Allocate memory on the device. | @@ -149,7 +149,7 @@ | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------| -| `cudaPointerGetAttributes` | | Returns attributes about a specified pointer. | +| `cudaPointerGetAttributes` | `hipPointerGetAttributes` | Returns attributes about a specified pointer. | **9. Peer Device Memory Access** @@ -227,6 +227,7 @@ | `cudaRuntimeGetVersion` | | Returns the CUDA Runtime version. | **17. C++ API Routines (7.0 contains, 7.5 doesn’t)** +> Will not support for HIP (probably) | **CUDA** | **HIP** | **CUDA description** | |-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|