Merge branch 'amd-develop' into amd-master
Change-Id: I3ba6305d6cd8de2381615a7ea54d7954e5514216
[ROCm/clr commit: 5e81056c48]
Esse commit está contido em:
Maneesh Gupta
@@ -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;
|
||||
|
||||
+113
-10
@@ -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**
|
||||
|
||||
|
||||
@@ -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)
|
||||
|
||||
@@ -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; deviceId<deviceCnt; deviceId++)
|
||||
{
|
||||
hipDeviceProp_t props;
|
||||
hipGetDeviceProperties(&props, deviceId);
|
||||
std::cout << "for gpu#" << deviceId << " " << props.name << std::endl;
|
||||
std::cout << " peer2peer supported : ";
|
||||
int PeerCnt=0;
|
||||
for (int i=0; i<deviceCnt; i++) {
|
||||
int isPeer;
|
||||
hipDeviceCanAccessPeer(&isPeer, i, deviceId);
|
||||
if (isPeer) {
|
||||
std::cout << "gpu#" << i << " ";
|
||||
++PeerCnt;
|
||||
}
|
||||
}
|
||||
if (PeerCnt==0)
|
||||
std::cout << "NONE" << " ";
|
||||
|
||||
std::cout << std::endl;
|
||||
std::cout << " peer2peer not supported : ";
|
||||
int nonPeerCnt=0;
|
||||
for (int i=0; i<deviceCnt; i++) {
|
||||
int isPeer;
|
||||
hipDeviceCanAccessPeer(&isPeer, i, deviceId);
|
||||
if (!isPeer && (i!=deviceId)) {
|
||||
std::cout << "gpu#" << i << " ";
|
||||
++nonPeerCnt;
|
||||
}
|
||||
}
|
||||
if (nonPeerCnt==0)
|
||||
std::cout << "NONE" << " ";
|
||||
|
||||
std::cout <<"\n"<<std::endl;
|
||||
}
|
||||
|
||||
std::cout << "\nNote: For non-supported peer2peer devices, memcopy will use/follow the normal behaviour (GPU1-->host 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; currentGpu<gpuCount; currentGpu++) {
|
||||
|
||||
for (peerGpu=0; peerGpu<gpuCount; peerGpu++){
|
||||
|
||||
if (currentGpu == peerGpu)
|
||||
continue;
|
||||
|
||||
float *currentGpuMem, *peerGpuMem;
|
||||
|
||||
hipSetDevice(currentGpu);
|
||||
hipMalloc((void**)¤tGpuMem, sizeof(float) * numMaxFloats);
|
||||
|
||||
hipSetDevice(peerGpu);
|
||||
hipMalloc((void**)&peerGpuMem, sizeof(float) * numMaxFloats);
|
||||
|
||||
enablePeer2Peer(currentGpu, peerGpu);
|
||||
|
||||
hipEvent_t start, stop;
|
||||
hipEventCreate(&start);
|
||||
hipEventCreate(&stop);
|
||||
CHECK_HIP_ERROR();
|
||||
|
||||
// Three passes, forward and backward both
|
||||
for (int pass = 0; pass < p_iterations; pass++)
|
||||
{
|
||||
// store the times temporarily to estimate latency
|
||||
//float times[nSizes];
|
||||
// Step through sizes forward on even passes and backward on odd
|
||||
for (int i = 0; i < nSizes; i++)
|
||||
{
|
||||
int sizeIndex;
|
||||
if ((pass % 2) == 0)
|
||||
sizeIndex = i;
|
||||
else
|
||||
sizeIndex = (nSizes - 1) - i;
|
||||
|
||||
const int thisSize = p_onesize ? p_onesize : sizes[sizeIndex];
|
||||
const int nbytes = sizeToBytes(thisSize);
|
||||
|
||||
hipDeviceSynchronize();
|
||||
|
||||
hipEventRecord(start, 0);
|
||||
|
||||
for (int j=0;j<p_beatsperiteration;j++) {
|
||||
hipMemcpy(peerGpuMem, currentGpuMem, nbytes, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
hipEventRecord(stop, 0);
|
||||
|
||||
hipEventSynchronize(stop);
|
||||
|
||||
float t = 0;
|
||||
hipEventElapsedTime(&t, start, stop);
|
||||
//times[sizeIndex] = t;
|
||||
|
||||
// Convert to GB/sec
|
||||
if (p_verbose)
|
||||
{
|
||||
std::cerr << "size " << sizeToString(thisSize) << " took " << t << " ms\n";
|
||||
}
|
||||
|
||||
double speed = (double(sizeToBytes(thisSize) * p_beatsperiteration) / (1000*1000)) / t;
|
||||
char sizeStr[256];
|
||||
if (p_beatsperiteration>1) {
|
||||
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; currentGpu<gpuCount; currentGpu++) {
|
||||
|
||||
for (peerGpu=0; peerGpu<gpuCount; peerGpu++){
|
||||
|
||||
if (currentGpu == peerGpu)
|
||||
continue;
|
||||
|
||||
float *currentGpuMem[2], *peerGpuMem[2];
|
||||
|
||||
hipSetDevice(currentGpu);
|
||||
hipMalloc((void**)¤tGpuMem[0], sizeof(float) * numMaxFloats);
|
||||
hipMalloc((void**)¤tGpuMem[1], sizeof(float) * numMaxFloats);
|
||||
|
||||
hipSetDevice(peerGpu);
|
||||
hipMalloc((void**)&peerGpuMem[0], sizeof(float) * numMaxFloats);
|
||||
hipMalloc((void**)&peerGpuMem[1], sizeof(float) * numMaxFloats);
|
||||
|
||||
enablePeer2Peer(currentGpu, peerGpu);
|
||||
|
||||
hipEvent_t start, stop;
|
||||
hipEventCreate(&start);
|
||||
hipEventCreate(&stop);
|
||||
CHECK_HIP_ERROR();
|
||||
|
||||
hipStreamCreate(&stream[0]);
|
||||
hipStreamCreate(&stream[1]);
|
||||
|
||||
// Three passes, forward and backward both
|
||||
for (int pass = 0; pass < p_iterations; pass++)
|
||||
{
|
||||
// store the times temporarily to estimate latency
|
||||
//float times[nSizes];
|
||||
// Step through sizes forward on even passes and backward on odd
|
||||
for (int i = 0; i < nSizes; i++)
|
||||
{
|
||||
int sizeIndex;
|
||||
if ((pass % 2) == 0)
|
||||
sizeIndex = i;
|
||||
else
|
||||
sizeIndex = (nSizes - 1) - i;
|
||||
|
||||
const int thisSize = p_onesize ? p_onesize : sizes[sizeIndex];
|
||||
const int nbytes = sizeToBytes(thisSize);
|
||||
|
||||
hipDeviceSynchronize();
|
||||
|
||||
hipEventRecord(start, 0);
|
||||
|
||||
for (int j=0;j<p_beatsperiteration;j++) {
|
||||
hipMemcpyAsync(peerGpuMem[0], currentGpuMem[0], nbytes, hipMemcpyDeviceToDevice, stream[0]);
|
||||
hipMemcpyAsync(currentGpuMem[1], peerGpuMem[1], nbytes, hipMemcpyDeviceToDevice, stream[1]);
|
||||
}
|
||||
|
||||
hipEventRecord(stop, 0);
|
||||
|
||||
hipEventSynchronize(stop);
|
||||
|
||||
float t = 0;
|
||||
hipEventElapsedTime(&t, start, stop);
|
||||
//times[sizeIndex] = t;
|
||||
|
||||
// Convert to GB/sec
|
||||
if (p_verbose)
|
||||
{
|
||||
std::cerr << "size " << sizeToString(thisSize) << " took " << t << " ms\n";
|
||||
}
|
||||
|
||||
double speed = (double(sizeToBytes(thisSize) * p_beatsperiteration) / (1000*1000)) / t;
|
||||
char sizeStr[256];
|
||||
if (p_beatsperiteration>1) {
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
|
||||
|
||||
@@ -33,10 +33,14 @@ THE SOFTWARE.
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#include <hc_am.hpp>
|
||||
#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();
|
||||
|
||||
Referência em uma Nova Issue
Bloquear um usuário