Merge 'master' into 'amd-master'
Change-Id: I79c383bed37cf6278ddb733ddf3bf3661e3dad35
[ROCm/hip commit: 355b722246]
Этот коммит содержится в:
Jenkins
@@ -1027,7 +1027,7 @@ sub simpleSubstitutions {
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_CONTEXT\b/hipErrorInvalidContext/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_DEVICE\b/hipErrorInvalidDevice/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_GRAPHICS_CONTEXT\b/hipErrorInvalidGraphicsContext/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_HANDLE\b/hipErrorInvalidResourceHandle/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_HANDLE\b/hipErrorInvalidHandle/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_IMAGE\b/hipErrorInvalidImage/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_PTX\b/hipErrorInvalidKernelFile/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_SOURCE\b/hipErrorInvalidSource/g;
|
||||
@@ -1042,13 +1042,15 @@ sub simpleSubstitutions {
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_NOT_MAPPED_AS_ARRAY\b/hipErrorNotMappedAsArray/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_NOT_MAPPED_AS_POINTER\b/hipErrorNotMappedAsPointer/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_NOT_READY\b/hipErrorNotReady/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_NOT_SUPPORTED\b/hipErrorNotSupported/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_NO_BINARY_FOR_GPU\b/hipErrorNoBinaryForGpu/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_NO_DEVICE\b/hipErrorNoDevice/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_OPERATING_SYSTEM\b/hipErrorOperatingSystem/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_OUT_OF_MEMORY\b/hipErrorMemoryAllocation/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_OUT_OF_MEMORY\b/hipErrorOutOfMemory/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED\b/hipErrorPeerAccessAlreadyEnabled/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_PEER_ACCESS_NOT_ENABLED\b/hipErrorPeerAccessNotEnabled/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_PEER_ACCESS_UNSUPPORTED\b/hipErrorPeerAccessUnsupported/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_PRIMARY_CONTEXT_ACTIVE\b/hipErrorSetOnActiveProcess/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_PROFILER_ALREADY_STARTED\b/hipErrorProfilerAlreadyStarted/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_PROFILER_ALREADY_STOPPED\b/hipErrorProfilerAlreadyStopped/g;
|
||||
$ft{'numeric_literal'} += s/\bCUDA_ERROR_PROFILER_DISABLED\b/hipErrorProfilerDisabled/g;
|
||||
@@ -1472,8 +1474,15 @@ sub simpleSubstitutions {
|
||||
$ft{'numeric_literal'} += s/\bcudaDevAttrTextureAlignment\b/hipDeviceAttributeTextureAlignment/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaDevAttrTotalConstantMemory\b/hipDeviceAttributeTotalConstantMemory/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaDevAttrWarpSize\b/hipDeviceAttributeWarpSize/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorAlreadyAcquired\b/hipErrorAlreadyAcquired/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorAlreadyMapped\b/hipErrorAlreadyMapped/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorArrayIsMapped\b/hipErrorArrayIsMapped/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorAssert\b/hipErrorAssert/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorCudartUnloading\b/hipErrorDeinitialized/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorDeviceAlreadyInUse\b/hipErrorContextAlreadyInUse/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorDeviceUninitilialized\b/hipErrorInvalidContext/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorECCUncorrectable\b/hipErrorECCNotCorrectable/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorFileNotFound\b/hipErrorFileNotFound/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorHostMemoryAlreadyRegistered\b/hipErrorHostMemoryAlreadyRegistered/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorHostMemoryNotRegistered\b/hipErrorHostMemoryNotRegistered/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorIllegalAddress\b/hipErrorIllegalAddress/g;
|
||||
@@ -1488,17 +1497,22 @@ sub simpleSubstitutions {
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorInvalidMemcpyDirection\b/hipErrorInvalidMemcpyDirection/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorInvalidPtx\b/hipErrorInvalidKernelFile/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorInvalidResourceHandle\b/hipErrorInvalidResourceHandle/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorInvalidSource\b/hipErrorInvalidSource/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorInvalidSymbol\b/hipErrorInvalidSymbol/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorInvalidValue\b/hipErrorInvalidValue/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorLaunchFailure\b/hipErrorLaunchFailure/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorLaunchOutOfResources\b/hipErrorLaunchOutOfResources/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorLaunchTimeout\b/hipErrorLaunchTimeOut/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorMapBufferObjectFailed\b/hipErrorMapFailed/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorMapBufferObjectFailed\b/hipErrorMapBufferObjectFailed/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorMemoryAllocation\b/hipErrorMemoryAllocation/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorMissingConfiguration\b/hipErrorMissingConfiguration/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorNoDevice\b/hipErrorNoDevice/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorNoKernelImageForDevice\b/hipErrorNoBinaryForGpu/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorNotMapped\b/hipErrorNotMapped/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorNotMappedAsArray\b/hipErrorNotMappedAsArray/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorNotMappedAsPointer\b/hipErrorNotMappedAsPointer/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorNotReady\b/hipErrorNotReady/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorNotSupported\b/hipErrorNotSupported/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorOperatingSystem\b/hipErrorOperatingSystem/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorPeerAccessAlreadyEnabled\b/hipErrorPeerAccessAlreadyEnabled/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorPeerAccessNotEnabled\b/hipErrorPeerAccessNotEnabled/g;
|
||||
@@ -1508,8 +1522,10 @@ sub simpleSubstitutions {
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorProfilerAlreadyStopped\b/hipErrorProfilerAlreadyStopped/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorProfilerDisabled\b/hipErrorProfilerDisabled/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorProfilerNotInitialized\b/hipErrorProfilerNotInitialized/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorSetOnActiveProcess\b/hipErrorSetOnActiveProcess/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorSharedObjectInitFailed\b/hipErrorSharedObjectInitFailed/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorSharedObjectSymbolNotFound\b/hipErrorSharedObjectSymbolNotFound/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorSymbolNotFound\b/hipErrorNotFound/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorUnknown\b/hipErrorUnknown/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorUnmapBufferObjectFailed\b/hipErrorUnmapFailed/g;
|
||||
$ft{'numeric_literal'} += s/\bcudaErrorUnsupportedLimit\b/hipErrorUnsupportedLimit/g;
|
||||
|
||||
@@ -382,7 +382,7 @@
|
||||
| typedef |`cudaError_enum` |***`hipError_t`*** |
|
||||
| 0 |*`CUDA_SUCCESS`* |*`hipSuccess`* |
|
||||
| 1 |*`CUDA_ERROR_INVALID_VALUE`* |*`hipErrorInvalidValue`* |
|
||||
| 2 |*`CUDA_ERROR_OUT_OF_MEMORY`* |*`hipErrorMemoryAllocation`* |
|
||||
| 2 |*`CUDA_ERROR_OUT_OF_MEMORY`* |*`hipErrorOutOfMemory`* |
|
||||
| 3 |*`CUDA_ERROR_NOT_INITIALIZED`* |*`hipErrorNotInitialized`* |
|
||||
| 4 |*`CUDA_ERROR_DEINITIALIZED`* |*`hipErrorDeinitialized`* |
|
||||
| 5 |*`CUDA_ERROR_PROFILER_DISABLED`* |*`hipErrorProfilerDisabled`* |
|
||||
@@ -416,7 +416,7 @@
|
||||
| 302 |*`CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND`* |*`hipErrorSharedObjectSymbolNotFound`* |
|
||||
| 303 |*`CUDA_ERROR_SHARED_OBJECT_INIT_FAILED`* |*`hipErrorSharedObjectInitFailed`* |
|
||||
| 304 |*`CUDA_ERROR_OPERATING_SYSTEM`* |*`hipErrorOperatingSystem`* |
|
||||
| 400 |*`CUDA_ERROR_INVALID_HANDLE`* |*`hipErrorInvalidResourceHandle`* |
|
||||
| 400 |*`CUDA_ERROR_INVALID_HANDLE`* |*`hipErrorInvalidHandle`* |
|
||||
| 401 |*`CUDA_ERROR_ILLEGAL_STATE`* | | 10.0 |
|
||||
| 500 |*`CUDA_ERROR_NOT_FOUND`* |*`hipErrorNotFound`* |
|
||||
| 600 |*`CUDA_ERROR_NOT_READY`* |*`hipErrorNotReady`* |
|
||||
@@ -426,7 +426,7 @@
|
||||
| 703 |*`CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING`* | |
|
||||
| 704 |*`CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED`* |*`hipErrorPeerAccessAlreadyEnabled`* |
|
||||
| 705 |*`CUDA_ERROR_PEER_ACCESS_NOT_ENABLED`* |*`hipErrorPeerAccessNotEnabled`* |
|
||||
| 708 |*`CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE`* | |
|
||||
| 708 |*`CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE`* |*`hipErrorSetOnActiveProcess`* |
|
||||
| 709 |*`CUDA_ERROR_CONTEXT_IS_DESTROYED`* | |
|
||||
| 710 |*`CUDA_ERROR_ASSERT`* |*`hipErrorAssert`* |
|
||||
| 711 |*`CUDA_ERROR_TOO_MANY_PEERS`* | |
|
||||
@@ -439,7 +439,7 @@
|
||||
| 718 |*`CUDA_ERROR_INVALID_PC`* | |
|
||||
| 719 |*`CUDA_ERROR_LAUNCH_FAILED`* |*`hipErrorLaunchFailure`* |
|
||||
| 800 |*`CUDA_ERROR_NOT_PERMITTED`* | |
|
||||
| 801 |*`CUDA_ERROR_NOT_SUPPORTED`* | |
|
||||
| 801 |*`CUDA_ERROR_NOT_SUPPORTED`* |*`hipErrorNotSupported`* |
|
||||
| 802 |*`CUDA_ERROR_SYSTEM_NOT_READY`* | | 10.0 |
|
||||
| 803 |*`CUDA_ERROR_SYSTEM_DRIVER_MISMATCH`* | | 10.1 |
|
||||
| 804 |*`CUDA_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE`* | | 10.1 |
|
||||
@@ -452,7 +452,7 @@
|
||||
| 906 |*`CUDA_ERROR_STREAM_CAPTURE_IMPLICIT`* | | 10.0 |
|
||||
| 907 |*`CUDA_ERROR_CAPTURED_EVENT`* | | 10.0 |
|
||||
| 908 |*`CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD`* | | 10.1 |
|
||||
| 999 |*`CUDA_ERROR_UNKNOWN`* | |
|
||||
| 999 |*`CUDA_ERROR_UNKNOWN`* |*`hipErrorUnknown`* |
|
||||
| enum |***`CUsharedconfig`*** |***`hipSharedMemConfig`*** |
|
||||
| typedef |***`CUsharedconfig_enum`*** |***`hipSharedMemConfig`*** |
|
||||
| 0x00 |*`CU_SHARED_MEM_CONFIG_DEFAULT_BANK_SIZE`* |*`hipSharedMemBankSizeDefault`* |
|
||||
|
||||
@@ -700,21 +700,17 @@
|
||||
| enum |***`cudaError`*** | |***`hipError_t`*** |
|
||||
| typedef |***`cudaError_t`*** | |***`hipError_t`*** |
|
||||
| 0 |*`cudaSuccess`* | |*`hipSuccess`* |
|
||||
| 1 |*`cudaErrorMissingConfiguration`* | |*`hipErrorMissingConfiguration`* | 1001 |
|
||||
| 1 |*`cudaErrorInvalidValue`* | |*`hipErrorInvalidValue`* | 1011 |
|
||||
| 2 |*`cudaErrorMemoryAllocation`* | |*`hipErrorMemoryAllocation`* | 1002 |
|
||||
| 3 |*`cudaErrorInitializationError`* | |*`hipErrorInitializationError`* | 1003 |
|
||||
| 4 |*`cudaErrorLaunchFailure`* | |*`hipErrorLaunchFailure`* | 1004 |
|
||||
| 5 |*`cudaErrorPriorLaunchFailure`* | |*`hipErrorPriorLaunchFailure`* | 1005 |
|
||||
| 6 |*`cudaErrorLaunchTimeout`* | |*`hipErrorLaunchTimeOut`* | 1006 |
|
||||
| 7 |*`cudaErrorLaunchOutOfResources`* | |*`hipErrorLaunchOutOfResources`* | 1007 |
|
||||
| 8 |*`cudaErrorInvalidDeviceFunction`* | |*`hipErrorInvalidDeviceFunction`* | 1008 |
|
||||
| 4 |*`cudaErrorCudartUnloading`* | |*`hipErrorDeinitialized`* |
|
||||
| 5 |*`cudaErrorProfilerDisabled`* | |*`hipErrorProfilerDisabled`* |
|
||||
| 6 |*`cudaErrorProfilerNotInitialized`* | |*`hipErrorProfilerNotInitialized`* |
|
||||
| 7 |*`cudaErrorProfilerAlreadyStarted`* | |*`hipErrorProfilerAlreadyStarted`* |
|
||||
| 8 |*`cudaErrorProfilerAlreadyStopped`* | |*`hipErrorProfilerAlreadyStopped`* |
|
||||
| 9 |*`cudaErrorInvalidConfiguration`* | |*`hipErrorInvalidConfiguration`* | 1009 |
|
||||
| 10 |*`cudaErrorInvalidDevice`* | |*`hipErrorInvalidDevice`* | 1010 |
|
||||
| 11 |*`cudaErrorInvalidValue`* | |*`hipErrorInvalidValue`* | 1011 |
|
||||
| 12 |*`cudaErrorInvalidPitchValue`* | | |
|
||||
| 13 |*`cudaErrorInvalidSymbol`* | |*`hipErrorInvalidSymbol`* | 701 |
|
||||
| 14 |*`cudaErrorMapBufferObjectFailed`* | |*`hipErrorMapFailed`* | 205 |
|
||||
| 15 |*`cudaErrorUnmapBufferObjectFailed`* | |*`hipErrorUnmapFailed`* | 206 |
|
||||
| 16 |*`cudaErrorInvalidHostPointer`* | | |
|
||||
| 17 |*`cudaErrorInvalidDevicePointer`* | |*`hipErrorInvalidDevicePointer`* | 1017 |
|
||||
| 18 |*`cudaErrorInvalidTexture`* | | |
|
||||
@@ -728,63 +724,70 @@
|
||||
| 26 |*`cudaErrorInvalidFilterSetting`* | | |
|
||||
| 27 |*`cudaErrorInvalidNormSetting`* | | |
|
||||
| 28 |*`cudaErrorMixedDeviceExecution`* | | |
|
||||
| 29 |*`cudaErrorCudartUnloading`* | |*`hipErrorDeinitialized`* | 4 |
|
||||
| 30 |*`cudaErrorUnknown`* | |*`hipErrorUnknown`* | 1030 |
|
||||
| 31 |*`cudaErrorNotYetImplemented`* | | |
|
||||
| 32 |*`cudaErrorMemoryValueTooLarge`* | | |
|
||||
| 33 |*`cudaErrorInvalidResourceHandle`* | |*`hipErrorInvalidResourceHandle`* | 1033 |
|
||||
| 34 |*`cudaErrorNotReady`* | |*`hipErrorNotReady`* | 1034 |
|
||||
| 35 |*`cudaErrorInsufficientDriver`* | |*`hipErrorInsufficientDriver`* |
|
||||
| 36 |*`cudaErrorSetOnActiveProcess`* | | |
|
||||
| 37 |*`cudaErrorInvalidSurface`* | | |
|
||||
| 38 |*`cudaErrorNoDevice`* | |*`hipErrorNoDevice`* | 1038 |
|
||||
| 39 |*`cudaErrorECCUncorrectable`* | |*`hipErrorECCNotCorrectable`* | 214 |
|
||||
| 40 |*`cudaErrorSharedObjectSymbolNotFound`* | |*`hipErrorSharedObjectSymbolNotFound`* | 302 |
|
||||
| 41 |*`cudaErrorSharedObjectInitFailed`* | |*`hipErrorSharedObjectInitFailed`* | 303 |
|
||||
| 42 |*`cudaErrorUnsupportedLimit`* | |*`hipErrorUnsupportedLimit`* | 215 |
|
||||
| 43 |*`cudaErrorDuplicateVariableName`* | | |
|
||||
| 44 |*`cudaErrorDuplicateTextureName`* | | |
|
||||
| 45 |*`cudaErrorDuplicateSurfaceName`* | | |
|
||||
| 46 |*`cudaErrorDevicesUnavailable`* | | |
|
||||
| 47 |*`cudaErrorInvalidKernelImage`* | |*`hipErrorInvalidImage`* | 200 |
|
||||
| 48 |*`cudaErrorNoKernelImageForDevice`* | |*`hipErrorNoBinaryForGpu`* | 209 |
|
||||
| 49 |*`cudaErrorIncompatibleDriverContext`* | | |
|
||||
| 50 |*`cudaErrorPeerAccessAlreadyEnabled`* | |*`hipErrorPeerAccessAlreadyEnabled`* | 1050 |
|
||||
| 51 |*`cudaErrorPeerAccessNotEnabled`* | |*`hipErrorPeerAccessNotEnabled`* | 1051 |
|
||||
| 52 | | | |
|
||||
| 53 | | | |
|
||||
| 54 |*`cudaErrorDeviceAlreadyInUse`* | | |
|
||||
| 55 |*`cudaErrorProfilerDisabled`* | |*`hipErrorProfilerDisabled`* | 5 |
|
||||
| 56 |*`cudaErrorProfilerNotInitialized`* | |*`hipErrorProfilerNotInitialized`* | 6 |
|
||||
| 57 |*`cudaErrorProfilerAlreadyStarted`* | |*`hipErrorProfilerAlreadyStarted`* | 7 |
|
||||
| 58 |*`cudaErrorProfilerAlreadyStopped`* | |*`hipErrorProfilerAlreadyStopped`* | 8 |
|
||||
| 59 |*`cudaErrorAssert`* | |*`hipErrorAssert`* | 1081 |
|
||||
| 60 |*`cudaErrorTooManyPeers`* | | |
|
||||
| 61 |*`cudaErrorHostMemoryAlreadyRegistered`* | |*`hipErrorHostMemoryAlreadyRegistered`* | 1061 |
|
||||
| 62 |*`cudaErrorHostMemoryNotRegistered`* | |*`hipErrorHostMemoryNotRegistered`* | 1062 |
|
||||
| 63 |*`cudaErrorOperatingSystem`* | |*`hipErrorOperatingSystem`* | 304 |
|
||||
| 64 |*`cudaErrorPeerAccessUnsupported`* | |*`hipErrorPeerAccessUnsupported`* | 217 |
|
||||
| 52 |*`cudaErrorMissingConfiguration`* | |*`hipErrorMissingConfiguration`* | 1001 |
|
||||
| 53 |*`cudaErrorPriorLaunchFailure`* | |*`hipErrorPriorLaunchFailure`* | 1005 |
|
||||
| 65 |*`cudaErrorLaunchMaxDepthExceeded`* | | |
|
||||
| 66 |*`cudaErrorLaunchFileScopedTex`* | | |
|
||||
| 67 |*`cudaErrorLaunchFileScopedSurf`* | | |
|
||||
| 68 |*`cudaErrorSyncDepthExceeded`* | | |
|
||||
| 69 |*`cudaErrorLaunchPendingCountExceeded`* | | |
|
||||
| 70 |*`cudaErrorNotPermitted`* | | |
|
||||
| 71 |*`cudaErrorNotSupported`* | | |
|
||||
| 72 |*`cudaErrorHardwareStackError`* | | |
|
||||
| 73 |*`cudaErrorIllegalInstruction`* | | |
|
||||
| 74 |*`cudaErrorMisalignedAddress`* | | |
|
||||
| 75 |*`cudaErrorInvalidAddressSpace`* | | |
|
||||
| 76 |*`cudaErrorInvalidPc`* | | |
|
||||
| 77 |*`cudaErrorIllegalAddress`* | |*`hipErrorIllegalAddress`* | 700 |
|
||||
| 78 |*`cudaErrorInvalidPtx`* | |*`hipErrorInvalidKernelFile`* | 218 |
|
||||
| 79 |*`cudaErrorInvalidGraphicsContext`* | |*`hipErrorInvalidGraphicsContext`* | 219 |
|
||||
| 80 |*`cudaErrorNvlinkUncorrectable`* | 8.0 | |
|
||||
| 81 |*`cudaErrorJitCompilerNotFound`* | 9.0 | |
|
||||
| 82 |*`cudaErrorCooperativeLaunchTooLarge`* | 9.0 | |
|
||||
| 83 |*`cudaErrorSystemNotReady`* | 10.0 | |
|
||||
| 84 |*`cudaErrorIllegalState`* | 10.0 | |
|
||||
| 98 |*`cudaErrorInvalidDeviceFunction`* | |*`hipErrorInvalidDeviceFunction`* | 1008 |
|
||||
| 100 |*`cudaErrorNoDevice`* | |*`hipErrorNoDevice`* | 1038 |
|
||||
| 101 |*`cudaErrorInvalidDevice`* | |*`hipErrorInvalidDevice`* | 1010 |
|
||||
| 127 |*`cudaErrorStartupFailure`* | 10.0 | |
|
||||
| 200 |*`cudaErrorInvalidKernelImage`* | |*`hipErrorInvalidImage`* |
|
||||
| 201 |*`cudaErrorDeviceUninitilialized`* | |*`hipErrorInvalidContext`* |
|
||||
| 205 |*`cudaErrorMapBufferObjectFailed`* | |*`hipErrorMapBufferObjectFailed`* | 1071 |
|
||||
| 206 |*`cudaErrorUnmapBufferObjectFailed`* | |*`hipErrorUnmapFailed`* |
|
||||
| 209 |*`cudaErrorNoKernelImageForDevice`* | |*`hipErrorNoBinaryForGpu`* |
|
||||
| 214 |*`cudaErrorECCUncorrectable`* | |*`hipErrorECCNotCorrectable`* |
|
||||
| 215 |*`cudaErrorUnsupportedLimit`* | |*`hipErrorUnsupportedLimit`* |
|
||||
| 216 |*`cudaErrorDeviceAlreadyInUse`* | | |
|
||||
| 217 |*`cudaErrorPeerAccessUnsupported`* | |*`hipErrorPeerAccessUnsupported`* |
|
||||
| 218 |*`cudaErrorInvalidPtx`* | |*`hipErrorInvalidKernelFile`* |
|
||||
| 219 |*`cudaErrorInvalidGraphicsContext`* | |*`hipErrorInvalidGraphicsContext`* |
|
||||
| 220 |*`cudaErrorNvlinkUncorrectable`* | 8.0 | |
|
||||
| 221 |*`cudaErrorJitCompilerNotFound`* | 9.0 | |
|
||||
| 300 |*`cudaErrorInvalidSource`* | 10.1 |*`hipErrorInvalidSource`* |
|
||||
| 301 |*`cudaErrorFileNotFound`* | 10.1 |*`hipErrorFileNotFound`* |
|
||||
| 302 |*`cudaErrorSharedObjectSymbolNotFound`* | |*`hipErrorSharedObjectSymbolNotFound`* |
|
||||
| 303 |*`cudaErrorSharedObjectInitFailed`* | |*`hipErrorSharedObjectInitFailed`* |
|
||||
| 304 |*`cudaErrorOperatingSystem`* | |*`hipErrorOperatingSystem`* |
|
||||
| 400 |*`cudaErrorInvalidResourceHandle`* | |*`hipErrorInvalidResourceHandle`* | 1033 |
|
||||
| 401 |*`cudaErrorIllegalState`* | 10.0 | |
|
||||
| 500 |*`cudaErrorSymbolNotFound`* | 10.1 |*`hipErrorNotFound`* |
|
||||
| 600 |*`cudaErrorNotReady`* | |*`hipErrorNotReady`* | 1034 |
|
||||
| 700 |*`cudaErrorIllegalAddress`* | |*`hipErrorIllegalAddress`* |
|
||||
| 701 |*`cudaErrorLaunchOutOfResources`* | |*`hipErrorLaunchOutOfResources`* | 1007 |
|
||||
| 702 |*`cudaErrorLaunchTimeout`* | |*`hipErrorLaunchTimeOut`* | 1006 |
|
||||
| 703 |*`cudaErrorLaunchIncompatibleTexturing`* | | |
|
||||
| 704 |*`cudaErrorPeerAccessAlreadyEnabled`* | |*`hipErrorPeerAccessAlreadyEnabled`* | 1050 |
|
||||
| 705 |*`cudaErrorPeerAccessNotEnabled`* | |*`hipErrorPeerAccessNotEnabled`* | 1051 |
|
||||
| 708 |*`cudaErrorSetOnActiveProcess`* | |*`hipErrorSetOnActiveProcess`* | 305 |
|
||||
| 709 |*`cudaErrorContextIsDestroyed`* | | |
|
||||
| 710 |*`cudaErrorAssert`* | |*`hipErrorAssert`* | 1081 |
|
||||
| 711 |*`cudaErrorTooManyPeers`* | | |
|
||||
| 712 |*`cudaErrorHostMemoryAlreadyRegistered`* | |*`hipErrorHostMemoryAlreadyRegistered`* | 1061 |
|
||||
| 713 |*`cudaErrorHostMemoryNotRegistered`* | |*`hipErrorHostMemoryNotRegistered`* | 1062 |
|
||||
| 714 |*`cudaErrorHardwareStackError`* | | |
|
||||
| 715 |*`cudaErrorIllegalInstruction`* | | |
|
||||
| 716 |*`cudaErrorMisalignedAddress`* | | |
|
||||
| 717 |*`cudaErrorInvalidAddressSpace`* | | |
|
||||
| 718 |*`cudaErrorInvalidPc`* | | |
|
||||
| 719 |*`cudaErrorLaunchFailure`* | |*`hipErrorLaunchFailure`* | 1004 |
|
||||
| 720 |*`cudaErrorCooperativeLaunchTooLarge`* | 9.0 | |
|
||||
| 800 |*`cudaErrorNotPermitted`* | | |
|
||||
| 801 |*`cudaErrorNotSupported`* | |*`hipErrorNotSupported`* | 1082 |
|
||||
| 802 |*`cudaErrorSystemNotReady`* | 10.0 | |
|
||||
| 803 |*`cudaErrorSystemDriverMismatch`* | 10.0 | |
|
||||
| 804 |*`cudaErrorCompatNotSupportedOnDevice`* | 10.0 | |
|
||||
| 900 |*`cudaErrorStreamCaptureUnsupported`* | 10.0 | |
|
||||
@@ -796,6 +799,7 @@
|
||||
| 906 |*`cudaErrorStreamCaptureImplicit`* | 10.0 | |
|
||||
| 907 |*`cudaErrorCapturedEvent`* | 10.0 | |
|
||||
| 908 |*`cudaErrorStreamCaptureWrongThread`* | 10.1 | |
|
||||
| 999 |*`cudaErrorUnknown`* | |*`hipErrorUnknown`* | 1030 |
|
||||
| 10000 |*`cudaErrorApiFailureBase`* | | |
|
||||
| enum |***`cudaFuncCache`*** | |***`hipFuncCache_t`*** |
|
||||
| 0 |*`cudaFuncCachePreferNone`* | |*`hipFuncCachePreferNone`* |
|
||||
|
||||
@@ -30,7 +30,7 @@ cl::opt<std::string> OutputFilename("o",
|
||||
cl::cat(ToolTemplateCategory));
|
||||
|
||||
cl::opt<std::string> OutputDir("o-dir",
|
||||
cl::desc("Output direcory"),
|
||||
cl::desc("Output directory"),
|
||||
cl::value_desc("directory"),
|
||||
cl::cat(ToolTemplateCategory));
|
||||
|
||||
@@ -45,17 +45,17 @@ cl::opt<bool> GeneratePython("python",
|
||||
cl::cat(ToolTemplateCategory));
|
||||
|
||||
cl::opt<std::string> OutputHipifyPerlDir("o-hipify-perl-dir",
|
||||
cl::desc("Output direcory for hipify-perl script"),
|
||||
cl::desc("Output directory for hipify-perl script"),
|
||||
cl::value_desc("directory"),
|
||||
cl::cat(ToolTemplateCategory));
|
||||
|
||||
cl::opt<std::string> OutputPythonMapDir("o-python-map-dir",
|
||||
cl::desc("Output direcory for Python map"),
|
||||
cl::desc("Output directory for Python map"),
|
||||
cl::value_desc("directory"),
|
||||
cl::cat(ToolTemplateCategory));
|
||||
|
||||
cl::opt<std::string> TemporaryDir("temp-dir",
|
||||
cl::desc("Temporary direcory"),
|
||||
cl::desc("Temporary directory"),
|
||||
cl::value_desc("directory"),
|
||||
cl::cat(ToolTemplateCategory));
|
||||
|
||||
|
||||
@@ -1105,161 +1105,195 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
|
||||
{"CUresult", {"hipError_t", "", CONV_TYPE, API_DRIVER}},
|
||||
{"cudaError_enum", {"hipError_t", "", CONV_TYPE, API_DRIVER}},
|
||||
// CUresult enum values
|
||||
// cudaSuccess = 0
|
||||
// cudaSuccess
|
||||
{"CUDA_SUCCESS", {"hipSuccess", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0
|
||||
// cudaErrorInvalidValue = 11
|
||||
// cudaErrorInvalidValue = 1, CUDA_ERROR_INVALID_VALUE = 1, hipErrorInvalidValue = 1011
|
||||
// TODO [HIP]: make hipErrorInvalidValue = 1
|
||||
// cudaErrorInvalidValue
|
||||
{"CUDA_ERROR_INVALID_VALUE", {"hipErrorInvalidValue", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 1
|
||||
// cudaErrorMemoryAllocation = 2
|
||||
{"CUDA_ERROR_OUT_OF_MEMORY", {"hipErrorMemoryAllocation", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 2
|
||||
// cudaErrorInitializationError = 3; hipErrorInitializationError = 1003
|
||||
// TODO: the same error code should be merged into 1
|
||||
// cudaErrorMemoryAllocation = 2, CUDA_ERROR_OUT_OF_MEMORY = 2, hipErrorOutOfMemory = 2, hipErrorMemoryAllocation = 1002
|
||||
// TODO [HIP]: remove hipErrorMemoryAllocation
|
||||
// cudaErrorMemoryAllocation
|
||||
{"CUDA_ERROR_OUT_OF_MEMORY", {"hipErrorOutOfMemory", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 2
|
||||
// cudaErrorInitializationError = 3, CUDA_ERROR_NOT_INITIALIZED = 3, hipErrorNotInitialized = 3, hipErrorInitializationError = 1003
|
||||
// TODO [HIP]: remove hipErrorInitializationError
|
||||
// cudaErrorInitializationError
|
||||
{"CUDA_ERROR_NOT_INITIALIZED", {"hipErrorNotInitialized", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 3
|
||||
// cudaErrorCudartUnloading = 29
|
||||
// TODO: double check, that these errors match
|
||||
// cudaErrorCudartUnloading
|
||||
{"CUDA_ERROR_DEINITIALIZED", {"hipErrorDeinitialized", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 4
|
||||
// cudaErrorProfilerDisabled = 55
|
||||
// cudaErrorProfilerDisabled
|
||||
{"CUDA_ERROR_PROFILER_DISABLED", {"hipErrorProfilerDisabled", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 5
|
||||
// cudaErrorProfilerNotInitialized = 56
|
||||
// cudaErrorProfilerNotInitialized
|
||||
// NOTE: Deprecated since CUDA 5.0
|
||||
{"CUDA_ERROR_PROFILER_NOT_INITIALIZED", {"hipErrorProfilerNotInitialized", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 6
|
||||
// NOTE: Deprecated since CUDA 5.0
|
||||
// cudaErrorProfilerAlreadyStarted = 57
|
||||
// cudaErrorProfilerAlreadyStarted
|
||||
{"CUDA_ERROR_PROFILER_ALREADY_STARTED", {"hipErrorProfilerAlreadyStarted", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 7
|
||||
// cudaErrorProfilerAlreadyStopped = 58
|
||||
// cudaErrorProfilerAlreadyStopped
|
||||
// NOTE: Deprecated since CUDA 5.0
|
||||
{"CUDA_ERROR_PROFILER_ALREADY_STOPPED", {"hipErrorProfilerAlreadyStopped", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 8
|
||||
// cudaErrorNoDevice = 38
|
||||
// cudaErrorNoDevice = 100, CUDA_ERROR_NO_DEVICE = 100, hipErrorNoDevice = 1038
|
||||
// TODO [HIP]: make hipErrorNoDevice = 100
|
||||
// cudaErrorNoDevice
|
||||
{"CUDA_ERROR_NO_DEVICE", {"hipErrorNoDevice", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 100
|
||||
// cudaErrorInvalidDevice = 10
|
||||
// cudaErrorInvalidDevice = 101, CUDA_ERROR_INVALID_DEVICE = 101, hipErrorInvalidDevice = 1010
|
||||
// TODO [HIP]: make hipErrorInvalidDevice = 101
|
||||
// cudaErrorInvalidDevice
|
||||
{"CUDA_ERROR_INVALID_DEVICE", {"hipErrorInvalidDevice", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 101
|
||||
// cudaErrorInvalidKernelImage = 47
|
||||
// TODO: double check the matching
|
||||
// cudaErrorInvalidKernelImage
|
||||
{"CUDA_ERROR_INVALID_IMAGE", {"hipErrorInvalidImage", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 200
|
||||
// no analogue
|
||||
// cudaErrorDeviceUninitilialized
|
||||
{"CUDA_ERROR_INVALID_CONTEXT", {"hipErrorInvalidContext", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 201
|
||||
// no analogue
|
||||
// NOTE: Deprecated since CUDA 3.2
|
||||
{"CUDA_ERROR_CONTEXT_ALREADY_CURRENT", {"hipErrorContextAlreadyCurrent", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 202
|
||||
// cudaErrorMapBufferObjectFailed = 14
|
||||
// TODO: double check the matching
|
||||
// cudaErrorMapBufferObjectFailed = 205, CUDA_ERROR_MAP_FAILED = 205, hipErrorMapFailed = 205, hipErrorMapBufferObjectFailed = 1071
|
||||
// TODO [HIP]: remove hipErrorMapBufferObjectFailed
|
||||
// TODO [HIPIFY]: rename hipErrorMapBufferObjectFailed to hipErrorMapFailed
|
||||
// cudaErrorMapBufferObjectFailed
|
||||
{"CUDA_ERROR_MAP_FAILED", {"hipErrorMapFailed", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 205
|
||||
// cudaErrorUnmapBufferObjectFailed = 15
|
||||
// TODO: double check the matching
|
||||
// cudaErrorUnmapBufferObjectFailed
|
||||
{"CUDA_ERROR_UNMAP_FAILED", {"hipErrorUnmapFailed", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 206
|
||||
// no analogue
|
||||
// cudaErrorArrayIsMapped
|
||||
{"CUDA_ERROR_ARRAY_IS_MAPPED", {"hipErrorArrayIsMapped", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 207
|
||||
// no analogue
|
||||
// cudaErrorAlreadyMapped
|
||||
{"CUDA_ERROR_ALREADY_MAPPED", {"hipErrorAlreadyMapped", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 208
|
||||
// cudaErrorNoKernelImageForDevice = 48
|
||||
// TODO: double check the matching
|
||||
// cudaErrorNoKernelImageForDevice
|
||||
{"CUDA_ERROR_NO_BINARY_FOR_GPU", {"hipErrorNoBinaryForGpu", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 209
|
||||
// no analogue
|
||||
// cudaErrorAlreadyAcquired
|
||||
{"CUDA_ERROR_ALREADY_ACQUIRED", {"hipErrorAlreadyAcquired", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 210
|
||||
// no analogue
|
||||
// cudaErrorNotMapped
|
||||
{"CUDA_ERROR_NOT_MAPPED", {"hipErrorNotMapped", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 211
|
||||
// no analogue
|
||||
// cudaErrorNotMappedAsArray
|
||||
{"CUDA_ERROR_NOT_MAPPED_AS_ARRAY", {"hipErrorNotMappedAsArray", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 212
|
||||
// no analogue
|
||||
// cudaErrorNotMappedAsPointer
|
||||
{"CUDA_ERROR_NOT_MAPPED_AS_POINTER", {"hipErrorNotMappedAsPointer", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 213
|
||||
// cudaErrorECCUncorrectable = 39
|
||||
// cudaErrorECCUncorrectable
|
||||
{"CUDA_ERROR_ECC_UNCORRECTABLE", {"hipErrorECCNotCorrectable", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 214
|
||||
// cudaErrorUnsupportedLimit = 42
|
||||
// cudaErrorUnsupportedLimit
|
||||
{"CUDA_ERROR_UNSUPPORTED_LIMIT", {"hipErrorUnsupportedLimit", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 215
|
||||
// no analogue
|
||||
// cudaErrorDeviceAlreadyInUse
|
||||
{"CUDA_ERROR_CONTEXT_ALREADY_IN_USE", {"hipErrorContextAlreadyInUse", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 216
|
||||
// cudaErrorPeerAccessUnsupported = 64
|
||||
// cudaErrorPeerAccessUnsupported
|
||||
{"CUDA_ERROR_PEER_ACCESS_UNSUPPORTED", {"hipErrorPeerAccessUnsupported", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 217
|
||||
// cudaErrorInvalidPtx = 78
|
||||
// cudaErrorInvalidPtx
|
||||
{"CUDA_ERROR_INVALID_PTX", {"hipErrorInvalidKernelFile", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 218
|
||||
// cudaErrorInvalidGraphicsContext = 79
|
||||
// cudaErrorInvalidGraphicsContext
|
||||
{"CUDA_ERROR_INVALID_GRAPHICS_CONTEXT", {"hipErrorInvalidGraphicsContext", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 219
|
||||
// cudaErrorNvlinkUncorrectable = 80
|
||||
// cudaErrorNvlinkUncorrectable
|
||||
{"CUDA_ERROR_NVLINK_UNCORRECTABLE", {"hipErrorNvlinkUncorrectable", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 220
|
||||
// cudaErrorJitCompilerNotFound = 81
|
||||
// cudaErrorJitCompilerNotFound
|
||||
{"CUDA_ERROR_JIT_COMPILER_NOT_FOUND", {"hipErrorJitCompilerNotFound", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 221
|
||||
// no analogue
|
||||
// cudaErrorInvalidSource
|
||||
{"CUDA_ERROR_INVALID_SOURCE", {"hipErrorInvalidSource", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 300
|
||||
// no analogue
|
||||
// cudaErrorFileNotFound
|
||||
{"CUDA_ERROR_FILE_NOT_FOUND", {"hipErrorFileNotFound", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 301
|
||||
// cudaErrorSharedObjectSymbolNotFound = 40
|
||||
// cudaErrorSharedObjectSymbolNotFound
|
||||
{"CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND", {"hipErrorSharedObjectSymbolNotFound", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 302
|
||||
// cudaErrorSharedObjectInitFailed = 41
|
||||
// cudaErrorSharedObjectInitFailed
|
||||
{"CUDA_ERROR_SHARED_OBJECT_INIT_FAILED", {"hipErrorSharedObjectInitFailed", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 303
|
||||
// cudaErrorOperatingSystem = 63
|
||||
// cudaErrorOperatingSystem
|
||||
{"CUDA_ERROR_OPERATING_SYSTEM", {"hipErrorOperatingSystem", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 304
|
||||
// cudaErrorInvalidResourceHandle = 33
|
||||
{"CUDA_ERROR_INVALID_HANDLE", {"hipErrorInvalidResourceHandle", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 400
|
||||
// cudaErrorIllegalState = 84
|
||||
// cudaErrorInvalidResourceHandle = 400, CUDA_ERROR_INVALID_HANDLE = 400, hipErrorInvalidHandle = 400, hipErrorInvalidResourceHandle = 1033
|
||||
// TODO [HIP]: remove hipErrorInvalidResourceHandle
|
||||
// TODO [HIPIFY]: rename hipErrorInvalidResourceHandle to hipErrorInvalidHandle
|
||||
// cudaErrorInvalidResourceHandle
|
||||
{"CUDA_ERROR_INVALID_HANDLE", {"hipErrorInvalidHandle", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 400
|
||||
// cudaErrorIllegalState
|
||||
{"CUDA_ERROR_ILLEGAL_STATE", {"hipErrorIllegalState", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 401
|
||||
// no analogue
|
||||
// cudaErrorSymbolNotFound
|
||||
{"CUDA_ERROR_NOT_FOUND", {"hipErrorNotFound", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 500
|
||||
// cudaErrorNotReady = 34
|
||||
// cudaErrorNotReady = 600, CUDA_ERROR_NOT_READY = 600, hipErrorNotReady = 1034
|
||||
// TODO [HIP]: make hipErrorNotReady = 600
|
||||
// cudaErrorNotReady
|
||||
{"CUDA_ERROR_NOT_READY", {"hipErrorNotReady", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 600
|
||||
// cudaErrorIllegalAddress = 77
|
||||
// cudaErrorIllegalAddress
|
||||
{"CUDA_ERROR_ILLEGAL_ADDRESS", {"hipErrorIllegalAddress", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 700
|
||||
// cudaErrorLaunchOutOfResources = 7
|
||||
// cudaErrorLaunchOutOfResources = 701, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701, hipErrorLaunchOutOfResources = 1007
|
||||
// TODO [HIP]: make hipErrorLaunchOutOfResources = 701
|
||||
// cudaErrorLaunchOutOfResources
|
||||
{"CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES", {"hipErrorLaunchOutOfResources", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 701
|
||||
// cudaErrorLaunchTimeout = 6
|
||||
// cudaErrorLaunchTimeout = 702, CUDA_ERROR_LAUNCH_TIMEOUT = 702, hipErrorLaunchTimeOut = 1006
|
||||
// TODO [HIP]: make hipErrorLaunchTimeOut = 702
|
||||
// cudaErrorLaunchTimeout
|
||||
{"CUDA_ERROR_LAUNCH_TIMEOUT", {"hipErrorLaunchTimeOut", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 702
|
||||
// no analogue
|
||||
// cudaErrorLaunchIncompatibleTexturing
|
||||
{"CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING", {"hipErrorLaunchIncompatibleTexturing", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 703
|
||||
// cudaErrorPeerAccessAlreadyEnabled = 50
|
||||
// cudaErrorPeerAccessAlreadyEnabled = 704, CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED = 704, hipErrorPeerAccessAlreadyEnabled = 1050
|
||||
// TODO [HIP]: make hipErrorPeerAccessAlreadyEnabled = 704
|
||||
// cudaErrorPeerAccessAlreadyEnabled
|
||||
{"CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED", {"hipErrorPeerAccessAlreadyEnabled", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 704
|
||||
// cudaErrorPeerAccessNotEnabled = 51
|
||||
// cudaErrorPeerAccessNotEnabled = 705, CUDA_ERROR_PEER_ACCESS_NOT_ENABLED = 705, hipErrorPeerAccessNotEnabled = 1051
|
||||
// TODO [HIP]: make hipErrorPeerAccessNotEnabled = 705
|
||||
// cudaErrorPeerAccessNotEnabled
|
||||
{"CUDA_ERROR_PEER_ACCESS_NOT_ENABLED", {"hipErrorPeerAccessNotEnabled", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 705
|
||||
// no analogue
|
||||
{"CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE", {"hipErrorPrimaryContextActive", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 708
|
||||
// no analogue
|
||||
// cudaErrorSetOnActiveProcess = 708, CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE = 708, hipErrorSetOnActiveProcess = 305
|
||||
// TODO [HIP]: make hipErrorSetOnActiveProcess = 708
|
||||
// cudaErrorSetOnActiveProcess
|
||||
{"CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE", {"hipErrorSetOnActiveProcess", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 708
|
||||
// cudaErrorContextIsDestroyed
|
||||
{"CUDA_ERROR_CONTEXT_IS_DESTROYED", {"hipErrorContextIsDestroyed", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 709
|
||||
// cudaErrorAssert = 59
|
||||
// cudaErrorAssert = 710, CUDA_ERROR_ASSERT = 710, hipErrorAssert = 1081
|
||||
// TODO [HIP]: make hipErrorAssert = 710
|
||||
// cudaErrorAssert
|
||||
{"CUDA_ERROR_ASSERT", {"hipErrorAssert", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 710
|
||||
// cudaErrorTooManyPeers = 60
|
||||
// cudaErrorTooManyPeers
|
||||
{"CUDA_ERROR_TOO_MANY_PEERS", {"hipErrorTooManyPeers", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 711
|
||||
// = cudaErrorHostMemoryAlreadyRegistered = 61
|
||||
// cudaErrorHostMemoryAlreadyRegistered = 712, CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED = 712, hipErrorHostMemoryAlreadyRegistered = 1061
|
||||
// TODO [HIP]: make hipErrorHostMemoryAlreadyRegistered = 712
|
||||
// cudaErrorHostMemoryAlreadyRegistered
|
||||
{"CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED", {"hipErrorHostMemoryAlreadyRegistered", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 712
|
||||
// cudaErrorHostMemoryNotRegistered = 62
|
||||
// cudaErrorHostMemoryNotRegistered = 713, CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED = 713, hipErrorHostMemoryNotRegistered = 1062
|
||||
// TODO [HIP]: make hipErrorHostMemoryNotRegistered = 713
|
||||
// cudaErrorHostMemoryNotRegistered
|
||||
{"CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED", {"hipErrorHostMemoryNotRegistered", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 713
|
||||
// cudaErrorHardwareStackError= 72
|
||||
// cudaErrorHardwareStackError
|
||||
{"CUDA_ERROR_HARDWARE_STACK_ERROR", {"hipErrorHardwareStackError", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 714
|
||||
// cudaErrorIllegalInstruction = 73
|
||||
// cudaErrorIllegalInstruction
|
||||
{"CUDA_ERROR_ILLEGAL_INSTRUCTION", {"hipErrorIllegalInstruction", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 715
|
||||
// cudaErrorMisalignedAddress= 74
|
||||
// cudaErrorMisalignedAddress
|
||||
{"CUDA_ERROR_MISALIGNED_ADDRESS", {"hipErrorMisalignedAddress", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 716
|
||||
// cudaErrorInvalidAddressSpace= 75
|
||||
// cudaErrorInvalidAddressSpace
|
||||
{"CUDA_ERROR_INVALID_ADDRESS_SPACE", {"hipErrorInvalidAddressSpace", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 717
|
||||
// cudaErrorInvalidPc = 76
|
||||
// cudaErrorInvalidPc
|
||||
{"CUDA_ERROR_INVALID_PC", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 718
|
||||
// cudaErrorLaunchFailure = 4
|
||||
// cudaErrorLaunchFailure = 719, CUDA_ERROR_LAUNCH_FAILED = 719, hipErrorLaunchFailure = 1004
|
||||
// TODO [HIP]: make hipErrorSetOnActiveProcess = 719
|
||||
// cudaErrorLaunchFailure
|
||||
{"CUDA_ERROR_LAUNCH_FAILED", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 719
|
||||
// cudaErrorNotPermitted = 70
|
||||
// cudaErrorNotPermitted
|
||||
{"CUDA_ERROR_NOT_PERMITTED", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 800
|
||||
// cudaErrorNotSupported = 71
|
||||
{"CUDA_ERROR_NOT_SUPPORTED", {"hipErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 801
|
||||
// cudaErrorSystemNotReady = 83
|
||||
// cudaErrorNotSupported = 801, CUDA_ERROR_NOT_SUPPORTED = 801, hipErrorNotSupported = 1082
|
||||
// TODO [HIP]: make hipErrorNotSupported = 801
|
||||
// cudaErrorNotSupported
|
||||
{"CUDA_ERROR_NOT_SUPPORTED", {"hipErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 801
|
||||
// cudaErrorSystemNotReady
|
||||
{"CUDA_ERROR_SYSTEM_NOT_READY", {"hipErrorSystemNotReady", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 802
|
||||
// cudaErrorSystemDriverMismatch = 803
|
||||
// cudaErrorSystemDriverMismatch
|
||||
{"CUDA_ERROR_SYSTEM_DRIVER_MISMATCH", {"hipErrorSystemDriverMismatch", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 803
|
||||
// cudaErrorCompatNotSupportedOnDevice = 804
|
||||
// cudaErrorCompatNotSupportedOnDevice
|
||||
{"CUDA_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE", {"hipErrorCompatNotSupportedOnDevice", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 804
|
||||
// cudaErrorStreamCaptureUnsupported = 900
|
||||
// cudaErrorStreamCaptureUnsupported
|
||||
{"CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED", {"hipErrorStreamCaptureUnsupported", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 900
|
||||
// cudaErrorStreamCaptureInvalidated = 901
|
||||
// cudaErrorStreamCaptureInvalidated
|
||||
{"CUDA_ERROR_STREAM_CAPTURE_INVALIDATED", {"hipErrorStreamCaptureInvalidated", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 901
|
||||
// cudaErrorStreamCaptureMerge = 902
|
||||
// cudaErrorStreamCaptureMerge
|
||||
{"CUDA_ERROR_STREAM_CAPTURE_MERGE", {"hipErrorStreamCaptureMerge", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 902
|
||||
// cudaErrorStreamCaptureUnmatched = 903
|
||||
// cudaErrorStreamCaptureUnmatched
|
||||
{"CUDA_ERROR_STREAM_CAPTURE_UNMATCHED", {"hipErrorStreamCaptureUnmatched", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 903
|
||||
// cudaErrorStreamCaptureUnjoined = 904
|
||||
// cudaErrorStreamCaptureUnjoined
|
||||
{"CUDA_ERROR_STREAM_CAPTURE_UNJOINED", {"hipErrorStreamCaptureUnjoined", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 904
|
||||
// cudaErrorStreamCaptureIsolation = 905
|
||||
// cudaErrorStreamCaptureIsolation
|
||||
{"CUDA_ERROR_STREAM_CAPTURE_ISOLATION", {"hipErrorStreamCaptureIsolation", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 905
|
||||
// cudaErrorStreamCaptureImplicit = 906
|
||||
// cudaErrorStreamCaptureImplicit
|
||||
{"CUDA_ERROR_STREAM_CAPTURE_IMPLICIT", {"hipErrorStreamCaptureImplicit", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 906
|
||||
// cudaErrorCapturedEvent = 907
|
||||
// cudaErrorCapturedEvent
|
||||
{"CUDA_ERROR_CAPTURED_EVENT", {"hipErrorCapturedEvent", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 907
|
||||
// cudaErrorStreamCaptureWrongThread = 908
|
||||
// cudaErrorStreamCaptureWrongThread
|
||||
{"CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD", {"hipErrorStreamCaptureWrongThread", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 908
|
||||
// cudaErrorUnknown = 30
|
||||
// cudaErrorUnknown = 999, CUDA_ERROR_UNKNOWN = 999, hipErrorUnknown = 1030
|
||||
// TODO [HIP]: make hipErrorUnknown = 999
|
||||
// cudaErrorUnknown
|
||||
{"CUDA_ERROR_UNKNOWN", {"hipErrorUnknown", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 999
|
||||
|
||||
// cudaSharedMemConfig
|
||||
|
||||
@@ -595,42 +595,51 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
|
||||
{"cudaError", {"hipError_t", "", CONV_TYPE, API_RUNTIME}},
|
||||
{"cudaError_t", {"hipError_t", "", CONV_TYPE, API_RUNTIME}},
|
||||
// cudaError enum values
|
||||
// CUDA_SUCCESS = 0
|
||||
// CUDA_SUCCESS
|
||||
{"cudaSuccess", {"hipSuccess", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 0
|
||||
// no analogue
|
||||
{"cudaErrorMissingConfiguration", {"hipErrorMissingConfiguration", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 1
|
||||
// CUDA_ERROR_OUT_OF_MEMORY = 2
|
||||
// cudaErrorInvalidValue = 1, CUDA_ERROR_INVALID_VALUE = 1, hipErrorInvalidValue = 1011
|
||||
// TODO [HIP]: make hipErrorInvalidValue = 1
|
||||
// CUDA_ERROR_INVALID_VALUE
|
||||
{"cudaErrorInvalidValue", {"hipErrorInvalidValue", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 1
|
||||
// cudaErrorMemoryAllocation = 2, CUDA_ERROR_OUT_OF_MEMORY = 2, hipErrorOutOfMemory = 2, hipErrorMemoryAllocation = 1002
|
||||
// TODO [HIP]: remove hipErrorMemoryAllocation
|
||||
// TODO [HIPIFY]: rename hipErrorMemoryAllocation to hipErrorOutOfMemory
|
||||
// CUDA_ERROR_OUT_OF_MEMORY
|
||||
{"cudaErrorMemoryAllocation", {"hipErrorMemoryAllocation", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 2
|
||||
// CUDA_ERROR_NOT_INITIALIZED = 3
|
||||
// cudaErrorInitializationError = 3, CUDA_ERROR_NOT_INITIALIZED = 3, hipErrorNotInitialized = 3, hipErrorInitializationError = 1003
|
||||
// TODO [HIP]: remove hipErrorInitializationError
|
||||
// TODO [HIPIFY]: rename hipErrorInitializationError to hipErrorNotInitialized
|
||||
// CUDA_ERROR_NOT_INITIALIZED
|
||||
{"cudaErrorInitializationError", {"hipErrorInitializationError", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 3
|
||||
// CUDA_ERROR_LAUNCH_FAILED = 719
|
||||
{"cudaErrorLaunchFailure", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 4
|
||||
// no analogue
|
||||
{"cudaErrorPriorLaunchFailure", {"hipErrorPriorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 5
|
||||
// CUDA_ERROR_LAUNCH_TIMEOUT = 702
|
||||
{"cudaErrorLaunchTimeout", {"hipErrorLaunchTimeOut", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 6
|
||||
// CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701
|
||||
{"cudaErrorLaunchOutOfResources", {"hipErrorLaunchOutOfResources", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 7
|
||||
// no analogue
|
||||
{"cudaErrorInvalidDeviceFunction", {"hipErrorInvalidDeviceFunction", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 8
|
||||
// CUDA_ERROR_DEINITIALIZED
|
||||
{"cudaErrorCudartUnloading", {"hipErrorDeinitialized", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 4
|
||||
// CUDA_ERROR_PROFILER_DISABLED
|
||||
{"cudaErrorProfilerDisabled", {"hipErrorProfilerDisabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 5
|
||||
// Deprecated since CUDA 5.0
|
||||
// CUDA_ERROR_PROFILER_NOT_INITIALIZED
|
||||
{"cudaErrorProfilerNotInitialized", {"hipErrorProfilerNotInitialized", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 6
|
||||
// Deprecated since CUDA 5.0
|
||||
// CUDA_ERROR_PROFILER_ALREADY_STARTED
|
||||
{"cudaErrorProfilerAlreadyStarted", {"hipErrorProfilerAlreadyStarted", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 7
|
||||
// Deprecated since CUDA 5.0
|
||||
// CUDA_ERROR_PROFILER_ALREADY_STOPPED
|
||||
{"cudaErrorProfilerAlreadyStopped", {"hipErrorProfilerAlreadyStopped", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 8
|
||||
// cudaErrorInvalidConfiguration = 9, hipErrorInvalidConfiguration = 1009
|
||||
// TODO [HIP]: make hipErrorInvalidConfiguration = 9
|
||||
// no analogue
|
||||
{"cudaErrorInvalidConfiguration", {"hipErrorInvalidConfiguration", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 9
|
||||
// CUDA_ERROR_INVALID_DEVICE = 101
|
||||
{"cudaErrorInvalidDevice", {"hipErrorInvalidDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 10
|
||||
// CUDA_ERROR_INVALID_VALUE = 1
|
||||
{"cudaErrorInvalidValue", {"hipErrorInvalidValue", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 11
|
||||
// no analogue
|
||||
{"cudaErrorInvalidPitchValue", {"hipErrorInvalidPitchValue", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 12
|
||||
// cudaErrorInvalidSymbol = 13, hipErrorInvalidSymbol = 701, but also cudaErrorLaunchOutOfResources = CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701
|
||||
// TODO [HIP]: make hipErrorInvalidSymbol = 13
|
||||
// no analogue
|
||||
{"cudaErrorInvalidSymbol", {"hipErrorInvalidSymbol", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 13
|
||||
// CUDA_ERROR_MAP_FAILED = 205
|
||||
// TODO: double check the matching
|
||||
{"cudaErrorMapBufferObjectFailed", {"hipErrorMapFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 14
|
||||
// CUDA_ERROR_UNMAP_FAILED = 206
|
||||
// TODO: double check the matching
|
||||
{"cudaErrorUnmapBufferObjectFailed", {"hipErrorUnmapFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 15
|
||||
// Deprecated since CUDA 10.1
|
||||
// no analogue
|
||||
{"cudaErrorInvalidHostPointer", {"hipErrorInvalidHostPointer", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 16
|
||||
// cudaErrorInvalidDevicePointer = 17, hipErrorInvalidDevicePointer = 1017
|
||||
// TODO [HIP]: make hipErrorInvalidDevicePointer = 17
|
||||
// Deprecated since CUDA 10.1
|
||||
// no analogue
|
||||
{"cudaErrorInvalidDevicePointer", {"hipErrorInvalidDevicePointer", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 17
|
||||
// no analogue
|
||||
@@ -639,6 +648,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
|
||||
{"cudaErrorInvalidTextureBinding", {"hipErrorInvalidTextureBinding", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 19
|
||||
// no analogue
|
||||
{"cudaErrorInvalidChannelDescriptor", {"hipErrorInvalidChannelDescriptor", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 20
|
||||
// cudaErrorInvalidMemcpyDirection = 21, hipErrorInvalidMemcpyDirection = 1021
|
||||
// TODO [HIP]: make hipErrorInvalidMemcpyDirection = 21
|
||||
// no analogue
|
||||
{"cudaErrorInvalidMemcpyDirection", {"hipErrorInvalidMemcpyDirection", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 21
|
||||
// no analogue
|
||||
@@ -655,36 +666,16 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
|
||||
{"cudaErrorInvalidNormSetting", {"hipErrorInvalidNormSetting", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 27
|
||||
// no analogue
|
||||
{"cudaErrorMixedDeviceExecution", {"hipErrorMixedDeviceExecution", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 28
|
||||
// CUDA_ERROR_DEINITIALIZED = 4
|
||||
{"cudaErrorCudartUnloading", {"hipErrorDeinitialized", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 29
|
||||
// CUDA_ERROR_UNKNOWN = 999
|
||||
{"cudaErrorUnknown", {"hipErrorUnknown", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 30
|
||||
// Deprecated since CUDA 4.1
|
||||
// no analogue
|
||||
{"cudaErrorNotYetImplemented", {"hipErrorNotYetImplemented", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 31
|
||||
// Deprecated since CUDA 3.1
|
||||
// no analogue
|
||||
{"cudaErrorMemoryValueTooLarge", {"hipErrorMemoryValueTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 32
|
||||
// CUDA_ERROR_INVALID_HANDLE = 400
|
||||
{"cudaErrorInvalidResourceHandle", {"hipErrorInvalidResourceHandle", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 33
|
||||
// CUDA_ERROR_NOT_READY = 600
|
||||
{"cudaErrorNotReady", {"hipErrorNotReady", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 34
|
||||
// no analogue
|
||||
{"cudaErrorInsufficientDriver", {"hipErrorInsufficientDriver", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 35
|
||||
// no analogue
|
||||
{"cudaErrorSetOnActiveProcess", {"hipErrorSetOnActiveProcess", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 36
|
||||
// no analogue
|
||||
{"cudaErrorInvalidSurface", {"hipErrorInvalidSurface", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 37
|
||||
// CUDA_ERROR_NO_DEVICE = 100
|
||||
{"cudaErrorNoDevice", {"hipErrorNoDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 38
|
||||
// CUDA_ERROR_ECC_UNCORRECTABLE = 214
|
||||
{"cudaErrorECCUncorrectable", {"hipErrorECCNotCorrectable", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 39
|
||||
// CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND = 302
|
||||
{"cudaErrorSharedObjectSymbolNotFound", {"hipErrorSharedObjectSymbolNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 40
|
||||
// CUDA_ERROR_SHARED_OBJECT_INIT_FAILED = 303
|
||||
{"cudaErrorSharedObjectInitFailed", {"hipErrorSharedObjectInitFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 41
|
||||
// CUDA_ERROR_UNSUPPORTED_LIMIT = 215
|
||||
{"cudaErrorUnsupportedLimit", {"hipErrorUnsupportedLimit", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 42
|
||||
// no analogue
|
||||
{"cudaErrorDuplicateVariableName", {"hipErrorDuplicateVariableName", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 43
|
||||
// no analogue
|
||||
@@ -693,41 +684,16 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
|
||||
{"cudaErrorDuplicateSurfaceName", {"hipErrorDuplicateSurfaceName", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 45
|
||||
// no analogue
|
||||
{"cudaErrorDevicesUnavailable", {"hipErrorDevicesUnavailable", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 46
|
||||
// CUDA_ERROR_INVALID_IMAGE = 200
|
||||
{"cudaErrorInvalidKernelImage", {"hipErrorInvalidImage", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 47
|
||||
// CUDA_ERROR_NO_BINARY_FOR_GPU = 209
|
||||
{"cudaErrorNoKernelImageForDevice", {"hipErrorNoBinaryForGpu", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 48
|
||||
// no analogue
|
||||
{"cudaErrorIncompatibleDriverContext", {"hipErrorIncompatibleDriverContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 49
|
||||
// CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED = 704
|
||||
{"cudaErrorPeerAccessAlreadyEnabled", {"hipErrorPeerAccessAlreadyEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 50
|
||||
// CUDA_ERROR_PEER_ACCESS_NOT_ENABLED = 705
|
||||
{"cudaErrorPeerAccessNotEnabled", {"hipErrorPeerAccessNotEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 51
|
||||
// cudaErrorMissingConfiguration = 52, hipErrorMissingConfiguration = 1001
|
||||
// TODO [HIP]: make hipErrorMissingConfiguration = 52
|
||||
// no analogue
|
||||
{"cudaErrorDeviceAlreadyInUse", {"hipErrorDeviceAlreadyInUse", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 54
|
||||
// CUDA_ERROR_PROFILER_DISABLED = 5
|
||||
{"cudaErrorProfilerDisabled", {"hipErrorProfilerDisabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 55
|
||||
// Deprecated since CUDA 5.0
|
||||
// CUDA_ERROR_PROFILER_NOT_INITIALIZED = 6
|
||||
{"cudaErrorProfilerNotInitialized", {"hipErrorProfilerNotInitialized", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 56
|
||||
// Deprecated since CUDA 5.0
|
||||
// CUDA_ERROR_PROFILER_ALREADY_STARTED = 7
|
||||
{"cudaErrorProfilerAlreadyStarted", {"hipErrorProfilerAlreadyStarted", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 57
|
||||
// Deprecated since CUDA 5.0
|
||||
// CUDA_ERROR_PROFILER_ALREADY_STOPPED = 8
|
||||
{"cudaErrorProfilerAlreadyStopped", {"hipErrorProfilerAlreadyStopped", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 58
|
||||
// CUDA_ERROR_ASSERT = 710
|
||||
{"cudaErrorAssert", {"hipErrorAssert", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 59
|
||||
// CUDA_ERROR_TOO_MANY_PEERS = 711
|
||||
{"cudaErrorTooManyPeers", {"hipErrorTooManyPeers", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 60
|
||||
// CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED = 712
|
||||
{"cudaErrorHostMemoryAlreadyRegistered", {"hipErrorHostMemoryAlreadyRegistered", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 61
|
||||
// CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED = 713
|
||||
{"cudaErrorHostMemoryNotRegistered", {"hipErrorHostMemoryNotRegistered", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 62
|
||||
// CUDA_ERROR_OPERATING_SYSTEM = 304
|
||||
{"cudaErrorOperatingSystem", {"hipErrorOperatingSystem", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 63
|
||||
// CUDA_ERROR_PEER_ACCESS_UNSUPPORTED = 217
|
||||
{"cudaErrorPeerAccessUnsupported", {"hipErrorPeerAccessUnsupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 64
|
||||
{"cudaErrorMissingConfiguration", {"hipErrorMissingConfiguration", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 52
|
||||
// cudaErrorPriorLaunchFailure = 53, hipErrorPriorLaunchFailure = 1005
|
||||
// TODO [HIP]: make hipErrorPriorLaunchFailure = 53
|
||||
// no analogue
|
||||
{"cudaErrorPriorLaunchFailure", {"hipErrorPriorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 53
|
||||
// no analogue
|
||||
{"cudaErrorLaunchMaxDepthExceeded", {"hipErrorLaunchMaxDepthExceeded", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 65
|
||||
// no analogue
|
||||
@@ -738,60 +704,174 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
|
||||
{"cudaErrorSyncDepthExceeded", {"hipErrorSyncDepthExceeded", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 68
|
||||
// no analogue
|
||||
{"cudaErrorLaunchPendingCountExceeded", {"hipErrorLaunchPendingCountExceeded", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 69
|
||||
// CUDA_ERROR_NOT_PERMITTED = 800
|
||||
{"cudaErrorNotPermitted", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 70
|
||||
// CUDA_ERROR_NOT_SUPPORTED = 801
|
||||
{"cudaErrorNotSupported", {"hipErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 71
|
||||
// CUDA_ERROR_HARDWARE_STACK_ERROR = 714
|
||||
{"cudaErrorHardwareStackError", {"hipErrorHardwareStackError", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 72
|
||||
// CUDA_ERROR_ILLEGAL_INSTRUCTION = 715
|
||||
{"cudaErrorIllegalInstruction", {"hipErrorIllegalInstruction", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 73
|
||||
// CUDA_ERROR_MISALIGNED_ADDRESS = 716
|
||||
{"cudaErrorMisalignedAddress", {"hipErrorMisalignedAddress", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 74
|
||||
// CUDA_ERROR_INVALID_ADDRESS_SPACE = 717
|
||||
{"cudaErrorInvalidAddressSpace", {"hipErrorInvalidAddressSpace", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 75
|
||||
// CUDA_ERROR_INVALID_PC = 718
|
||||
{"cudaErrorInvalidPc", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 76
|
||||
// CUDA_ERROR_ILLEGAL_ADDRESS = 700
|
||||
{"cudaErrorIllegalAddress", {"hipErrorIllegalAddress", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 77
|
||||
// CUDA_ERROR_INVALID_PTX = 218
|
||||
{"cudaErrorInvalidPtx", {"hipErrorInvalidKernelFile", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 78
|
||||
// CUDA_ERROR_INVALID_GRAPHICS_CONTEXT = 219
|
||||
{"cudaErrorInvalidGraphicsContext", {"hipErrorInvalidGraphicsContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 79
|
||||
// CUDA_ERROR_NVLINK_UNCORRECTABLE = 220
|
||||
{"cudaErrorNvlinkUncorrectable", {"hipErrorNvlinkUncorrectable", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 80
|
||||
// cudaErrorInvalidDeviceFunction = 98, hipErrorInvalidDeviceFunction = 1008
|
||||
// TODO [HIP]: make hipErrorInvalidDeviceFunction = 98
|
||||
// no analogue
|
||||
{"cudaErrorJitCompilerNotFound", {"hipErrorJitCompilerNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 81
|
||||
{"cudaErrorInvalidDeviceFunction", {"hipErrorInvalidDeviceFunction", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 98
|
||||
// cudaErrorNoDevice = 100, CUDA_ERROR_NO_DEVICE = 100, hipErrorNoDevice = 1038
|
||||
// TODO [HIP]: make hipErrorNoDevice = 100
|
||||
// CUDA_ERROR_NO_DEVICE
|
||||
{"cudaErrorNoDevice", {"hipErrorNoDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 100
|
||||
// cudaErrorInvalidDevice = 101, CUDA_ERROR_INVALID_DEVICE = 101, hipErrorInvalidDevice = 1010
|
||||
// TODO [HIP]: make hipErrorInvalidDevice = 101
|
||||
// CUDA_ERROR_INVALID_DEVICE
|
||||
{"cudaErrorInvalidDevice", {"hipErrorInvalidDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 101
|
||||
// no analogue
|
||||
{"cudaErrorCooperativeLaunchTooLarge", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 82
|
||||
// CUDA_ERROR_SYSTEM_NOT_READY = 802
|
||||
{"cudaErrorSystemNotReady", {"hipErrorSystemNotReady", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 83
|
||||
// CUDA_ERROR_ILLEGAL_STATE = 401
|
||||
{"cudaErrorIllegalState", {"hipErrorIllegalState", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 84
|
||||
{"cudaErrorStartupFailure", {"hipErrorStartupFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 127
|
||||
// CUDA_ERROR_INVALID_IMAGE
|
||||
{"cudaErrorInvalidKernelImage", {"hipErrorInvalidImage", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 200
|
||||
// CUDA_ERROR_INVALID_CONTEXT
|
||||
{"cudaErrorDeviceUninitilialized", {"hipErrorInvalidContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 201
|
||||
// cudaErrorMapBufferObjectFailed = 205, CUDA_ERROR_MAP_FAILED = 205, hipErrorMapFailed = 205, hipErrorMapBufferObjectFailed = 1071
|
||||
// TODO [HIP]: remove hipErrorMapBufferObjectFailed
|
||||
// TODO [HIPIFY]: rename hipErrorMapBufferObjectFailed to hipErrorMapFailed
|
||||
// CUDA_ERROR_MAP_FAILED
|
||||
{"cudaErrorMapBufferObjectFailed", {"hipErrorMapBufferObjectFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 205
|
||||
// CUDA_ERROR_UNMAP_FAILED
|
||||
{"cudaErrorUnmapBufferObjectFailed", {"hipErrorUnmapFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 206
|
||||
// CUDA_ERROR_ARRAY_IS_MAPPED
|
||||
{"cudaErrorArrayIsMapped", {"hipErrorArrayIsMapped", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 207
|
||||
// CUDA_ERROR_ALREADY_MAPPED
|
||||
{"cudaErrorAlreadyMapped", {"hipErrorAlreadyMapped", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 208
|
||||
// CUDA_ERROR_NO_BINARY_FOR_GPU
|
||||
{"cudaErrorNoKernelImageForDevice", {"hipErrorNoBinaryForGpu", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 209
|
||||
// CUDA_ERROR_ALREADY_ACQUIRED
|
||||
{"cudaErrorAlreadyAcquired", {"hipErrorAlreadyAcquired", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 210
|
||||
// CUDA_ERROR_NOT_MAPPED
|
||||
{"cudaErrorNotMapped", {"hipErrorNotMapped", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 211
|
||||
// CUDA_ERROR_NOT_MAPPED_AS_ARRAY
|
||||
{"cudaErrorNotMappedAsArray", {"hipErrorNotMappedAsArray", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 212
|
||||
// CUDA_ERROR_NOT_MAPPED_AS_POINTER
|
||||
{"cudaErrorNotMappedAsPointer", {"hipErrorNotMappedAsPointer", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 213
|
||||
// CUDA_ERROR_ECC_UNCORRECTABLE
|
||||
{"cudaErrorECCUncorrectable", {"hipErrorECCNotCorrectable", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 214
|
||||
// CUDA_ERROR_UNSUPPORTED_LIMIT
|
||||
{"cudaErrorUnsupportedLimit", {"hipErrorUnsupportedLimit", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 215
|
||||
// CUDA_ERROR_CONTEXT_ALREADY_IN_USE
|
||||
{"cudaErrorDeviceAlreadyInUse", {"hipErrorContextAlreadyInUse", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 216
|
||||
// CUDA_ERROR_PEER_ACCESS_UNSUPPORTED
|
||||
{"cudaErrorPeerAccessUnsupported", {"hipErrorPeerAccessUnsupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 217
|
||||
// CUDA_ERROR_INVALID_PTX
|
||||
{"cudaErrorInvalidPtx", {"hipErrorInvalidKernelFile", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 218
|
||||
// CUDA_ERROR_INVALID_GRAPHICS_CONTEXT
|
||||
{"cudaErrorInvalidGraphicsContext", {"hipErrorInvalidGraphicsContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 219
|
||||
// CUDA_ERROR_NVLINK_UNCORRECTABLE
|
||||
{"cudaErrorNvlinkUncorrectable", {"hipErrorNvlinkUncorrectable", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 220
|
||||
// CUDA_ERROR_JIT_COMPILER_NOT_FOUND
|
||||
{"cudaErrorJitCompilerNotFound", {"hipErrorJitCompilerNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 221
|
||||
// CUDA_ERROR_INVALID_SOURCE
|
||||
{"cudaErrorInvalidSource", {"hipErrorInvalidSource", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 300
|
||||
// CUDA_ERROR_FILE_NOT_FOUND
|
||||
{"cudaErrorFileNotFound", {"hipErrorFileNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 301
|
||||
// CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND
|
||||
{"cudaErrorSharedObjectSymbolNotFound", {"hipErrorSharedObjectSymbolNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 302
|
||||
// CUDA_ERROR_SHARED_OBJECT_INIT_FAILED
|
||||
{"cudaErrorSharedObjectInitFailed", {"hipErrorSharedObjectInitFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 303
|
||||
// CUDA_ERROR_OPERATING_SYSTEM
|
||||
{"cudaErrorOperatingSystem", {"hipErrorOperatingSystem", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 304
|
||||
// cudaErrorInvalidResourceHandle = 400, CUDA_ERROR_INVALID_HANDLE = 400, hipErrorInvalidHandle = 400, hipErrorInvalidResourceHandle = 1033
|
||||
// TODO [HIP]: remove hipErrorInvalidResourceHandle
|
||||
// TODO [HIPIFY]: rename hipErrorInvalidResourceHandle to hipErrorInvalidHandle
|
||||
// CUDA_ERROR_INVALID_HANDLE
|
||||
{"cudaErrorInvalidResourceHandle", {"hipErrorInvalidResourceHandle", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 400
|
||||
// CUDA_ERROR_ILLEGAL_STATE
|
||||
{"cudaErrorIllegalState", {"hipErrorIllegalState", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 401
|
||||
// CUDA_ERROR_NOT_FOUND
|
||||
{"cudaErrorSymbolNotFound", {"hipErrorNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 500
|
||||
// cudaErrorNotReady = 600, CUDA_ERROR_NOT_READY = 600, hipErrorNotReady = 1034
|
||||
// TODO [HIP]: make hipErrorNotReady = 600
|
||||
// CUDA_ERROR_NOT_READY
|
||||
{"cudaErrorNotReady", {"hipErrorNotReady", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 600
|
||||
// CUDA_ERROR_ILLEGAL_ADDRESS
|
||||
{"cudaErrorIllegalAddress", {"hipErrorIllegalAddress", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 700
|
||||
// cudaErrorLaunchOutOfResources = 701, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701, hipErrorLaunchOutOfResources = 1007
|
||||
// TODO [HIP]: make hipErrorLaunchOutOfResources = 701
|
||||
// CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES
|
||||
{"cudaErrorLaunchOutOfResources", {"hipErrorLaunchOutOfResources", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 701
|
||||
// cudaErrorLaunchTimeout = 702, CUDA_ERROR_LAUNCH_TIMEOUT = 702, hipErrorLaunchTimeOut = 1006
|
||||
// TODO [HIP]: make hipErrorLaunchTimeOut = 702
|
||||
// CUDA_ERROR_LAUNCH_TIMEOUT
|
||||
{"cudaErrorLaunchTimeout", {"hipErrorLaunchTimeOut", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 702
|
||||
// CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING
|
||||
{"cudaErrorLaunchIncompatibleTexturing", {"hipErrorLaunchIncompatibleTexturing", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 703
|
||||
// cudaErrorPeerAccessAlreadyEnabled = 704, CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED = 704, hipErrorPeerAccessAlreadyEnabled = 1050
|
||||
// TODO [HIP]: make hipErrorPeerAccessAlreadyEnabled = 704
|
||||
// CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED
|
||||
{"cudaErrorPeerAccessAlreadyEnabled", {"hipErrorPeerAccessAlreadyEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 704
|
||||
// cudaErrorPeerAccessNotEnabled = 705, CUDA_ERROR_PEER_ACCESS_NOT_ENABLED = 705, hipErrorPeerAccessNotEnabled = 1051
|
||||
// TODO [HIP]: make hipErrorPeerAccessNotEnabled = 705
|
||||
// CUDA_ERROR_PEER_ACCESS_NOT_ENABLED
|
||||
{"cudaErrorPeerAccessNotEnabled", {"hipErrorPeerAccessNotEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 705
|
||||
// cudaErrorSetOnActiveProcess = 708, CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE = 708, hipErrorSetOnActiveProcess = 305
|
||||
// TODO [HIP]: make hipErrorSetOnActiveProcess = 708
|
||||
// CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE
|
||||
{"cudaErrorSetOnActiveProcess", {"hipErrorSetOnActiveProcess", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 708
|
||||
// CUDA_ERROR_CONTEXT_IS_DESTROYED
|
||||
{"cudaErrorContextIsDestroyed", {"hipErrorContextIsDestroyed", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 709
|
||||
// cudaErrorAssert = 710, CUDA_ERROR_ASSERT = 710, hipErrorAssert = 1081
|
||||
// TODO [HIP]: make hipErrorAssert = 710
|
||||
// CUDA_ERROR_ASSERT
|
||||
{"cudaErrorAssert", {"hipErrorAssert", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 710
|
||||
// CUDA_ERROR_TOO_MANY_PEERS
|
||||
{"cudaErrorTooManyPeers", {"hipErrorTooManyPeers", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 711
|
||||
// cudaErrorHostMemoryAlreadyRegistered = 712, CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED = 712, hipErrorHostMemoryAlreadyRegistered = 1061
|
||||
// TODO [HIP]: make hipErrorHostMemoryAlreadyRegistered = 712
|
||||
// CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED
|
||||
{"cudaErrorHostMemoryAlreadyRegistered", {"hipErrorHostMemoryAlreadyRegistered", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 712
|
||||
// cudaErrorHostMemoryNotRegistered = 713, CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED = 713, hipErrorHostMemoryNotRegistered = 1062
|
||||
// TODO [HIP]: make hipErrorHostMemoryNotRegistered = 713
|
||||
// CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED
|
||||
{"cudaErrorHostMemoryNotRegistered", {"hipErrorHostMemoryNotRegistered", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 713
|
||||
// CUDA_ERROR_HARDWARE_STACK_ERROR
|
||||
{"cudaErrorHardwareStackError", {"hipErrorHardwareStackError", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 714
|
||||
// CUDA_ERROR_ILLEGAL_INSTRUCTION
|
||||
{"cudaErrorIllegalInstruction", {"hipErrorIllegalInstruction", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 715
|
||||
// CUDA_ERROR_MISALIGNED_ADDRESS
|
||||
{"cudaErrorMisalignedAddress", {"hipErrorMisalignedAddress", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 716
|
||||
// CUDA_ERROR_INVALID_ADDRESS_SPACE
|
||||
{"cudaErrorInvalidAddressSpace", {"hipErrorInvalidAddressSpace", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 717
|
||||
// CUDA_ERROR_INVALID_PC
|
||||
{"cudaErrorInvalidPc", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 718
|
||||
// cudaErrorLaunchFailure = 719, CUDA_ERROR_LAUNCH_FAILED = 719, hipErrorLaunchFailure = 1004
|
||||
// TODO [HIP]: make hipErrorSetOnActiveProcess = 719
|
||||
// CUDA_ERROR_LAUNCH_FAILED
|
||||
{"cudaErrorLaunchFailure", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 719
|
||||
// no analogue
|
||||
{"cudaErrorCooperativeLaunchTooLarge", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 720
|
||||
// CUDA_ERROR_NOT_PERMITTED
|
||||
{"cudaErrorNotPermitted", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 800
|
||||
// cudaErrorNotSupported = 801, CUDA_ERROR_NOT_SUPPORTED = 801, hipErrorNotSupported = 1082
|
||||
// TODO [HIP]: make hipErrorNotSupported = 801
|
||||
// CUDA_ERROR_NOT_SUPPORTED
|
||||
{"cudaErrorNotSupported", {"hipErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 801
|
||||
// CUDA_ERROR_SYSTEM_NOT_READY
|
||||
{"cudaErrorSystemNotReady", {"hipErrorSystemNotReady", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 802
|
||||
// CUDA_ERROR_SYSTEM_DRIVER_MISMATCH
|
||||
{"cudaErrorSystemDriverMismatch", {"hipErrorSystemDriverMismatch", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 803
|
||||
// CUDA_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE
|
||||
{"cudaErrorCompatNotSupportedOnDevice", {"hipErrorCompatNotSupportedOnDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 804
|
||||
// no analogue
|
||||
{"cudaErrorStartupFailure", {"hipErrorStartupFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 127
|
||||
// CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED = 900
|
||||
// CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED
|
||||
{"cudaErrorStreamCaptureUnsupported", {"hipErrorStreamCaptureUnsupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 900
|
||||
// CUDA_ERROR_STREAM_CAPTURE_INVALIDATED = 901
|
||||
// CUDA_ERROR_STREAM_CAPTURE_INVALIDATED
|
||||
{"cudaErrorStreamCaptureInvalidated", {"hipErrorStreamCaptureInvalidated", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 901
|
||||
// CUDA_ERROR_STREAM_CAPTURE_MERGE = 902
|
||||
// CUDA_ERROR_STREAM_CAPTURE_MERGE
|
||||
{"cudaErrorStreamCaptureMerge", {"hipErrorStreamCaptureMerge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 902
|
||||
// CUDA_ERROR_STREAM_CAPTURE_UNMATCHED = 903
|
||||
// CUDA_ERROR_STREAM_CAPTURE_UNMATCHED
|
||||
{"cudaErrorStreamCaptureUnmatched", {"hipErrorStreamCaptureUnmatched", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 903
|
||||
// CUDA_ERROR_STREAM_CAPTURE_UNJOINED = 904
|
||||
// CUDA_ERROR_STREAM_CAPTURE_UNJOINED
|
||||
{"cudaErrorStreamCaptureUnjoined", {"hipErrorStreamCaptureUnjoined", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 904
|
||||
// CUDA_ERROR_STREAM_CAPTURE_ISOLATION = 905
|
||||
// CUDA_ERROR_STREAM_CAPTURE_ISOLATION
|
||||
{"cudaErrorStreamCaptureIsolation", {"hipErrorStreamCaptureIsolation", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 905
|
||||
// CUDA_ERROR_STREAM_CAPTURE_IMPLICIT = 906
|
||||
// CUDA_ERROR_STREAM_CAPTURE_IMPLICIT
|
||||
{"cudaErrorStreamCaptureImplicit", {"hipErrorStreamCaptureImplicit", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 906
|
||||
// CUDA_ERROR_CAPTURED_EVENT = 907
|
||||
// CUDA_ERROR_CAPTURED_EVENT
|
||||
{"cudaErrorCapturedEvent", {"hipErrorCapturedEvent", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 907
|
||||
// CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD = 908
|
||||
// CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD
|
||||
{"cudaErrorStreamCaptureWrongThread", {"hipErrorStreamCaptureWrongThread", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 908
|
||||
// cudaErrorUnknown = 999, CUDA_ERROR_UNKNOWN = 999, hipErrorUnknown = 1030
|
||||
// TODO [HIP]: make hipErrorUnknown = 999
|
||||
// CUDA_ERROR_UNKNOWN
|
||||
{"cudaErrorUnknown", {"hipErrorUnknown", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 999
|
||||
// Deprecated since CUDA 4.1
|
||||
{"cudaErrorApiFailureBase", {"hipErrorApiFailureBase", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 10000
|
||||
|
||||
|
||||
@@ -794,7 +794,7 @@ void *__amdgcn_get_dynamicgroupbaseptr() {
|
||||
return __get_dynamicgroupbaseptr();
|
||||
}
|
||||
|
||||
#if defined(__HCC__) && (__hcc_minor__ < 3)
|
||||
#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
|
||||
// hip.amdgcn.bc - sync threads
|
||||
#define __CLK_LOCAL_MEM_FENCE 0x01
|
||||
typedef unsigned __cl_mem_fence_flags;
|
||||
|
||||
@@ -80,7 +80,7 @@ extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(vo
|
||||
__device__ inline static __local void* __to_local(unsigned x) { return (__local void*)x; }
|
||||
#endif //__HIP_DEVICE_COMPILE__
|
||||
|
||||
#if defined(__HCC__) && (__hcc_minor__ < 3)
|
||||
#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
|
||||
// __llvm_fence* functions from device-libs/irif/src/fence.ll
|
||||
extern "C" __device__ void __llvm_fence_acq_sg(void);
|
||||
extern "C" __device__ void __llvm_fence_acq_wg(void);
|
||||
|
||||
@@ -176,5 +176,4 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
|
||||
numBlocks, dimBlocks, sharedMemBytes,
|
||||
stream, &config[0]);
|
||||
}
|
||||
|
||||
#pragma GCC visibility pop
|
||||
|
||||
@@ -27,6 +27,26 @@ THE SOFTWARE.
|
||||
// std::false_type, std result_of and std::true_type.
|
||||
#include <utility> // For std::declval.
|
||||
|
||||
#ifdef __has_include // Check if __has_include is present
|
||||
# if __has_include(<version>) // Check for version header
|
||||
# include <version>
|
||||
# if defined(__cpp_lib_is_invocable) && !defined(HIP_HAS_INVOCABLE)
|
||||
# define HIP_HAS_INVOCABLE __cpp_lib_is_invocable
|
||||
# endif
|
||||
# if defined(__cpp_lib_result_of_sfinae) && !defined(HIP_HAS_RESULT_OF_SFINAE)
|
||||
# define HIP_HAS_RESULT_OF_SFINAE __cpp_lib_result_of_sfinae
|
||||
# endif
|
||||
# endif
|
||||
#endif
|
||||
|
||||
#ifndef HIP_HAS_INVOCABLE
|
||||
#define HIP_HAS_INVOCABLE 0
|
||||
#endif
|
||||
|
||||
#ifndef HIP_HAS_RESULT_OF_SFINAE
|
||||
#define HIP_HAS_RESULT_OF_SFINAE 0
|
||||
#endif
|
||||
|
||||
namespace std { // TODO: these should be removed as soon as possible.
|
||||
#if (__cplusplus < 201406L)
|
||||
#if (__cplusplus < 201402L)
|
||||
@@ -48,58 +68,56 @@ namespace hip_impl {
|
||||
template <typename...>
|
||||
using void_t_ = void;
|
||||
|
||||
#if (__cplusplus < 201402L)
|
||||
template <FunctionalProcedure F, unsigned int n = 0u, typename = void>
|
||||
struct is_callable_impl : is_callable_impl<F, n + 1u> {};
|
||||
#if HIP_HAS_INVOCABLE
|
||||
template <typename, typename = void>
|
||||
struct is_callable_impl;
|
||||
|
||||
// Pointer to member function, call through non-pointer.
|
||||
template <FunctionalProcedure F, typename C, typename... Ts>
|
||||
struct is_callable_impl<
|
||||
F(C, Ts...), 0u,
|
||||
void_t_<decltype((std::declval<C>().*std::declval<F>())(std::declval<Ts>()...))> >
|
||||
: std::true_type {};
|
||||
|
||||
// Pointer to member function, call through pointer.
|
||||
template <FunctionalProcedure F, typename C, typename... Ts>
|
||||
struct is_callable_impl<
|
||||
F(C, Ts...), 1u,
|
||||
void_t_<decltype(((*std::declval<C>()).*std::declval<F>())(std::declval<Ts>()...))> >
|
||||
: std::true_type {};
|
||||
|
||||
// Pointer to member data, call through non-pointer, no args.
|
||||
template <FunctionalProcedure F, typename C>
|
||||
struct is_callable_impl<F(C), 2u, void_t_<decltype(std::declval<C>().*std::declval<F>())> >
|
||||
: std::true_type {};
|
||||
|
||||
// Pointer to member data, call through pointer, no args.
|
||||
template <FunctionalProcedure F, typename C>
|
||||
struct is_callable_impl<F(C), 3u, void_t_<decltype(*std::declval<C>().*std::declval<F>())> >
|
||||
: std::true_type {};
|
||||
|
||||
// General call, n args.
|
||||
template <FunctionalProcedure F, typename... Ts>
|
||||
struct is_callable_impl<F(Ts...), 4u, void_t_<decltype(std::declval<F>()(std::declval<Ts>()...))> >
|
||||
: std::true_type {};
|
||||
|
||||
// Not callable.
|
||||
template <FunctionalProcedure F>
|
||||
struct is_callable_impl<F, 5u> : std::false_type {};
|
||||
#elif (__cplusplus < 201703L)
|
||||
struct is_callable_impl<F(Ts...)> : std::is_invocable<F, Ts...> {};
|
||||
#elif HIP_HAS_RESULT_OF_SFINAE
|
||||
template <typename, typename = void>
|
||||
struct is_callable_impl : std::false_type {};
|
||||
|
||||
template <FunctionalProcedure F, typename... Ts>
|
||||
struct is_callable_impl<F(Ts...), void_t_<typename std::result_of<F(Ts...)>::type > > : std::true_type {};
|
||||
#else
|
||||
template <class Base, class T, class Derived>
|
||||
auto simple_invoke(T Base::*pmd, Derived&& ref)
|
||||
-> decltype(static_cast<Derived&&>(ref).*pmd);
|
||||
|
||||
template <class PMD, class Pointer>
|
||||
auto simple_invoke(PMD&& pmd, Pointer&& ptr)
|
||||
-> decltype((*static_cast<Pointer&&>(ptr)).*static_cast<PMD&&>(pmd));
|
||||
|
||||
// C++17
|
||||
template <class Base, class T, class Derived>
|
||||
auto simple_invoke(T Base::*pmd, const std::reference_wrapper<Derived>& ref)
|
||||
-> decltype(ref.get().*pmd);
|
||||
|
||||
template <class Base, class T, class Derived, class... Args>
|
||||
auto simple_invoke(T Base::*pmf, Derived&& ref, Args&&... args)
|
||||
-> decltype((static_cast<Derived&&>(ref).*pmf)(static_cast<Args&&>(args)...));
|
||||
|
||||
template <class PMF, class Pointer, class... Args>
|
||||
auto simple_invoke(PMF&& pmf, Pointer&& ptr, Args&&... args)
|
||||
-> decltype(((*static_cast<Pointer&&>(ptr)).*static_cast<PMF&&>(pmf))(static_cast<Args&&>(args)...));
|
||||
|
||||
template <class Base, class T, class Derived, class... Args>
|
||||
auto simple_invoke(T Base::*pmf, const std::reference_wrapper<Derived>& ref, Args&&... args)
|
||||
-> decltype((ref.get().*pmf)(static_cast<Args&&>(args)...));
|
||||
|
||||
template<class F, class... Ts>
|
||||
auto simple_invoke(F&& f, Ts&&... xs)
|
||||
-> decltype(f(static_cast<Ts&&>(xs)...));
|
||||
|
||||
template <typename, typename = void>
|
||||
struct is_callable_impl : std::false_type {};
|
||||
|
||||
template <FunctionalProcedure F, typename... Ts>
|
||||
struct is_callable_impl<F(Ts...), void_t_<std::invoke_result<F(Ts...)> > > : std::true_type {};
|
||||
struct is_callable_impl<F(Ts...), void_t_<decltype(simple_invoke(std::declval<F>(), std::declval<Ts>()...))> >
|
||||
: std::true_type {};
|
||||
|
||||
#endif
|
||||
|
||||
template <typename Call>
|
||||
struct is_callable : is_callable_impl<Call> {};
|
||||
|
||||
|
||||
@@ -186,7 +186,7 @@ __device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double
|
||||
__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
|
||||
hipDoubleComplex ret;
|
||||
ret.x = z.x;
|
||||
ret.y = z.y;
|
||||
ret.y = -z.y;
|
||||
return ret;
|
||||
}
|
||||
|
||||
|
||||
@@ -391,7 +391,7 @@ hipError_t hipSetDevice(int deviceId);
|
||||
* This device is used implicitly for HIP runtime APIs called by this thread.
|
||||
* hipGetDevice returns in * @p device the default device for the calling host thread.
|
||||
*
|
||||
* @returns #hipSuccess
|
||||
* @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue
|
||||
*
|
||||
* @see hipSetDevice, hipGetDevicesizeBytes
|
||||
*/
|
||||
@@ -883,7 +883,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
|
||||
|
||||
* #hipEventDefault : Default flag. The event will use active synchronization and will support
|
||||
timing. Blocking synchronization provides lowest possible latency at the expense of dedicating a
|
||||
CPU to poll on the eevent.
|
||||
CPU to poll on the event.
|
||||
* #hipEventBlockingSync : The event will use blocking synchronization : if hipEventSynchronize is
|
||||
called on this event, the thread will block until the event completes. This can increase latency
|
||||
for the synchroniation but can result in lower power and more resources for other CPU threads.
|
||||
@@ -935,7 +935,7 @@ hipError_t hipEventCreate(hipEvent_t* event);
|
||||
* If hipEventRecord() has been previously called on this event, then this call will overwrite any
|
||||
* existing state in event.
|
||||
*
|
||||
* If this function is called on a an event that is currently being recorded, results are undefined
|
||||
* If this function is called on an event that is currently being recorded, results are undefined
|
||||
* - either outstanding recording may save state into the event, and the order is not guaranteed.
|
||||
*
|
||||
* @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventSynchronize,
|
||||
@@ -1277,7 +1277,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height
|
||||
* @param[in] height Requested pitched allocation height
|
||||
*
|
||||
* If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
|
||||
* The intended usage of pitch is as a separate parameter of the allocation, used to compute addresses within the 2D array.
|
||||
* The intended usage of pitch is as a separate parameter of the allocation, used to compute addresses within the 2D array.
|
||||
* Given the row and column of an array element of type T, the address is computed as:
|
||||
* T* pElement = (T*)((char*)BaseAddress + Row * Pitch) + Column;
|
||||
*
|
||||
@@ -1361,6 +1361,10 @@ hipError_t hipHostFree(void* ptr);
|
||||
*/
|
||||
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind);
|
||||
|
||||
// TODO: Add description
|
||||
hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
|
||||
hipMemcpyKind kind, hipStream_t stream);
|
||||
|
||||
/**
|
||||
* @brief Copy data from Host to Device
|
||||
*
|
||||
@@ -1498,7 +1502,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName,
|
||||
#else
|
||||
hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*);
|
||||
|
||||
#ifdef __cplusplus //Start : Not supported in gcc
|
||||
#ifdef __cplusplus //Start : Not supported in gcc
|
||||
namespace hip_impl {
|
||||
inline
|
||||
__attribute__((visibility("hidden")))
|
||||
@@ -2885,6 +2889,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigne
|
||||
unsigned int sharedMemBytes, hipStream_t stream,
|
||||
void** kernelParams, void** extra);
|
||||
|
||||
|
||||
/**
|
||||
* @brief launches kernel f with launch parameters and shared memory on stream with arguments passed
|
||||
* to kernelparams or extra, where thread blocks can cooperate and synchronize as they execute
|
||||
|
||||
@@ -40,6 +40,7 @@ THE SOFTWARE.
|
||||
#endif
|
||||
|
||||
#if defined(__cplusplus) && defined(__clang__)
|
||||
#include <iosfwd>
|
||||
#include <type_traits>
|
||||
|
||||
namespace hip_impl {
|
||||
@@ -68,6 +69,23 @@ THE SOFTWARE.
|
||||
}
|
||||
};
|
||||
|
||||
friend
|
||||
inline
|
||||
std::ostream& operator<<(std::ostream& os,
|
||||
const Scalar_accessor& x) noexcept {
|
||||
return os << x.data[idx];
|
||||
}
|
||||
friend
|
||||
inline
|
||||
std::istream& operator>>(std::istream& is,
|
||||
Scalar_accessor& x) noexcept {
|
||||
T tmp;
|
||||
is >> tmp;
|
||||
x.data[idx] = tmp;
|
||||
|
||||
return is;
|
||||
}
|
||||
|
||||
// Idea from https://t0rakka.silvrback.com/simd-scalar-accessor
|
||||
Vector data;
|
||||
|
||||
@@ -76,6 +94,17 @@ THE SOFTWARE.
|
||||
__host__ __device__
|
||||
operator T() const volatile noexcept { return data[idx]; }
|
||||
|
||||
__host__ __device__
|
||||
operator T&() noexcept {
|
||||
return reinterpret_cast<
|
||||
T (&)[sizeof(Vector) / sizeof(T)]>(data)[idx];
|
||||
}
|
||||
__host__ __device__
|
||||
operator volatile T&() volatile noexcept {
|
||||
return reinterpret_cast<
|
||||
volatile T (&)[sizeof(Vector) / sizeof(T)]>(data)[idx];
|
||||
}
|
||||
|
||||
__host__ __device__
|
||||
Address operator&() const noexcept { return Address{this}; }
|
||||
|
||||
@@ -115,22 +144,40 @@ THE SOFTWARE.
|
||||
return *this;
|
||||
}
|
||||
|
||||
// TODO: convertibility is too restrictive, constraint should be on
|
||||
// the operator being invocable with a value of type U.
|
||||
template<
|
||||
typename U,
|
||||
typename std::enable_if<
|
||||
std::is_convertible<U, T>{}>::type* = nullptr>
|
||||
__host__ __device__
|
||||
Scalar_accessor& operator+=(T x) noexcept {
|
||||
Scalar_accessor& operator+=(U x) noexcept {
|
||||
data[idx] += x;
|
||||
return *this;
|
||||
}
|
||||
template<
|
||||
typename U,
|
||||
typename std::enable_if<
|
||||
std::is_convertible<U, T>{}>::type* = nullptr>
|
||||
__host__ __device__
|
||||
Scalar_accessor& operator-=(T x) noexcept {
|
||||
Scalar_accessor& operator-=(U x) noexcept {
|
||||
data[idx] -= x;
|
||||
return *this;
|
||||
}
|
||||
|
||||
template<
|
||||
typename U,
|
||||
typename std::enable_if<
|
||||
std::is_convertible<U, T>{}>::type* = nullptr>
|
||||
__host__ __device__
|
||||
Scalar_accessor& operator*=(T x) noexcept {
|
||||
Scalar_accessor& operator*=(U x) noexcept {
|
||||
data[idx] *= x;
|
||||
return *this;
|
||||
}
|
||||
template<
|
||||
typename U,
|
||||
typename std::enable_if<
|
||||
std::is_convertible<U, T>{}>::type* = nullptr>
|
||||
__host__ __device__
|
||||
Scalar_accessor& operator/=(T x) noexcept {
|
||||
data[idx] /= x;
|
||||
@@ -138,50 +185,56 @@ THE SOFTWARE.
|
||||
}
|
||||
template<
|
||||
typename U = T,
|
||||
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
|
||||
typename std::enable_if<std::is_convertible<U, T>{} &&
|
||||
std::is_integral<U>{}>::type* = nullptr>
|
||||
__host__ __device__
|
||||
Scalar_accessor& operator%=(T x) noexcept {
|
||||
Scalar_accessor& operator%=(U x) noexcept {
|
||||
data[idx] %= x;
|
||||
return *this;
|
||||
}
|
||||
|
||||
template<
|
||||
typename U = T,
|
||||
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
|
||||
typename std::enable_if<std::is_convertible<U, T>{} &&
|
||||
std::is_integral<U>{}>::type* = nullptr>
|
||||
__host__ __device__
|
||||
Scalar_accessor& operator>>=(T x) noexcept {
|
||||
Scalar_accessor& operator>>=(U x) noexcept {
|
||||
data[idx] >>= x;
|
||||
return *this;
|
||||
}
|
||||
template<
|
||||
typename U = T,
|
||||
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
|
||||
typename std::enable_if<std::is_convertible<U, T>{} &&
|
||||
std::is_integral<U>{}>::type* = nullptr>
|
||||
__host__ __device__
|
||||
Scalar_accessor& operator<<=(T x) noexcept {
|
||||
Scalar_accessor& operator<<=(U x) noexcept {
|
||||
data[idx] <<= x;
|
||||
return *this;
|
||||
}
|
||||
template<
|
||||
typename U = T,
|
||||
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
|
||||
typename std::enable_if<std::is_convertible<U, T>{} &&
|
||||
std::is_integral<U>{}>::type* = nullptr>
|
||||
__host__ __device__
|
||||
Scalar_accessor& operator&=(T x) noexcept {
|
||||
Scalar_accessor& operator&=(U x) noexcept {
|
||||
data[idx] &= x;
|
||||
return *this;
|
||||
}
|
||||
template<
|
||||
typename U = T,
|
||||
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
|
||||
typename std::enable_if<std::is_convertible<U, T>{} &&
|
||||
std::is_integral<U>{}>::type* = nullptr>
|
||||
__host__ __device__
|
||||
Scalar_accessor& operator|=(T x) noexcept {
|
||||
Scalar_accessor& operator|=(U x) noexcept {
|
||||
data[idx] |= x;
|
||||
return *this;
|
||||
}
|
||||
template<
|
||||
typename U = T,
|
||||
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
|
||||
typename std::enable_if<std::is_convertible<U, T>{} &&
|
||||
std::is_integral<U>{}>::type* = nullptr>
|
||||
__host__ __device__
|
||||
Scalar_accessor& operator^=(T x) noexcept {
|
||||
Scalar_accessor& operator^=(U x) noexcept {
|
||||
data[idx] ^= x;
|
||||
return *this;
|
||||
}
|
||||
@@ -198,6 +251,8 @@ THE SOFTWARE.
|
||||
Native_vec_ data;
|
||||
hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
|
||||
};
|
||||
|
||||
using value_type = T;
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
@@ -209,6 +264,8 @@ THE SOFTWARE.
|
||||
hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
|
||||
hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
|
||||
};
|
||||
|
||||
using value_type = T;
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
@@ -367,6 +424,8 @@ THE SOFTWARE.
|
||||
T z;
|
||||
};
|
||||
};
|
||||
|
||||
using value_type = T;
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
@@ -380,6 +439,8 @@ THE SOFTWARE.
|
||||
hip_impl::Scalar_accessor<T, Native_vec_, 2> z;
|
||||
hip_impl::Scalar_accessor<T, Native_vec_, 3> w;
|
||||
};
|
||||
|
||||
using value_type = T;
|
||||
};
|
||||
|
||||
template<typename T, unsigned int rank>
|
||||
|
||||
@@ -22,7 +22,7 @@ THE SOFTWARE.
|
||||
|
||||
#ifndef HIP_INCLUDE_HIP_HIP_EXT_H
|
||||
#define HIP_INCLUDE_HIP_HIP_EXT_H
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#include "hip/hip_runtime.h"
|
||||
#ifdef __HCC__
|
||||
|
||||
// Forward declarations:
|
||||
@@ -60,12 +60,11 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a
|
||||
|
||||
#endif // #ifdef __HCC__
|
||||
|
||||
|
||||
/**
|
||||
* @brief launches kernel f with launch parameters and shared memory on stream with arguments passed
|
||||
to kernelparams or extra
|
||||
*
|
||||
* @param [in[ f Kernel to launch.
|
||||
* @param [in[ f Kernel to launch.
|
||||
* @param [in] gridDimX X grid dimension specified in work-items
|
||||
* @param [in] gridDimY Y grid dimension specified in work-items
|
||||
* @param [in] gridDimZ Z grid dimension specified in work-items
|
||||
@@ -88,7 +87,6 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a
|
||||
*
|
||||
* @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please
|
||||
refer to hip_porting_driver_api.md for sample usage.
|
||||
|
||||
* HIP/ROCm actually updates the start event when the associated kernel completes.
|
||||
*/
|
||||
HIP_PUBLIC_API
|
||||
@@ -111,8 +109,62 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
|
||||
hipEvent_t stopEvent = nullptr)
|
||||
__attribute__((deprecated("use hipExtModuleLaunchKernel instead")));
|
||||
|
||||
#if !__HIP_VDI__ && defined(__cplusplus)
|
||||
|
||||
namespace hip_impl {
|
||||
inline
|
||||
__attribute__((visibility("hidden")))
|
||||
void hipExtLaunchKernelGGLImpl(
|
||||
std::uintptr_t function_address,
|
||||
const dim3& numBlocks,
|
||||
const dim3& dimBlocks,
|
||||
std::uint32_t sharedMemBytes,
|
||||
hipStream_t stream,
|
||||
hipEvent_t startEvent,
|
||||
hipEvent_t stopEvent,
|
||||
std::uint32_t flags,
|
||||
void** kernarg) {
|
||||
|
||||
const auto& kd = hip_impl::get_program_state()
|
||||
.kernel_descriptor(function_address, target_agent(stream));
|
||||
|
||||
hipExtModuleLaunchKernel(kd, numBlocks.x * dimBlocks.x,
|
||||
numBlocks.y * dimBlocks.y,
|
||||
numBlocks.z * dimBlocks.z,
|
||||
dimBlocks.x, dimBlocks.y, dimBlocks.z,
|
||||
sharedMemBytes, stream, nullptr, kernarg,
|
||||
startEvent, stopEvent, flags);
|
||||
}
|
||||
} // namespace hip_impl
|
||||
|
||||
template <typename... Args, typename F = void (*)(Args...)>
|
||||
inline
|
||||
void hipExtLaunchKernelGGL(F kernel, const dim3& numBlocks,
|
||||
const dim3& dimBlocks, std::uint32_t sharedMemBytes,
|
||||
hipStream_t stream, hipEvent_t startEvent,
|
||||
hipEvent_t stopEvent, std::uint32_t flags,
|
||||
Args... args) {
|
||||
hip_impl::hip_init();
|
||||
auto kernarg =
|
||||
hip_impl::make_kernarg(kernel, std::tuple<Args...>{std::move(args)...});
|
||||
std::size_t kernarg_size = kernarg.size();
|
||||
|
||||
void* config[]{
|
||||
HIP_LAUNCH_PARAM_BUFFER_POINTER,
|
||||
kernarg.data(),
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE,
|
||||
&kernarg_size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
|
||||
hip_impl::hipExtLaunchKernelGGLImpl(reinterpret_cast<std::uintptr_t>(kernel),
|
||||
numBlocks, dimBlocks, sharedMemBytes,
|
||||
stream, startEvent, stopEvent, flags,
|
||||
&config[0]);
|
||||
}
|
||||
#endif // !__HIP_VDI__ && defined(__cplusplus)
|
||||
|
||||
// doxygen end AMD-specific features
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
#endif // #ifdef HIP_INCLUDE_HIP_HIP_EXT_H
|
||||
#endif // #iidef HIP_INCLUDE_HIP_HIP_EXT_H
|
||||
|
||||
@@ -66,24 +66,24 @@ typedef enum hipMemcpyKind {
|
||||
#define HIP_LIBRARY_PATCH_LEVEL PATCH_LEVEL
|
||||
|
||||
// hipTextureAddressMode
|
||||
#define hipTextureAddressMode cudaTextureAddressMode
|
||||
typedef enum cudaTextureAddressMode hipTextureAddressMode;
|
||||
#define hipAddressModeWrap cudaAddressModeWrap
|
||||
#define hipAddressModeClamp cudaAddressModeClamp
|
||||
#define hipAddressModeMirror cudaAddressModeMirror
|
||||
#define hipAddressModeBorder cudaAddressModeBorder
|
||||
|
||||
// hipTextureFilterMode
|
||||
#define hipTextureFilterMode cudaTextureFilterMode
|
||||
typedef enum cudaTextureFilterMode hipTextureFilterMode;
|
||||
#define hipFilterModePoint cudaFilterModePoint
|
||||
#define hipFilterModeLinear cudaFilterModeLinear
|
||||
|
||||
// hipTextureReadMode
|
||||
#define hipTextureReadMode cudaTextureReadMode
|
||||
typedef enum cudaTextureReadMode hipTextureReadMode;
|
||||
#define hipReadModeElementType cudaReadModeElementType
|
||||
#define hipReadModeNormalizedFloat cudaReadModeNormalizedFloat
|
||||
|
||||
// hipChannelFormatKind
|
||||
#define hipChannelFormatKind cudaChannelFormatKind
|
||||
typedef enum cudaChannelFormatKind hipChannelFormatKind;
|
||||
#define hipChannelFormatKindSigned cudaChannelFormatKindSigned
|
||||
#define hipChannelFormatKindUnsigned cudaChannelFormatKindUnsigned
|
||||
#define hipChannelFormatKindFloat cudaChannelFormatKindFloat
|
||||
@@ -167,7 +167,7 @@ typedef cudaIpcMemHandle_t hipIpcMemHandle_t;
|
||||
typedef enum cudaLimit hipLimit_t;
|
||||
typedef enum cudaFuncCache hipFuncCache_t;
|
||||
typedef CUcontext hipCtx_t;
|
||||
typedef cudaSharedMemConfig hipSharedMemConfig;
|
||||
typedef enum cudaSharedMemConfig hipSharedMemConfig;
|
||||
typedef CUfunc_cache hipFuncCache;
|
||||
typedef CUjit_option hipJitOption;
|
||||
typedef CUdevice hipDevice_t;
|
||||
@@ -177,8 +177,8 @@ typedef CUdeviceptr hipDeviceptr_t;
|
||||
typedef struct cudaArray hipArray;
|
||||
typedef struct cudaArray* hipArray_t;
|
||||
typedef struct cudaArray* hipArray_const_t;
|
||||
typedef cudaFuncAttributes hipFuncAttributes;
|
||||
typedef CUfunction_attribute hipFunction_attribute;
|
||||
typedef struct cudaFuncAttributes hipFuncAttributes;
|
||||
#define hipFunction_attribute CUfunction_attribute
|
||||
#define hip_Memcpy2D CUDA_MEMCPY2D
|
||||
#define hipMemcpy3DParms cudaMemcpy3DParms
|
||||
#define hipArrayDefault cudaArrayDefault
|
||||
@@ -196,8 +196,8 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t;
|
||||
#define hipTextureType3D cudaTextureType3D
|
||||
#define hipDeviceMapHost cudaDeviceMapHost
|
||||
|
||||
#define hipExtent cudaExtent
|
||||
#define hipPitchedPtr cudaPitchedPtr
|
||||
typedef struct cudaExtent hipExtent;
|
||||
typedef struct cudaPitchedPtr hipPitchedPtr;
|
||||
#define make_hipExtent make_cudaExtent
|
||||
#define make_hipPos make_cudaPos
|
||||
#define make_hipPitchedPtr make_cudaPitchedPtr
|
||||
@@ -205,10 +205,10 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t;
|
||||
#define hipStreamDefault cudaStreamDefault
|
||||
#define hipStreamNonBlocking cudaStreamNonBlocking
|
||||
|
||||
#define hipChannelFormatDesc cudaChannelFormatDesc
|
||||
#define hipResourceDesc cudaResourceDesc
|
||||
#define hipTextureDesc cudaTextureDesc
|
||||
#define hipResourceViewDesc cudaResourceViewDesc
|
||||
typedef struct cudaChannelFormatDesc hipChannelFormatDesc;
|
||||
typedef struct cudaResourceDesc hipResourceDesc;
|
||||
typedef struct cudaTextureDesc hipTextureDesc;
|
||||
typedef struct cudaResourceViewDesc hipResourceViewDesc;
|
||||
// adding code for hipmemSharedConfig
|
||||
#define hipSharedMemBankSizeDefault cudaSharedMemBankSizeDefault
|
||||
#define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte
|
||||
@@ -394,7 +394,7 @@ inline static enum cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind ki
|
||||
}
|
||||
}
|
||||
|
||||
inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode(
|
||||
inline static enum cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode(
|
||||
hipTextureAddressMode kind) {
|
||||
switch (kind) {
|
||||
case hipAddressModeWrap:
|
||||
@@ -410,7 +410,7 @@ inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMo
|
||||
}
|
||||
}
|
||||
|
||||
inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode(
|
||||
inline static enum cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode(
|
||||
hipTextureFilterMode kind) {
|
||||
switch (kind) {
|
||||
case hipFilterModePoint:
|
||||
@@ -422,7 +422,7 @@ inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode(
|
||||
}
|
||||
}
|
||||
|
||||
inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) {
|
||||
inline static enum cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) {
|
||||
switch (kind) {
|
||||
case hipReadModeElementType:
|
||||
return cudaReadModeElementType;
|
||||
@@ -433,7 +433,7 @@ inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTex
|
||||
}
|
||||
}
|
||||
|
||||
inline static cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind(
|
||||
inline static enum cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind(
|
||||
hipChannelFormatKind kind) {
|
||||
switch (kind) {
|
||||
case hipChannelFormatKindSigned:
|
||||
@@ -510,14 +510,14 @@ inline static hipError_t hipMallocManaged(void** ptr, size_t size, unsigned int
|
||||
return hipCUDAErrorTohipError(cudaMallocManaged(ptr, size, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc,
|
||||
inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
|
||||
size_t width, size_t height,
|
||||
unsigned int flags __dparm(hipArrayDefault)) {
|
||||
return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc,
|
||||
struct hipExtent extent, unsigned int flags) {
|
||||
inline static hipError_t hipMalloc3DArray(hipArray** array, const hipChannelFormatDesc* desc,
|
||||
hipExtent extent, unsigned int flags) {
|
||||
return hipCUDAErrorTohipError(cudaMalloc3DArray(array, desc, extent, flags));
|
||||
}
|
||||
|
||||
@@ -610,6 +610,16 @@ inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes,
|
||||
cudaMemcpy(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind)));
|
||||
}
|
||||
|
||||
inline hipError_t hipMemcpyWithStream(void* dst, const void* src,
|
||||
size_t sizeBytes, hipMemcpyKind copyKind,
|
||||
hipStream_t stream) {
|
||||
cudaError_t error = cudaMemcpyAsync(dst, src, sizeBytes,
|
||||
hipMemcpyKindToCudaMemcpyKind(copyKind),
|
||||
stream);
|
||||
if (error != cudaSuccess) return hipCUDAErrorTohipError(error);
|
||||
|
||||
return hipCUDAErrorTohipError(cudaStreamSynchronize(stream));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes,
|
||||
hipMemcpyKind copyKind, hipStream_t stream __dparm(0)) {
|
||||
@@ -1048,7 +1058,6 @@ inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attribut
|
||||
return err;
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipMemGetInfo(size_t* free, size_t* total) {
|
||||
return hipCUDAErrorTohipError(cudaMemGetInfo(free, total));
|
||||
}
|
||||
@@ -1073,7 +1082,6 @@ inline static hipError_t hipEventDestroy(hipEvent_t event) {
|
||||
return hipCUDAErrorTohipError(cudaEventDestroy(event));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) {
|
||||
return hipCUDAErrorTohipError(cudaStreamCreateWithFlags(stream, flags));
|
||||
}
|
||||
@@ -1360,79 +1368,15 @@ inline static hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gri
|
||||
kernelParams, extra));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t cacheConfig) {
|
||||
return hipCUDAErrorTohipError(cudaFuncSetCacheConfig(func, cacheConfig));
|
||||
}
|
||||
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#endif
|
||||
|
||||
#ifdef __CUDACC__
|
||||
|
||||
template <class T>
|
||||
inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int* blockSize, T func,
|
||||
size_t dynamicSMemSize = 0,
|
||||
int blockSizeLimit = 0) {
|
||||
cudaError_t cerror;
|
||||
cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit);
|
||||
return hipCUDAErrorTohipError(cerror);
|
||||
}
|
||||
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTexture(size_t* offset, const struct texture<T, dim, readMode>& tex,
|
||||
const void* devPtr, size_t size = UINT_MAX) {
|
||||
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex,
|
||||
const void* devPtr, const struct hipChannelFormatDesc& desc,
|
||||
size_t size = UINT_MAX) {
|
||||
inline static hipError_t hipBindTexture(size_t* offset, struct textureReference* tex, const void* devPtr,
|
||||
const hipChannelFormatDesc* desc, size_t size __dparm(UINT_MAX)){
|
||||
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode>* tex) {
|
||||
return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode> &tex) {
|
||||
return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
|
||||
}
|
||||
|
||||
inline static hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr,
|
||||
const hipChannelFormatDesc* desc, size_t size = UINT_MAX){
|
||||
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
|
||||
hipArray_const_t array,
|
||||
const struct hipChannelFormatDesc& desc) {
|
||||
return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> *tex,
|
||||
hipArray_const_t array,
|
||||
const struct hipChannelFormatDesc* desc) {
|
||||
return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
|
||||
hipArray_const_t array) {
|
||||
return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array));
|
||||
}
|
||||
|
||||
template <class T>
|
||||
inline static hipChannelFormatDesc hipCreateChannelDesc() {
|
||||
return cudaCreateChannelDesc<T>();
|
||||
}
|
||||
|
||||
inline static hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w,
|
||||
hipChannelFormatKind f) {
|
||||
return cudaCreateChannelDesc(x, y, z, w, hipChannelFormatKindToCudaChannelFormatKind(f));
|
||||
@@ -1464,7 +1408,7 @@ inline static hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDe
|
||||
return hipCUDAErrorTohipError(cudaGetTextureObjectResourceDesc( pResDesc, textureObject));
|
||||
}
|
||||
|
||||
inline static hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* texref)
|
||||
inline static hipError_t hipGetTextureAlignmentOffset(size_t* offset, const struct textureReference* texref)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaGetTextureAlignmentOffset(offset,texref));
|
||||
}
|
||||
@@ -1473,6 +1417,70 @@ inline static hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaGetChannelDesc(desc,array));
|
||||
}
|
||||
|
||||
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#endif
|
||||
|
||||
#ifdef __CUDACC__
|
||||
|
||||
template <class T>
|
||||
inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int* blockSize, T func,
|
||||
size_t dynamicSMemSize = 0,
|
||||
int blockSizeLimit = 0) {
|
||||
cudaError_t cerror;
|
||||
cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit);
|
||||
return hipCUDAErrorTohipError(cerror);
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTexture(size_t* offset, const struct texture<T, dim, readMode>& tex,
|
||||
const void* devPtr, size_t size = UINT_MAX) {
|
||||
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex,
|
||||
const void* devPtr, const hipChannelFormatDesc& desc,
|
||||
size_t size = UINT_MAX) {
|
||||
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode>* tex) {
|
||||
return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode> &tex) {
|
||||
return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
|
||||
hipArray_const_t array,
|
||||
const hipChannelFormatDesc& desc) {
|
||||
return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> *tex,
|
||||
hipArray_const_t array,
|
||||
const hipChannelFormatDesc* desc) {
|
||||
return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
|
||||
hipArray_const_t array) {
|
||||
return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array));
|
||||
}
|
||||
|
||||
template <class T>
|
||||
inline static hipChannelFormatDesc hipCreateChannelDesc() {
|
||||
return cudaCreateChannelDesc<T>();
|
||||
}
|
||||
#endif //__CUDACC__
|
||||
|
||||
#endif // HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H
|
||||
|
||||
@@ -33,18 +33,16 @@ hipError_t hipGetDevice(int* deviceId) {
|
||||
HIP_INIT_API(hipGetDevice, deviceId);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
if (deviceId == nullptr)
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
|
||||
auto ctx = ihipGetTlsDefaultCtx();
|
||||
|
||||
if (deviceId != nullptr) {
|
||||
if (ctx == nullptr) {
|
||||
e = hipErrorInvalidDevice; // TODO, check error code.
|
||||
*deviceId = -1;
|
||||
} else {
|
||||
*deviceId = ctx->getDevice()->_deviceId;
|
||||
}
|
||||
if (ctx == nullptr) {
|
||||
e = hipErrorInvalidDevice; // TODO, check error code.
|
||||
*deviceId = -1;
|
||||
} else {
|
||||
e = hipErrorInvalidValue;
|
||||
*deviceId = ctx->getDevice()->_deviceId;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
|
||||
@@ -143,14 +143,14 @@ hipError_t hipEventDestroy(hipEvent_t event) {
|
||||
hipError_t hipEventSynchronize(hipEvent_t event) {
|
||||
HIP_INIT_SPECIAL_API(hipEventSynchronize, TRACE_SYNC, event);
|
||||
|
||||
if (!(event->_flags & hipEventReleaseToSystem)) {
|
||||
tprintf(DB_WARN,
|
||||
if (event){
|
||||
if (!(event->_flags & hipEventReleaseToSystem)) {
|
||||
tprintf(DB_WARN,
|
||||
"hipEventSynchronize on event without system-scope fence ; consider creating with "
|
||||
"hipEventReleaseToSystem\n");
|
||||
}
|
||||
auto ecd = event->locked_copyCrit();
|
||||
}
|
||||
auto ecd = event->locked_copyCrit();
|
||||
|
||||
if (event) {
|
||||
if (ecd._state == hipEventStatusUnitialized) {
|
||||
return ihipLogStatus(hipErrorInvalidResourceHandle);
|
||||
} else if (ecd._state == hipEventStatusCreated) {
|
||||
|
||||
@@ -310,6 +310,8 @@ void ihipStream_t::locked_wait() {
|
||||
hc::completion_future marker;
|
||||
{
|
||||
LockedAccessor_StreamCrit_t crit(_criticalData);
|
||||
// skipping marker since stream is empty
|
||||
if (crit->_av.get_is_empty()) return;
|
||||
marker = crit->_av.create_marker(hc::no_scope);
|
||||
}
|
||||
|
||||
@@ -1497,21 +1499,18 @@ hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream) {
|
||||
return e;
|
||||
}
|
||||
|
||||
bool ihipStreamCallbackHandler(hsa_signal_value_t value, void* cbArgs) {
|
||||
void ihipStreamCallbackHandler(ihipStreamCallback_t* cb) {
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
ihipStreamCallback_t* cb = static_cast<ihipStreamCallback_t*> (cbArgs);
|
||||
|
||||
if(cb->comFuture.valid())
|
||||
cb->comFuture.wait();
|
||||
// Synchronize stream
|
||||
tprintf(DB_SYNC, "ihipStreamCallbackHandler wait on stream %s\n",
|
||||
ToString(cb->_stream).c_str());
|
||||
GET_TLS();
|
||||
e = ihipStreamSynchronize(tls, cb->_stream);
|
||||
|
||||
// Call registered callback function
|
||||
cb->_callback(cb->_stream, e, cb->_userData);
|
||||
|
||||
hsa_signal_store_screlease(cb->_signal,0);
|
||||
|
||||
delete cb;
|
||||
return false;
|
||||
}
|
||||
|
||||
//---
|
||||
|
||||
@@ -650,9 +650,7 @@ class ihipStreamCallback_t {
|
||||
: _stream(stream), _callback(callback), _userData(userData) {
|
||||
};
|
||||
hipStream_t _stream;
|
||||
hsa_signal_t _signal;
|
||||
hipStreamCallback_t _callback;
|
||||
hc::completion_future comFuture;
|
||||
void* _userData;
|
||||
};
|
||||
|
||||
@@ -970,7 +968,7 @@ hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod,
|
||||
|
||||
hipStream_t ihipSyncAndResolveStream(hipStream_t, bool lockAcquired = 0);
|
||||
hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream);
|
||||
bool ihipStreamCallbackHandler(hsa_signal_value_t value, void* cbArgs);
|
||||
void ihipStreamCallbackHandler(ihipStreamCallback_t* cb);
|
||||
|
||||
// Stream printf functions:
|
||||
inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s) {
|
||||
|
||||
@@ -27,6 +27,7 @@ THE SOFTWARE.
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
|
||||
#include <functional>
|
||||
#include <fstream>
|
||||
|
||||
__device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
|
||||
@@ -35,23 +36,369 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
|
||||
// Internal HIP APIS:
|
||||
namespace hip_internal {
|
||||
|
||||
hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
|
||||
hipStream_t stream) {
|
||||
hipError_t e = hipSuccess;
|
||||
namespace {
|
||||
inline
|
||||
const char* hsa_to_string(hsa_status_t err) noexcept
|
||||
{
|
||||
const char* r{};
|
||||
|
||||
// Return success if number of bytes to copy is 0
|
||||
if (sizeBytes == 0) return e;
|
||||
if (hsa_status_string(err, &r) != HSA_STATUS_SUCCESS) return r;
|
||||
|
||||
return "Unknown.";
|
||||
}
|
||||
|
||||
template<std::size_t m, std::size_t n>
|
||||
inline
|
||||
void throwing_result_check(hsa_status_t res, const char (&file)[m],
|
||||
const char (&function)[n], int line) {
|
||||
if (res == HSA_STATUS_SUCCESS) return;
|
||||
if (res == HSA_STATUS_INFO_BREAK) return;
|
||||
|
||||
throw std::runtime_error{"Failed in file " + (file +
|
||||
(", in function \"" + (function +
|
||||
("\", on line " + std::to_string(line))))) +
|
||||
", with error: " + hsa_to_string(res)};
|
||||
}
|
||||
|
||||
inline
|
||||
hsa_agent_t cpu_agent() noexcept {
|
||||
static hsa_agent_t cpu{[]() {
|
||||
hsa_agent_t r{};
|
||||
throwing_result_check(
|
||||
hsa_iterate_agents([](hsa_agent_t x, void* pr) {
|
||||
hsa_device_type_t t{};
|
||||
hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &t);
|
||||
|
||||
if (t != HSA_DEVICE_TYPE_CPU) return HSA_STATUS_SUCCESS;
|
||||
|
||||
*static_cast<hsa_agent_t *>(pr) = x;
|
||||
|
||||
return HSA_STATUS_INFO_BREAK;
|
||||
}, &r), __FILE__, __func__, __LINE__);
|
||||
|
||||
return r;
|
||||
}()};
|
||||
|
||||
return cpu;
|
||||
}
|
||||
|
||||
inline
|
||||
hsa_device_type_t type(hsa_agent_t x) noexcept
|
||||
{
|
||||
hsa_device_type_t r{};
|
||||
throwing_result_check(hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &r),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
const auto is_large_BAR{[](){
|
||||
std::unique_ptr<void, void (*)(void*)> hsa{
|
||||
(throwing_result_check(hsa_init(), __FILE__, __func__, __LINE__),
|
||||
nullptr),
|
||||
[](void*) { hsa_shut_down(); }};
|
||||
bool r{true};
|
||||
|
||||
throwing_result_check(hsa_iterate_agents([](hsa_agent_t x, void* pr) {
|
||||
if (x.handle == cpu_agent().handle) return HSA_STATUS_SUCCESS;
|
||||
|
||||
throwing_result_check(
|
||||
hsa_agent_iterate_regions(x, [](hsa_region_t y, void* p) {
|
||||
hsa_region_segment_t seg{};
|
||||
throwing_result_check(
|
||||
hsa_region_get_info(y, HSA_REGION_INFO_SEGMENT, &seg),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
if (seg != HSA_REGION_SEGMENT_GLOBAL) {
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
uint32_t flags{};
|
||||
throwing_result_check(hsa_region_get_info(
|
||||
y, HSA_REGION_INFO_GLOBAL_FLAGS, &flags),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) {
|
||||
hsa_amd_memory_pool_access_t tmp{};
|
||||
throwing_result_check(
|
||||
hsa_amd_agent_memory_pool_get_info(
|
||||
cpu_agent(),
|
||||
hsa_amd_memory_pool_t{y.handle},
|
||||
HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS,
|
||||
&tmp),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
*static_cast<bool*>(p) &=
|
||||
tmp != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED;
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}, pr), __FILE__, __func__, __LINE__);
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}, &r), __FILE__, __func__, __LINE__);
|
||||
|
||||
return r;
|
||||
}()};
|
||||
|
||||
inline
|
||||
hsa_amd_pointer_info_t info(const void* p) noexcept
|
||||
{
|
||||
hsa_amd_pointer_info_t r{sizeof(hsa_amd_pointer_info_t)};
|
||||
throwing_result_check(
|
||||
hsa_amd_pointer_info(
|
||||
const_cast<void*>(p), &r, nullptr, nullptr, nullptr),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
r.size = is_large_BAR || (type(r.agentOwner) == HSA_DEVICE_TYPE_CPU) ?
|
||||
UINT32_MAX : sizeof(hsa_amd_pointer_info_t);
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
constexpr size_t staging_sz{4 * 1024 * 1024}; // 2 Pages.
|
||||
|
||||
thread_local const std::unique_ptr<void, void (*)(void *)> staging_buffer{
|
||||
[]() {
|
||||
hsa_region_t r{};
|
||||
throwing_result_check(hsa_agent_iterate_regions(
|
||||
cpu_agent(), [](hsa_region_t x, void *p) {
|
||||
hsa_region_segment_t seg{};
|
||||
throwing_result_check(
|
||||
hsa_region_get_info(x, HSA_REGION_INFO_SEGMENT, &seg),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
if (seg != HSA_REGION_SEGMENT_GLOBAL) return HSA_STATUS_SUCCESS;
|
||||
|
||||
uint32_t flags{};
|
||||
throwing_result_check(hsa_region_get_info(
|
||||
x, HSA_REGION_INFO_GLOBAL_FLAGS, &flags),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) {
|
||||
*static_cast<hsa_region_t *>(p) = x;
|
||||
|
||||
return HSA_STATUS_INFO_BREAK;
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}, &r), __FILE__, __func__, __LINE__);
|
||||
|
||||
void *tp{};
|
||||
throwing_result_check(hsa_memory_allocate(r, staging_sz, &tp),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
return tp;
|
||||
}(),
|
||||
[](void *ptr) { hsa_memory_free(ptr); }};
|
||||
|
||||
thread_local hsa_signal_t copy_signal{[]() {
|
||||
hsa_agent_t cpu{cpu_agent()};
|
||||
hsa_signal_t sgn{};
|
||||
throwing_result_check(hsa_signal_create(1, 1, &cpu, &sgn),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
return sgn;
|
||||
}()};
|
||||
} // Unnamed namespace.
|
||||
|
||||
inline
|
||||
void do_copy(void* __restrict dst, const void* __restrict src, std::size_t n,
|
||||
hsa_agent_t da, hsa_agent_t sa) {
|
||||
hsa_signal_silent_store_relaxed(copy_signal, 1);
|
||||
throwing_result_check(
|
||||
hsa_amd_memory_async_copy(dst, da, src, sa, n, 0, nullptr, copy_signal),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
while (hsa_signal_wait_relaxed(copy_signal, HSA_SIGNAL_CONDITION_EQ, 0,
|
||||
UINT64_MAX, HSA_WAIT_STATE_ACTIVE));
|
||||
}
|
||||
|
||||
inline
|
||||
void do_std_memcpy(
|
||||
void* __restrict dst, const void* __restrict src, std::size_t n) {
|
||||
std::memcpy(dst, src, n);
|
||||
|
||||
return std::atomic_thread_fence(std::memory_order_seq_cst);
|
||||
}
|
||||
|
||||
inline
|
||||
void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n,
|
||||
hsa_amd_pointer_info_t si) {
|
||||
if (si.size == UINT32_MAX) return do_std_memcpy(dst, src, n);
|
||||
|
||||
const auto di{info(dst)};
|
||||
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
|
||||
dst = static_cast<char*>(di.agentBaseAddress) +
|
||||
(static_cast<char*>(dst) -
|
||||
static_cast<char*>(di.hostBaseAddress));
|
||||
do_copy(dst, src, n, si.agentOwner, si.agentOwner);
|
||||
}
|
||||
else if (n <= staging_sz) {
|
||||
do_copy(staging_buffer.get(), src, n, si.agentOwner, si.agentOwner);
|
||||
std::memcpy(dst, staging_buffer.get(), n);
|
||||
}
|
||||
else {
|
||||
std::unique_ptr<void, void (*)(void*)> lck{
|
||||
dst, [](void* p) { hsa_amd_memory_unlock(p); }};
|
||||
|
||||
throwing_result_check(hsa_amd_memory_lock(dst, n, &si.agentOwner, 1,
|
||||
const_cast<void**>(&dst)),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
do_copy(dst, src, n, si.agentOwner, si.agentOwner);
|
||||
}
|
||||
}
|
||||
|
||||
inline
|
||||
void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n,
|
||||
hsa_amd_pointer_info_t di) {
|
||||
if (di.size == UINT32_MAX) return do_std_memcpy(dst, src, n);
|
||||
|
||||
const auto si{info(const_cast<void*>(src))};
|
||||
|
||||
if (si.type == HSA_EXT_POINTER_TYPE_LOCKED) {
|
||||
src = static_cast<char*>(si.agentBaseAddress) +
|
||||
(static_cast<const char*>(src) -
|
||||
static_cast<char*>(si.hostBaseAddress));
|
||||
do_copy(dst, src, n, di.agentOwner, di.agentOwner);
|
||||
}
|
||||
else if (n <= staging_sz) {
|
||||
std::memcpy(staging_buffer.get(), src, n);
|
||||
do_copy(dst, staging_buffer.get(), n, di.agentOwner, di.agentOwner);
|
||||
}
|
||||
else {
|
||||
std::unique_ptr<void, void (*)(void*)> lck{
|
||||
const_cast<void*>(src), [](void* p) { hsa_amd_memory_unlock(p); }};
|
||||
|
||||
throwing_result_check(hsa_amd_memory_lock(const_cast<void*>(src), n,
|
||||
&di.agentOwner, 1,
|
||||
const_cast<void**>(&src)),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
do_copy(dst, src, n, di.agentOwner, di.agentOwner);
|
||||
}
|
||||
}
|
||||
|
||||
inline
|
||||
void generic_copy(void* __restrict dst, const void* __restrict src, size_t n,
|
||||
hsa_amd_pointer_info_t di, hsa_amd_pointer_info_t si) {
|
||||
if (di.size == UINT32_MAX && si.size == UINT32_MAX) {
|
||||
return do_std_memcpy(dst, src, n);
|
||||
}
|
||||
|
||||
std::unique_ptr<void, void (*)(void*)> lck0{
|
||||
nullptr, [](void* p) { hsa_amd_memory_unlock(p); }};
|
||||
std::unique_ptr<void, void (*)(void*)> lck1{nullptr, lck0.get_deleter()};
|
||||
|
||||
switch (si.type) {
|
||||
case HSA_EXT_POINTER_TYPE_HSA:
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_HSA) {
|
||||
hsa_memory_copy(dst, src, n);
|
||||
return; // TODO: do_copy(dst, src, n, di.agentOwner, si.agentOwner);
|
||||
}
|
||||
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN ||
|
||||
di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
|
||||
return d2h_copy(dst, src, n, si);
|
||||
}
|
||||
break;
|
||||
case HSA_EXT_POINTER_TYPE_LOCKED:
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
|
||||
std::memcpy(dst, si.hostBaseAddress, n);
|
||||
|
||||
return;
|
||||
}
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
|
||||
std::memcpy(di.hostBaseAddress, si.hostBaseAddress, n);
|
||||
|
||||
return;
|
||||
}
|
||||
src = si.agentBaseAddress;
|
||||
si.agentOwner = di.agentOwner;
|
||||
break;
|
||||
case HSA_EXT_POINTER_TYPE_UNKNOWN:
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
|
||||
std::memcpy(dst, src, n);
|
||||
|
||||
return;
|
||||
}
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
|
||||
std::memcpy(di.hostBaseAddress, src, n);
|
||||
|
||||
return;
|
||||
}
|
||||
return h2d_copy(dst, src, n, di);
|
||||
default: do_copy(dst, src, n, di.agentOwner, si.agentOwner); break;
|
||||
}
|
||||
}
|
||||
|
||||
inline
|
||||
void memcpy_impl(void* __restrict dst, const void* __restrict src, size_t n,
|
||||
hipMemcpyKind k) noexcept {
|
||||
switch (k) {
|
||||
case hipMemcpyHostToHost: std::memcpy(dst, src, n); break;
|
||||
case hipMemcpyHostToDevice:
|
||||
return is_large_BAR ? do_std_memcpy(dst, src, n)
|
||||
: h2d_copy(dst, src, n, info(dst));
|
||||
case hipMemcpyDeviceToHost:
|
||||
return is_large_BAR ? do_std_memcpy(dst, src, n)
|
||||
: d2h_copy(dst, src, n, info(src));
|
||||
case hipMemcpyDeviceToDevice:
|
||||
return do_copy(dst, src, n, info(dst).agentOwner,
|
||||
info(const_cast<void*>(src)).agentOwner);
|
||||
default: return generic_copy(dst, src, n, info(dst), info(src));
|
||||
}
|
||||
}
|
||||
|
||||
hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes,
|
||||
hipMemcpyKind kind, hipStream_t stream) {
|
||||
if (sizeBytes == 0) return hipSuccess;
|
||||
if (!dst || !src) return hipErrorInvalidValue;
|
||||
|
||||
if (!(stream = ihipSyncAndResolveStream(stream))) {
|
||||
return hipErrorInvalidValue;
|
||||
}
|
||||
|
||||
try {
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
|
||||
if (!stream) return hipErrorInvalidValue;
|
||||
|
||||
stream->locked_copyAsync(dst, src, sizeBytes, kind);
|
||||
}
|
||||
catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
catch (const ihipException& ex) {
|
||||
return ex._code;
|
||||
}
|
||||
catch (const std::exception& ex) {
|
||||
std::cerr << ex.what() << std::endl;
|
||||
throw;
|
||||
}
|
||||
catch (...) {
|
||||
return hipErrorUnknown;
|
||||
}
|
||||
|
||||
return hipSuccess;
|
||||
}
|
||||
|
||||
hipError_t memcpySync(void* dst, const void* src, size_t sizeBytes,
|
||||
hipMemcpyKind kind, hipStream_t stream) {
|
||||
if (sizeBytes == 0) return hipSuccess;
|
||||
if (!dst || !src) return hipErrorInvalidValue;
|
||||
|
||||
try {
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
|
||||
if (!stream) return hipErrorInvalidValue;
|
||||
|
||||
LockedAccessor_StreamCrit_t cs{stream->criticalData()};
|
||||
cs->_av.wait();
|
||||
|
||||
memcpy_impl(dst, src, sizeBytes, kind);
|
||||
}
|
||||
catch (const ihipException& ex) {
|
||||
return ex._code;
|
||||
}
|
||||
catch (const std::exception& ex) {
|
||||
std::cerr << ex.what() << std::endl;
|
||||
throw;
|
||||
}
|
||||
catch (...) {
|
||||
return hipErrorUnknown;
|
||||
@@ -896,20 +1243,8 @@ hipError_t hipMemcpyToSymbol(void* dst, const void* src, size_t count,
|
||||
|
||||
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst);
|
||||
|
||||
if (dst == nullptr) {
|
||||
return ihipLogStatus(hipErrorInvalidSymbol);
|
||||
}
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
if (kind == hipMemcpyHostToDevice || kind == hipMemcpyDefault ||
|
||||
kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) {
|
||||
stream->locked_copySync((char*)dst+offset, (void*)src, count, kind, false);
|
||||
} else {
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
return ihipLogStatus(
|
||||
hipMemcpy(static_cast<char*>(dst) + offset, src, count, kind));
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count,
|
||||
@@ -920,20 +1255,8 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count,
|
||||
|
||||
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst);
|
||||
|
||||
if (dst == nullptr) {
|
||||
return ihipLogStatus(hipErrorInvalidSymbol);
|
||||
}
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
if (kind == hipMemcpyDefault || kind == hipMemcpyDeviceToHost ||
|
||||
kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) {
|
||||
stream->locked_copySync((void*)dst, (char*)src+offset, count, kind, false);
|
||||
} else {
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
return ihipLogStatus(
|
||||
hipMemcpy(dst, static_cast<const char*>(src) + offset, count, kind));
|
||||
}
|
||||
|
||||
|
||||
@@ -995,120 +1318,49 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* src, size_t count,
|
||||
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpy, (TRACE_MCMD), dst, src, sizeBytes, kind);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
// Return success if number of bytes to copy is 0
|
||||
if (sizeBytes == 0) return ihipLogStatus(e);
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
|
||||
if(dst==NULL || src==NULL)
|
||||
{
|
||||
e=hipErrorInvalidValue;
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
try {
|
||||
stream->locked_copySync(dst, src, sizeBytes, kind);
|
||||
} catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind,
|
||||
hipStreamNull));
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpyHtoD, (TRACE_MCMD), dst, src, sizeBytes);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
if (sizeBytes == 0) return ihipLogStatus(e);
|
||||
|
||||
if(dst==NULL || src==NULL){
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
|
||||
try {
|
||||
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToDevice, false);
|
||||
} catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
|
||||
hipMemcpyHostToDevice,
|
||||
hipStreamNull));
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpyDtoH, (TRACE_MCMD), dst, src, sizeBytes);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
if (sizeBytes == 0) return ihipLogStatus(e);
|
||||
|
||||
if(dst==NULL || src==NULL){
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
|
||||
try {
|
||||
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToHost, false);
|
||||
} catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
|
||||
hipMemcpyDeviceToHost,
|
||||
hipStreamNull));
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
if (sizeBytes == 0) return ihipLogStatus(e);
|
||||
|
||||
if(dst==NULL || src==NULL){
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
|
||||
try {
|
||||
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToDevice, false);
|
||||
} catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
|
||||
hipMemcpyDeviceToDevice,
|
||||
hipStreamNull));
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpyHtoH, (TRACE_MCMD), dst, src, sizeBytes);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
if (sizeBytes == 0) return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
|
||||
hipMemcpyHostToHost,
|
||||
hipStreamNull));
|
||||
}
|
||||
|
||||
if(dst==NULL || src==NULL){
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
hipError_t hipMemcpyWithStream(void* dst, void* src, size_t sizeBytes,
|
||||
hipMemcpyKind kind, hipStream_t stream) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpyWithStream, (TRACE_MCMD), dst, src, sizeBytes,
|
||||
kind, stream);
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
try {
|
||||
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToHost, false);
|
||||
} catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind,
|
||||
stream));
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
|
||||
|
||||
@@ -31,7 +31,7 @@ THE SOFTWARE.
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
// Stream
|
||||
//
|
||||
#if defined(__HCC__) && (__hcc_minor__ < 3)
|
||||
#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
|
||||
enum queue_priority
|
||||
{
|
||||
priority_high = 0,
|
||||
@@ -73,7 +73,7 @@ hipError_t ihipStreamCreate(TlsData *tls, hipStream_t* stream, unsigned int flag
|
||||
// Obtain mutex access to the device critical data, release by destructor
|
||||
LockedAccessor_CtxCrit_t ctxCrit(ctx->criticalData());
|
||||
|
||||
#if defined(__HCC__) && (__hcc_minor__ < 3)
|
||||
#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
|
||||
auto istream = new ihipStream_t(ctx, acc.create_view(), flags);
|
||||
#else
|
||||
auto istream = new ihipStream_t(ctx, acc.create_view(Kalmar::execute_any_order, Kalmar::queuing_mode_automatic, (Kalmar::queue_priority)priority), flags);
|
||||
@@ -242,7 +242,7 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
|
||||
} else if (stream == hipStreamNull) {
|
||||
return ihipLogStatus(hipErrorInvalidResourceHandle);
|
||||
} else {
|
||||
#if defined(__HCC__) && (__hcc_minor__ < 3)
|
||||
#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
|
||||
*priority = 0;
|
||||
#else
|
||||
LockedAccessor_StreamCrit_t crit(stream->criticalData());
|
||||
@@ -254,137 +254,14 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
|
||||
|
||||
|
||||
//---
|
||||
void setCallbackPacket(hsa_queue_t* queue,
|
||||
uint64_t& index, uint64_t& nextIndex,
|
||||
hsa_barrier_and_packet_t** barrier1,
|
||||
hsa_barrier_and_packet_t** barrier2){
|
||||
|
||||
uint64_t tempIndex = 0;
|
||||
uint32_t mask = queue->size - 1;
|
||||
hsa_barrier_and_packet_t* tempBarrier;
|
||||
|
||||
// Check for empty packets
|
||||
do{
|
||||
tempIndex = hsa_queue_load_write_index_scacquire(queue);
|
||||
tempBarrier = &(((hsa_barrier_and_packet_t*)(queue->base_address))[tempIndex & mask]);
|
||||
}while(!(tempBarrier->header & HSA_PACKET_TYPE_INVALID));
|
||||
|
||||
// Reserve two packets for two barriers
|
||||
index = hsa_queue_add_write_index_scacquire(queue, 2);
|
||||
|
||||
if(index > mask)
|
||||
{
|
||||
index = 0;
|
||||
nextIndex = 1;
|
||||
}
|
||||
else if(index == mask)
|
||||
nextIndex = 0;
|
||||
else
|
||||
nextIndex = index + 1;
|
||||
|
||||
tempBarrier = new hsa_barrier_and_packet_t;
|
||||
memset(tempBarrier, 0, sizeof(hsa_barrier_and_packet_t));
|
||||
tempBarrier->header = HSA_PACKET_TYPE_INVALID;
|
||||
|
||||
// Barrier 1
|
||||
*barrier1 = &(((hsa_barrier_and_packet_t*)(queue->base_address))[index & mask]);
|
||||
memcpy(*barrier1,tempBarrier,sizeof(hsa_barrier_and_packet_t));
|
||||
|
||||
// Barrier 2
|
||||
*barrier2 = &(((hsa_barrier_and_packet_t*)(queue->base_address))[nextIndex & mask]);
|
||||
memcpy(*barrier2,tempBarrier,sizeof(hsa_barrier_and_packet_t));
|
||||
|
||||
delete tempBarrier;
|
||||
}
|
||||
|
||||
hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
|
||||
unsigned int flags) {
|
||||
|
||||
HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags);
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
if(stream == hipStreamNull)
|
||||
{
|
||||
ihipCtx_t* device = ihipGetTlsDefaultCtx();
|
||||
stream = device->_defaultStream;
|
||||
}
|
||||
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
|
||||
// Lock the stream
|
||||
LockedAccessor_StreamCrit_t crit(stream->criticalData());
|
||||
|
||||
// Device synchronization
|
||||
hc::completion_future marker = crit->_av.create_marker(hc::system_scope);
|
||||
|
||||
// 1. Lock the queue
|
||||
hsa_queue_t* lockedQ = static_cast<hsa_queue_t*> (crit->_av.acquire_locked_hsa_queue());
|
||||
|
||||
if(lockedQ == nullptr)
|
||||
{
|
||||
// No queue attached to stream hence exiting early
|
||||
return ihipLogStatus(hipErrorMissingConfiguration);
|
||||
}
|
||||
|
||||
// 2. Allocate a singals
|
||||
hsa_signal_t signal;
|
||||
hsa_status_t status = hsa_signal_create(1, 0, NULL, &signal);
|
||||
|
||||
if(status != HSA_STATUS_SUCCESS)
|
||||
{
|
||||
crit->_av.release_locked_hsa_queue();
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
hsa_signal_t depSignal;
|
||||
status = hsa_signal_create(1, 0, NULL, &depSignal);
|
||||
|
||||
if(status != HSA_STATUS_SUCCESS)
|
||||
{
|
||||
crit->_av.release_locked_hsa_queue();
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
// 3. Store callback details, will destroy allocation in callback handler
|
||||
// Create a thread in detached mode to handle callback
|
||||
ihipStreamCallback_t* cb = new ihipStreamCallback_t(stream, callback, userData);
|
||||
if(cb == nullptr)
|
||||
{
|
||||
crit->_av.release_locked_hsa_queue();
|
||||
return ihipLogStatus(hipErrorMemoryAllocation);
|
||||
}
|
||||
cb->_signal = depSignal;
|
||||
cb->comFuture = marker ;
|
||||
|
||||
// 4. Create barrier packets
|
||||
uint64_t index ;
|
||||
uint64_t nextIndex;
|
||||
|
||||
hsa_barrier_and_packet_t* barrier;
|
||||
hsa_barrier_and_packet_t* depBarrier;
|
||||
|
||||
setCallbackPacket(lockedQ, index, nextIndex, &barrier, &depBarrier);
|
||||
|
||||
barrier->completion_signal = signal;
|
||||
|
||||
depBarrier->dep_signal[0] = depSignal;
|
||||
|
||||
uint16_t header = (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE)| 1 << HSA_PACKET_HEADER_BARRIER;
|
||||
|
||||
// 5. Update packet header,
|
||||
// Intentionally updated second barrier header before first in order to avoid race
|
||||
depBarrier->header = header;
|
||||
barrier->header = header;
|
||||
|
||||
// 6. Trigger the doorbell
|
||||
nextIndex = nextIndex + 1;
|
||||
hsa_queue_store_write_index_screlease(lockedQ, nextIndex);
|
||||
hsa_signal_store_relaxed(lockedQ->doorbell_signal, index+1);
|
||||
|
||||
// 7. Release queue
|
||||
crit->_av.release_locked_hsa_queue();
|
||||
|
||||
// 8. Register signal callback
|
||||
hsa_amd_signal_async_handler(signal, HSA_SIGNAL_CONDITION_EQ, 0, ihipStreamCallbackHandler, cb);
|
||||
std::thread(ihipStreamCallbackHandler, cb).detach();
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
@@ -0,0 +1,31 @@
|
||||
/*
|
||||
Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "is_callable_test.hpp"
|
||||
#include <test_common.h>
|
||||
|
||||
int main() { passed(); }
|
||||
@@ -0,0 +1,34 @@
|
||||
/*
|
||||
Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
// Ensure fallback path is chosen
|
||||
#define HIP_HAS_RESULT_OF_SFINAE 0
|
||||
#define HIP_HAS_INVOCABLE 0
|
||||
#include "is_callable_test.hpp"
|
||||
#include <test_common.h>
|
||||
|
||||
int main() { passed(); }
|
||||
@@ -0,0 +1,31 @@
|
||||
/*
|
||||
Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++14 EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "is_callable_test.hpp"
|
||||
#include <test_common.h>
|
||||
|
||||
int main() { passed(); }
|
||||
@@ -0,0 +1,31 @@
|
||||
/*
|
||||
Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++17 EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "is_callable_test.hpp"
|
||||
#include <test_common.h>
|
||||
|
||||
int main() { passed(); }
|
||||
@@ -0,0 +1,209 @@
|
||||
/*
|
||||
Copyright (c) 2017 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip/hcc_detail/helpers.hpp>
|
||||
#include <memory>
|
||||
|
||||
using hip_impl::is_callable;
|
||||
|
||||
template<int N>
|
||||
struct callable_rank : callable_rank<N-1>
|
||||
{};
|
||||
|
||||
template<>
|
||||
struct callable_rank<0>
|
||||
{};
|
||||
|
||||
struct test1
|
||||
{
|
||||
struct is_callable_class
|
||||
{
|
||||
void operator()(int) const
|
||||
{
|
||||
}
|
||||
};
|
||||
struct callable_test_param {};
|
||||
|
||||
void is_callable_function(int)
|
||||
{
|
||||
}
|
||||
|
||||
struct is_callable_rank_class
|
||||
{
|
||||
void operator()(int, callable_rank<3>) const
|
||||
{
|
||||
}
|
||||
|
||||
void operator()(int, callable_rank<4>) const
|
||||
{
|
||||
}
|
||||
};
|
||||
|
||||
static_assert(is_callable<is_callable_class(int)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_class(long)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_class(double)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_class(const int&)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_class(const long&)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_class(const double&)>::value, "Not callable");
|
||||
static_assert(not is_callable<is_callable_class(callable_test_param)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_class()>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_class(int, int)>::value, "callable failed");
|
||||
|
||||
typedef void (*is_callable_function_pointer)(int);
|
||||
static_assert(is_callable<is_callable_function_pointer(int)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_function_pointer(long)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_function_pointer(double)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_function_pointer(const int&)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_function_pointer(const long&)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_function_pointer(const double&)>::value, "Not callable");
|
||||
static_assert(not is_callable<is_callable_function_pointer(callable_test_param)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_function_pointer()>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_function_pointer(int, int)>::value, "callable failed");
|
||||
|
||||
static_assert(is_callable<is_callable_rank_class(int, callable_rank<3>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(long, callable_rank<3>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(double, callable_rank<3>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const int&, callable_rank<3>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const long&, callable_rank<3>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const double&, callable_rank<3>)>::value, "Not callable");
|
||||
|
||||
static_assert(is_callable<is_callable_rank_class(int, callable_rank<4>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(long, callable_rank<4>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(double, callable_rank<4>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const int&, callable_rank<4>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const long&, callable_rank<4>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const double&, callable_rank<4>)>::value, "Not callable");
|
||||
|
||||
static_assert(is_callable<is_callable_rank_class(int, callable_rank<5>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(long, callable_rank<5>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(double, callable_rank<5>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const int&, callable_rank<5>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const long&, callable_rank<5>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const double&, callable_rank<5>)>::value, "Not callable");
|
||||
|
||||
static_assert(is_callable<is_callable_rank_class(int, callable_rank<6>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(long, callable_rank<6>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(double, callable_rank<6>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const int&, callable_rank<6>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const long&, callable_rank<6>)>::value, "Not callable");
|
||||
static_assert(is_callable<is_callable_rank_class(const double&, callable_rank<6>)>::value, "Not callable");
|
||||
|
||||
static_assert(not is_callable<is_callable_rank_class(int, callable_rank<1>)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_rank_class(long, callable_rank<1>)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_rank_class(double, callable_rank<1>)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_rank_class(const int&, callable_rank<1>)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_rank_class(const long&, callable_rank<1>)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_rank_class(const double&, callable_rank<1>)>::value, "callable failed");
|
||||
|
||||
static_assert(not is_callable<is_callable_rank_class(callable_test_param, callable_test_param)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_rank_class(callable_rank<3>, callable_test_param)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_rank_class(callable_rank<4>, callable_test_param)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_rank_class(callable_test_param, callable_rank<3>)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_rank_class(callable_test_param, callable_rank<4>)>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_rank_class()>::value, "callable failed");
|
||||
static_assert(not is_callable<is_callable_rank_class(int, int)>::value, "callable failed");
|
||||
};
|
||||
|
||||
struct test2
|
||||
{
|
||||
typedef int(callable_rank<0>::*fn)(int);
|
||||
|
||||
static_assert(is_callable<fn(callable_rank<0>&, int)>::value, "Failed");
|
||||
static_assert(is_callable<fn(callable_rank<1>&, int)>::value, "Failed");
|
||||
static_assert(not is_callable<fn(callable_rank<0>&)>::value, "Failed");
|
||||
static_assert(not is_callable<fn(callable_rank<0> const&, int)>::value, "Failed");
|
||||
};
|
||||
|
||||
struct test3
|
||||
{
|
||||
typedef int(callable_rank<0>::*fn)(int);
|
||||
|
||||
typedef callable_rank<0>* T;
|
||||
typedef callable_rank<1>* DT;
|
||||
typedef const callable_rank<0>* CT;
|
||||
typedef std::unique_ptr<callable_rank<0>> ST;
|
||||
|
||||
static_assert(is_callable<fn(T&, int)>::value, "Failed");
|
||||
static_assert(is_callable<fn(DT&, int)>::value, "Failed");
|
||||
static_assert(is_callable<fn(const T&, int)>::value, "Failed");
|
||||
static_assert(is_callable<fn(T&&, int)>::value, "Failed");
|
||||
static_assert(is_callable<fn(ST, int)>::value, "Failed");
|
||||
static_assert(not is_callable<fn(CT&, int)>::value, "Failed");
|
||||
|
||||
};
|
||||
|
||||
struct test4
|
||||
{
|
||||
typedef int(callable_rank<0>::*fn);
|
||||
|
||||
static_assert(not is_callable<fn()>::value, "Failed");
|
||||
};
|
||||
|
||||
struct test5
|
||||
{
|
||||
typedef int(callable_rank<0>::*fn);
|
||||
|
||||
static_assert(is_callable<fn(callable_rank<0>&)>::value, "Failed");
|
||||
static_assert(is_callable<fn(callable_rank<0>&&)>::value, "Failed");
|
||||
static_assert(is_callable<fn(const callable_rank<0>&)>::value, "Failed");
|
||||
static_assert(is_callable<fn(callable_rank<1>&)>::value, "Failed");
|
||||
};
|
||||
|
||||
struct test6
|
||||
{
|
||||
typedef int(callable_rank<0>::*fn);
|
||||
|
||||
typedef callable_rank<0>* T;
|
||||
typedef callable_rank<1>* DT;
|
||||
typedef const callable_rank<0>* CT;
|
||||
typedef std::unique_ptr<callable_rank<0>> ST;
|
||||
|
||||
static_assert(is_callable<fn(T&)>::value, "Failed");
|
||||
static_assert(is_callable<fn(DT&)>::value, "Failed");
|
||||
static_assert(is_callable<fn(const T&)>::value, "Failed");
|
||||
static_assert(is_callable<fn(T&&)>::value, "Failed");
|
||||
static_assert(is_callable<fn(ST)>::value, "Failed");
|
||||
static_assert(is_callable<fn(CT&)>::value, "Failed");
|
||||
|
||||
};
|
||||
|
||||
struct test7
|
||||
{
|
||||
typedef void(*fp)(callable_rank<0>&, int);
|
||||
|
||||
static_assert(is_callable<fp(callable_rank<0>&, int)>::value, "Failed");
|
||||
static_assert(is_callable<fp(callable_rank<1>&, int)>::value, "Failed");
|
||||
static_assert(not is_callable<fp(const callable_rank<0>&, int)>::value, "Failed");
|
||||
static_assert(not is_callable<fp()>::value, "Failed");
|
||||
static_assert(not is_callable<fp(callable_rank<0>&)>::value, "Failed");
|
||||
};
|
||||
|
||||
struct test8
|
||||
{
|
||||
typedef void(&fp)(callable_rank<0>&, int);
|
||||
|
||||
static_assert(is_callable<fp(callable_rank<0>&, int)>::value, "Failed");
|
||||
static_assert(is_callable<fp(callable_rank<1>&, int)>::value, "Failed");
|
||||
static_assert(not is_callable<fp(const callable_rank<0>&, int)>::value, "Failed");
|
||||
static_assert(not is_callable<fp()>::value, "Failed");
|
||||
static_assert(not is_callable<fp(callable_rank<0>&)>::value, "Failed");
|
||||
};
|
||||
@@ -34,6 +34,7 @@ THE SOFTWARE.
|
||||
#include <assert.h>
|
||||
#include <iostream>
|
||||
#include <string>
|
||||
#include <sstream>
|
||||
#include <type_traits>
|
||||
|
||||
using namespace std;
|
||||
@@ -157,6 +158,24 @@ bool TestVectorType() {
|
||||
if (f1 == f2) return false;
|
||||
if (!(f1 != f2)) return false;
|
||||
|
||||
using T = typename V::value_type;
|
||||
|
||||
const T& x = f1.x;
|
||||
T& y = f2.x;
|
||||
const volatile T& z = f3.x;
|
||||
volatile T& w = f2.x;
|
||||
|
||||
if (x != T{3}) return false;
|
||||
if (y != T{4}) return false;
|
||||
if (z != T{3}) return false;
|
||||
if (w != T{4}) return false;
|
||||
|
||||
stringstream str;
|
||||
str << f1.x;
|
||||
str >> f2.x;
|
||||
|
||||
if (f1.x != f2.x) return false;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
|
||||
@@ -149,6 +149,20 @@ bool TestVectorType() {
|
||||
if (f1 == f2) return false;
|
||||
if (!(f1 != f2)) return false;
|
||||
|
||||
#if 0 // TODO: investigate on GFX8
|
||||
using T = typename V::value_type;
|
||||
|
||||
const T& x = f1.x;
|
||||
T& y = f2.x;
|
||||
const volatile T& z = f3.x;
|
||||
volatile T& w = f2.x;
|
||||
|
||||
if (x != T{3}) return false;
|
||||
if (y != T{4}) return false;
|
||||
if (z != T{3}) return false;
|
||||
if (w != T{4}) return false;
|
||||
#endif
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
|
||||
@@ -18,8 +18,10 @@
|
||||
* */
|
||||
|
||||
/* HIT_START
|
||||
* BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t EXCLUDE_HIP_PLATFORM nvcc
|
||||
* BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_NVCC__ -I%hip-path/include -I/usr/local/cuda/include %S/%s -o %T/hipMalloc_nv -L/usr/local/cuda/lib64 -lcudart EXCLUDE_HIP_PLATFORM hcc
|
||||
* BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: hipMalloc_nv EXCLUDE_HIP_PLATFORM hcc
|
||||
* TEST: hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvcc
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
@@ -21,7 +21,7 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s test_common.cpp
|
||||
* BUILD: %t %s test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
@@ -34,41 +34,43 @@ OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWA
|
||||
|
||||
using namespace std;
|
||||
|
||||
const string directed_dir = "directed_tests" + string(PATH_SEPERATOR_STR) + "hipEnvVar";
|
||||
const string dir = "." + string(PATH_SEPERATOR_STR) + "hipEnvVar";
|
||||
|
||||
int getDeviceNumber() {
|
||||
FILE* in;
|
||||
char buff[512];
|
||||
string str;
|
||||
std::this_thread::sleep_for(std::chrono::milliseconds(10));
|
||||
if (!(in = popen("./directed_tests/hipEnvVar -c", "r"))) {
|
||||
// Check at same level
|
||||
if (!(in = popen("./hipEnvVar -c", "r"))) {
|
||||
FILE* in = popen((directed_dir + " -c").c_str(), "r");
|
||||
if(fgets(buff, 512, in) == NULL){
|
||||
pclose(in);
|
||||
//Check at same level
|
||||
in = popen((dir + " -c").c_str(), "r");
|
||||
if(fgets(buff, 512, in) == NULL){
|
||||
pclose(in);
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
while (fgets(buff, 512, in) != NULL) {
|
||||
cout << buff;
|
||||
}
|
||||
cout << buff;
|
||||
pclose(in);
|
||||
return atoi(buff);
|
||||
}
|
||||
|
||||
// Query the current device ID remotely to hipEnvVar
|
||||
void getDevicePCIBusNumRemote(int deviceID, char* pciBusID) {
|
||||
FILE* in;
|
||||
string str = "./directed_tests/hipEnvVar -d ";
|
||||
str += std::to_string(deviceID);
|
||||
void getDevicePCIBusNumRemote(int deviceID, char* pciBusID) {
|
||||
std::this_thread::sleep_for(std::chrono::milliseconds(10));
|
||||
if (!(in = popen(str.c_str(), "r"))) {
|
||||
// Check at same level
|
||||
if (!(in = popen("./hipEnvVar -d ", "r"))) {
|
||||
exit(1);
|
||||
FILE* in = popen((directed_dir + " -d " + std::to_string(deviceID)).c_str(), "r");
|
||||
if(fgets(pciBusID, 100, in) == NULL){
|
||||
pclose(in);
|
||||
//Check at same level
|
||||
in = popen((dir + " -d").c_str(), "r");
|
||||
if(fgets(pciBusID, 100, in) == NULL){
|
||||
pclose(in);
|
||||
return;
|
||||
}
|
||||
|
||||
}
|
||||
while (fgets(pciBusID, 100, in) != NULL) {
|
||||
cout << pciBusID;
|
||||
}
|
||||
cout << pciBusID;
|
||||
pclose(in);
|
||||
return;
|
||||
}
|
||||
|
||||
// Query the current device ID locally on AMD path
|
||||
@@ -81,9 +83,8 @@ void getDevicePCIBusNum(int deviceID, char* pciBusID) {
|
||||
}
|
||||
|
||||
int main() {
|
||||
unsetenv("HIP_VISIBLE_DEVICES");
|
||||
unsetenv("CUDA_VISIBLE_DEVICES");
|
||||
|
||||
unsetenv(HIP_VISIBLE_DEVICES_STR);
|
||||
unsetenv(CUDA_VISIBLE_DEVICES_STR);
|
||||
std::vector<std::string> devPCINum;
|
||||
char pciBusID[100];
|
||||
// collect the device pci bus ID for all devices
|
||||
@@ -126,8 +127,8 @@ int main() {
|
||||
setenv("CUDA_VISIBLE_DEVICES", "0,1,2", 1);
|
||||
assert(getDeviceNumber() == 3);
|
||||
// test if CUDA_VISIBLE_DEVICES will be accepted by the runtime
|
||||
unsetenv("HIP_VISIBLE_DEVICES");
|
||||
unsetenv("CUDA_VISIBLE_DEVICES");
|
||||
unsetenv(HIP_VISIBLE_DEVICES_STR);
|
||||
unsetenv(CUDA_VISIBLE_DEVICES_STR);
|
||||
setenv("CUDA_VISIBLE_DEVICES", "0,1,2", 1);
|
||||
assert(getDeviceNumber() == 3);
|
||||
}
|
||||
|
||||
@@ -22,7 +22,7 @@ THE SOFTWARE.
|
||||
// Test the HCC-specific API extensions for HIP:
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s
|
||||
* BUILD: %t %s EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t EXCLUDE_HIP_PLATFORM all
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
Обычный файл → Исполняемый файл
@@ -35,6 +35,7 @@ THE SOFTWARE.
|
||||
#include <memory>
|
||||
#include <iostream>
|
||||
#include <iterator>
|
||||
#include <vector>
|
||||
|
||||
static constexpr auto NUM_THREADS{128};
|
||||
static constexpr auto NUM_BLOCKS{32};
|
||||
|
||||
@@ -0,0 +1,106 @@
|
||||
/*
|
||||
Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s EXCLUDE_HIP_PLATFORM all
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <test_common.h>
|
||||
|
||||
extern "C" __device__ HIP_vector_base<long, 2>::Native_vec_ __ockl_call_host_function(
|
||||
ulong fptr, ulong arg0, ulong arg1, ulong arg2, ulong arg3, ulong arg4, ulong arg5, ulong arg6);
|
||||
|
||||
static void callee(ulong* output, ulong* input) {
|
||||
output[0] = input[0] + 1;
|
||||
output[1] = input[1] + input[2];
|
||||
}
|
||||
|
||||
__global__ void kernel(ulong fptr, ulong* retval0, ulong* retval1) {
|
||||
uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
|
||||
ulong arg0 = (ulong)fptr;
|
||||
ulong arg1 = tid;
|
||||
ulong arg2 = 42;
|
||||
ulong arg3 = tid % 23;
|
||||
ulong arg4 = 0;
|
||||
ulong arg5 = 0;
|
||||
ulong arg6 = 0;
|
||||
ulong arg7 = 0;
|
||||
|
||||
long2 result = {0, 0};
|
||||
if (tid % 71 != 1) {
|
||||
result.data = __ockl_call_host_function(arg0, arg1, arg2, arg3, arg4, arg5, arg6, arg7);
|
||||
retval0[tid] = result.x;
|
||||
retval1[tid] = result.y;
|
||||
}
|
||||
}
|
||||
|
||||
static bool test() {
|
||||
uint num_blocks = 5;
|
||||
uint threads_per_block = 1000;
|
||||
uint num_threads = num_blocks * threads_per_block;
|
||||
|
||||
void* retval0_void;
|
||||
HIPCHECK(hipHostMalloc(&retval0_void, 8 * num_threads));
|
||||
auto retval0 = reinterpret_cast<uint64_t*>(retval0_void);
|
||||
for (uint i = 0; i != num_threads; ++i) {
|
||||
retval0[i] = 0x23232323;
|
||||
}
|
||||
|
||||
void* retval1_void;
|
||||
HIPCHECK(hipHostMalloc(&retval1_void, 8 * num_threads));
|
||||
auto retval1 = reinterpret_cast<uint64_t*>(retval1_void);
|
||||
for (uint i = 0; i != num_threads; ++i) {
|
||||
retval1[i] = 0x23232323;
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(kernel, dim3(num_blocks), dim3(threads_per_block), 0, 0, (ulong)callee,
|
||||
retval0, retval1);
|
||||
|
||||
hipStreamSynchronize(0);
|
||||
|
||||
for (uint i = 0; i != num_threads; ++i) {
|
||||
ulong value = retval0[i];
|
||||
if (i % 71 == 1) {
|
||||
if (value != 0x23232323) {
|
||||
printf("failed\n");
|
||||
return false;
|
||||
}
|
||||
} else {
|
||||
if (value != i + 1) {
|
||||
printf("failed\n");
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
if (!test()) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
printf("passed\n");
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,125 @@
|
||||
/*
|
||||
Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s EXCLUDE_HIP_PLATFORM all
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <test_common.h>
|
||||
|
||||
// This is NOT a real printf test. It is a test for calling a host function
|
||||
// which happens to be a wrapper around system printf.
|
||||
|
||||
extern "C" __device__ HIP_vector_base<long, 2>::Native_vec_ __ockl_call_host_function(
|
||||
ulong fptr, ulong arg0, ulong arg1, ulong arg2, ulong arg3, ulong arg4, ulong arg5, ulong arg6);
|
||||
|
||||
// FuncCall service function that expects three arguments bundled in the
|
||||
// request: the format string, and two uint64_t arguments.
|
||||
void print_things_0(ulong* output, ulong* input) {
|
||||
auto fmt = reinterpret_cast<const char*>(input);
|
||||
auto arg0 = input[2];
|
||||
auto arg1 = input[3];
|
||||
output[0] = fprintf(stdout, fmt, arg0, arg1);
|
||||
}
|
||||
|
||||
__global__ void kernel0(ulong fptr, ulong* retval) {
|
||||
uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
|
||||
ulong arg0 = fptr;
|
||||
|
||||
const char* str = "(%lu -> %lu)\n";
|
||||
ulong arg1 = 0;
|
||||
for (int ii = 0; ii != 8; ++ii) {
|
||||
arg1 |= (ulong)str[ii] << (8 * ii);
|
||||
}
|
||||
ulong arg2 = 0;
|
||||
for (int ii = 0; ii != 7; ++ii) {
|
||||
arg2 |= (ulong)str[ii + 8] << (8 * ii);
|
||||
}
|
||||
|
||||
ulong arg3 = 42;
|
||||
ulong arg4 = tid;
|
||||
ulong arg5 = 0;
|
||||
ulong arg6 = 0;
|
||||
ulong arg7 = 0;
|
||||
|
||||
long2 result = {0, 0};
|
||||
result.data = __ockl_call_host_function(arg0, arg1, arg2, arg3, arg4, arg5, arg6, arg7);
|
||||
*retval = result.x;
|
||||
}
|
||||
|
||||
// FuncCall service function that expects two arguments bundled in the request:
|
||||
// a kernel "name" and a uint64_t thread ID. The format string is built into the
|
||||
// service function itself.
|
||||
void print_things_1(ulong* output, const ulong* input) {
|
||||
auto name = reinterpret_cast<const char*>(input[0]);
|
||||
auto tid = input[1];
|
||||
output[0] = fprintf(stdout, "kernel: %s; tid: %lu\n", name, tid);
|
||||
}
|
||||
|
||||
__global__ void kernel1(ulong fptr, ulong name, ulong* retval) {
|
||||
uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
|
||||
ulong arg0 = fptr;
|
||||
ulong arg1 = name;
|
||||
ulong arg2 = tid;
|
||||
ulong arg3 = 0;
|
||||
ulong arg4 = 0;
|
||||
ulong arg5 = 0;
|
||||
ulong arg6 = 0;
|
||||
ulong arg7 = 0;
|
||||
|
||||
long2 result = {0, 0};
|
||||
result.data = __ockl_call_host_function(arg0, arg1, arg2, arg3, arg4, arg5, arg6, arg7);
|
||||
*retval = result.x;
|
||||
}
|
||||
|
||||
static bool test() {
|
||||
void* retval_void;
|
||||
HIPCHECK(hipHostMalloc(&retval_void, 8));
|
||||
auto retval = reinterpret_cast<uint64_t*>(retval_void);
|
||||
*retval = 0x23232323;
|
||||
|
||||
hipLaunchKernelGGL(kernel0, dim3(1), dim3(1), 0, 0, (ulong)print_things_0, retval);
|
||||
hipStreamSynchronize(0);
|
||||
if (*retval != strlen("(42 -> 0)\n")) {
|
||||
return false;
|
||||
}
|
||||
|
||||
*retval = 0x23232323;
|
||||
const char* name = "kernel1";
|
||||
hipLaunchKernelGGL(kernel1, dim3(1), dim3(1), 0, 0, (ulong)print_things_1, (ulong)name, retval);
|
||||
hipStreamSynchronize(0);
|
||||
if (*retval != strlen("kernel: kernel1; tid: 0\n")) {
|
||||
return false;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
if (!test()) {
|
||||
printf("failed\n");
|
||||
return 1;
|
||||
}
|
||||
printf("passed\n");
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,62 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
// Test the Grid_Launch syntax.
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hip_ext.h"
|
||||
#include "test_common.h"
|
||||
|
||||
void test(size_t N) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N);
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipExtLaunchKernelGGL(HipTest::vectorADD, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, 0, nullptr, nullptr, 0,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
|
||||
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
test(N);
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -32,7 +32,6 @@ THE SOFTWARE.
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
#include <hip/hip_ext.h>
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN * sizeof(float)
|
||||
|
||||
@@ -1,145 +0,0 @@
|
||||
#include <stdio.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <unistd.h>
|
||||
#include "test_common.h"
|
||||
#include <atomic>
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
enum class ExecState
|
||||
{
|
||||
EXEC_NOT_STARTED,
|
||||
EXEC_STARTED,
|
||||
EXEC_CB_STARTED,
|
||||
EXEC_CB_FINISHED,
|
||||
EXEC_FINISHED
|
||||
};
|
||||
|
||||
struct UserData
|
||||
{
|
||||
size_t size;
|
||||
int* ptr;
|
||||
};
|
||||
|
||||
// Global variable to check exection order
|
||||
std::atomic<ExecState> gData(ExecState::EXEC_NOT_STARTED);
|
||||
|
||||
|
||||
void myCallback(hipStream_t stream, hipError_t status, void* user_data)
|
||||
{
|
||||
if(gData.load() != ExecState::EXEC_STARTED)
|
||||
return; // Error hence return early
|
||||
|
||||
gData.store(ExecState::EXEC_CB_STARTED);
|
||||
|
||||
UserData* data = reinterpret_cast<UserData*>(user_data);
|
||||
printf("Callback started\n");
|
||||
|
||||
sleep(1);
|
||||
|
||||
printf("Callback ending.\n");
|
||||
gData.store(ExecState::EXEC_CB_FINISHED);
|
||||
}
|
||||
|
||||
bool test(int count)
|
||||
{
|
||||
printf("\n============ Test iteration %d =============\n",count);
|
||||
// Stream
|
||||
hipStream_t stream;
|
||||
bool result = true;
|
||||
|
||||
gData.store(ExecState::EXEC_STARTED);
|
||||
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
// Array size
|
||||
size_t size = 10000;
|
||||
|
||||
// Device array
|
||||
int *data = NULL;
|
||||
HIPCHECK(hipMalloc((void**)&data, sizeof(int) * size));
|
||||
|
||||
// Initialize device array to -1
|
||||
HIPCHECK(hipMemset(data, -1, sizeof(int) * size));
|
||||
|
||||
// Host array
|
||||
int *host = NULL;
|
||||
HIPCHECK(hipHostMalloc((void**)&host, sizeof(int) * size));
|
||||
|
||||
// Print host ptr address
|
||||
printf("In main thread\n");
|
||||
|
||||
// Initialize user_data for callback
|
||||
UserData arg;
|
||||
arg.size = size;
|
||||
arg.ptr = host;
|
||||
|
||||
// Synchronize device
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
// Asynchronous copy from device to host
|
||||
HIPCHECK(hipMemcpyAsync(host, data, sizeof(int) * size, hipMemcpyDeviceToHost, stream));
|
||||
|
||||
// Asynchronous memset on device
|
||||
HIPCHECK(hipMemsetAsync(data, 0, sizeof(int) * size, stream));
|
||||
|
||||
// Add callback - should happen after hipMemsetAsync()
|
||||
HIPCHECK(hipStreamAddCallback(stream, myCallback, &arg, 0));
|
||||
|
||||
printf("Will wait in main thread until callback completes\n");
|
||||
|
||||
//This should synchronize the stream (including the callback)
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
|
||||
if(gData.load() != ExecState::EXEC_CB_FINISHED)
|
||||
{
|
||||
std::cout<<"Callback is not finished\n";
|
||||
return false;
|
||||
}
|
||||
printf("Callback completed will resume main thread execution\n");
|
||||
|
||||
if(host[size/2] != -1)
|
||||
{
|
||||
// Print some host data that just got copied
|
||||
printf("Pseudo host data printing (should be -1): %d\n", host[size/2]);
|
||||
result = false;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemcpy(host, data, sizeof(int)*size, hipMemcpyDeviceToHost));
|
||||
|
||||
if(host[size-1] != 0)
|
||||
{
|
||||
printf("Pseudo host data printing (should be 0): %d\n", host[size-1]);
|
||||
result = false;
|
||||
}
|
||||
|
||||
HIPCHECK(hipFree(data));
|
||||
HIPCHECK(hipHostFree(host));
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
gData.store(ExecState::EXEC_FINISHED);
|
||||
return result;
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
// Test involves multithreading hence running multiple times
|
||||
// to make sure consitency in the behavior
|
||||
bool status = true;
|
||||
|
||||
for(int i=0; i < 10; i++){
|
||||
status = test(i+1);
|
||||
if(status == false)
|
||||
{
|
||||
failed("Test Failed!\n");
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
if(status == true) passed();
|
||||
return 0;
|
||||
}
|
||||
@@ -1,133 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
#include "test_common.h"
|
||||
#include <mutex>
|
||||
#include <condition_variable>
|
||||
#include <stdlib.h>
|
||||
|
||||
|
||||
//Globals
|
||||
const int workloadCount = 1000000;
|
||||
std::mutex gMutx;
|
||||
bool callbackCompleted = false;
|
||||
std::condition_variable condVar;
|
||||
|
||||
// Device function
|
||||
__global__ void increment(int *data,int N)
|
||||
{
|
||||
int i = blockIdx.x*blockDim.x + threadIdx.x;
|
||||
|
||||
if(i < N)
|
||||
data[i] = 1 + data[i];
|
||||
}
|
||||
|
||||
struct USER_DATA
|
||||
{
|
||||
int* result; // Data received from device
|
||||
int* copyOfOriginalData; // Copy of initial data which will be used for validation
|
||||
};
|
||||
|
||||
// Callback
|
||||
void callback(hipStream_t event, hipError_t status, void *userData)
|
||||
{
|
||||
USER_DATA *data = (USER_DATA *)userData;
|
||||
|
||||
if(!(data == nullptr || data->result == nullptr || data->copyOfOriginalData == nullptr))
|
||||
{
|
||||
for(int i=0;i<workloadCount;i++)
|
||||
{
|
||||
if(data->result[i] != data->copyOfOriginalData[i]+1)
|
||||
{
|
||||
std::cout<<"Error value : "<<data->result[i]<<"| Expected value :"<<data->copyOfOriginalData[i]+1<<std::endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
callbackCompleted = true;
|
||||
condVar.notify_all();
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
int *hData = nullptr;
|
||||
int *dData = nullptr;
|
||||
int *hResultData = nullptr;
|
||||
int devCount = 0;
|
||||
int size = workloadCount * sizeof(int);
|
||||
|
||||
// query device count
|
||||
HIPCHECK(hipGetDeviceCount(&devCount));
|
||||
|
||||
// Allocate
|
||||
// Host allocation
|
||||
hData = (int *)malloc(size);
|
||||
hResultData = (int *)malloc(size);
|
||||
|
||||
if(hData == nullptr || hResultData == nullptr)
|
||||
{
|
||||
HIPCHECK(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
// Initialize host data
|
||||
for(int i =0; i<workloadCount; i++)
|
||||
{
|
||||
hData[i] = rand()%workloadCount;
|
||||
}
|
||||
|
||||
// Device allocation
|
||||
HIPCHECK(hipMalloc(&dData,size));
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(dData,hData,size,hipMemcpyHostToDevice,0));
|
||||
|
||||
dim3 block(256);
|
||||
dim3 grid((workloadCount + block.x-1) / block.x);
|
||||
hipLaunchKernelGGL(increment, grid, block, 0, 0, dData, workloadCount);
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(hResultData, dData, size, hipMemcpyDeviceToHost, 0));
|
||||
|
||||
hipStream_t stream;
|
||||
USER_DATA *inputParam = (USER_DATA*)malloc(sizeof(USER_DATA));
|
||||
|
||||
if(inputParam == nullptr) return 0;
|
||||
|
||||
inputParam->result = hResultData;
|
||||
inputParam->copyOfOriginalData = hData;
|
||||
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
HIPCHECK(hipStreamAddCallback(stream,callback,inputParam,0));
|
||||
|
||||
// Wait for stream add callback to complete
|
||||
std::unique_lock<std::mutex> l(gMutx);
|
||||
|
||||
while(!callbackCompleted)
|
||||
condVar.wait(l);
|
||||
|
||||
// Will destroy device memory hence explicite hipFree is not needed
|
||||
HIPCHECK(hipDeviceReset());
|
||||
|
||||
free(hData);
|
||||
free(hResultData);
|
||||
passed();
|
||||
}
|
||||
@@ -1,178 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <thread>
|
||||
#include <unistd.h>
|
||||
#include <mutex>
|
||||
#include <condition_variable>
|
||||
#include "test_common.h"
|
||||
|
||||
// Will indicate completion of callback added as part of hipStreamAddCallback
|
||||
struct signal
|
||||
{
|
||||
int completedThreads;
|
||||
std::mutex mu;
|
||||
std::condition_variable cv;
|
||||
};
|
||||
|
||||
struct workload
|
||||
{
|
||||
int _workloadId;
|
||||
int _deviceID;
|
||||
|
||||
int *copyOf_hData; // copy of host data which will be used to validation
|
||||
int *hData; // will contain host data
|
||||
int *dData; // device data will be stored
|
||||
hipStream_t _stream;// stream on which data will be processed
|
||||
|
||||
bool success; // start will be stored
|
||||
};
|
||||
|
||||
// Global data
|
||||
int numWorkloads = 8;
|
||||
const int perWorkloadSize = 1000000;
|
||||
signal completionSignal;
|
||||
|
||||
// Kernel run on device
|
||||
__global__ void increment(int *data, int N)
|
||||
{
|
||||
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
if (i < N)
|
||||
data[i] = data[i]+1;
|
||||
}
|
||||
|
||||
/*
|
||||
* Method validates processed data array with saved copy and notifies via conditional variable to all waiting threads
|
||||
*/
|
||||
void Analyze(hipStream_t event, hipError_t status, void *data)
|
||||
{
|
||||
HIPCHECK(status);
|
||||
|
||||
workload *W = (workload *) data;
|
||||
|
||||
if(W != NULL)
|
||||
{
|
||||
W->success = true;
|
||||
|
||||
for (int i=0; i< perWorkloadSize; ++i)
|
||||
{
|
||||
W->success &= (W->copyOf_hData[i] == (W->hData[i]+1));
|
||||
if(!W->success)
|
||||
{
|
||||
std::cout<<"\nExpected Data :"<<(W->hData[i]+1)<<" Current Data :"<<W->copyOf_hData[i]<<std::endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
std::lock_guard<std::mutex> guard(completionSignal.mu);
|
||||
completionSignal.completedThreads += 1;
|
||||
|
||||
completionSignal.cv.notify_all();
|
||||
}
|
||||
|
||||
/*
|
||||
* Thread routine to launch workloads into separate stream
|
||||
*/
|
||||
void LaunchWorkload(void *args)
|
||||
{
|
||||
workload *W = (workload *) args;
|
||||
|
||||
if (W == nullptr) return;
|
||||
|
||||
std::srand(std::time(nullptr));
|
||||
|
||||
HIPCHECK(hipSetDevice(W->_deviceID));
|
||||
|
||||
// Allocate memory
|
||||
HIPCHECK(hipMalloc(&W->dData,perWorkloadSize*sizeof(int)));
|
||||
|
||||
size_t s = perWorkloadSize*sizeof(int);
|
||||
HIPCHECK(hipMemset(W->dData,0,s))
|
||||
|
||||
W->hData = (int *) malloc(perWorkloadSize*sizeof(int));
|
||||
W->copyOf_hData = (int *) malloc(perWorkloadSize*sizeof(int));
|
||||
|
||||
// Initialize host array
|
||||
for(int i =0;i<perWorkloadSize;i++)
|
||||
{
|
||||
W->hData[i] = W->copyOf_hData[i] = std::rand() % perWorkloadSize;
|
||||
}
|
||||
|
||||
HIPCHECK(hipStreamCreate(&W->_stream));
|
||||
|
||||
dim3 block(256);
|
||||
dim3 grid((perWorkloadSize + block.x-1) / block.x);
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(W->dData,W->hData,perWorkloadSize*sizeof(int),hipMemcpyHostToDevice,W->_stream));
|
||||
|
||||
hipLaunchKernelGGL(increment, grid, block,0, W->_stream,W->dData, perWorkloadSize);
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(W->copyOf_hData,W->dData,perWorkloadSize*sizeof(int),hipMemcpyDeviceToHost,W->_stream));
|
||||
|
||||
HIPCHECK(hipStreamAddCallback(W->_stream, Analyze, W,0))
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
|
||||
int numDevice = 0;
|
||||
|
||||
HIPCHECK(hipGetDeviceCount(&numDevice));
|
||||
|
||||
std::thread workerThread[numWorkloads];
|
||||
|
||||
workload *workloads;
|
||||
workloads = (workload *) malloc(numWorkloads * sizeof(workload));
|
||||
|
||||
for(int i =0;i<numWorkloads; i++)
|
||||
{
|
||||
workloads[i]._workloadId = i;
|
||||
workloads[i]._deviceID = i%numDevice;
|
||||
|
||||
// launch threads
|
||||
workerThread[i] = std::thread(LaunchWorkload,&workloads[i]);
|
||||
}
|
||||
|
||||
// should wait for joining all threads
|
||||
for(int i =0; i<numWorkloads; i++)
|
||||
{
|
||||
workerThread[i].join();
|
||||
}
|
||||
|
||||
// wait util all callbacks are done
|
||||
std::unique_lock<std::mutex> l(completionSignal.mu);
|
||||
while(completionSignal.completedThreads != numWorkloads)
|
||||
{
|
||||
completionSignal.cv.wait(l);
|
||||
}
|
||||
|
||||
//clean-up
|
||||
hipDeviceReset();
|
||||
free(workloads);
|
||||
|
||||
passed();
|
||||
return 0;
|
||||
}
|
||||
@@ -53,7 +53,7 @@ int runTest(int argc, char** argv) {
|
||||
|
||||
hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice);
|
||||
|
||||
struct hipResourceDesc resDesc;
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
@@ -62,7 +62,7 @@ int runTest(int argc, char** argv) {
|
||||
hipCreateSurfaceObject(&surfaceObject, &resDesc);
|
||||
|
||||
hipMallocArray(&hipOutArray, &channelDesc, width, height);
|
||||
struct hipResourceDesc resOutDesc;
|
||||
hipResourceDesc resOutDesc;
|
||||
memset(&resOutDesc, 0, sizeof(resOutDesc));
|
||||
resOutDesc.resType = hipResourceTypeArray;
|
||||
resOutDesc.res.array.array = hipOutArray;
|
||||
|
||||
@@ -33,7 +33,15 @@ unsigned threadsPerBlock = 256;
|
||||
int p_gpuDevice = 0;
|
||||
unsigned p_verbose = 0;
|
||||
int p_tests = -1; /*which tests to run. Interpretation is left to each test. default:all*/
|
||||
|
||||
#ifdef _WIN64
|
||||
const char* HIP_VISIBLE_DEVICES_STR = "HIP_VISIBLE_DEVICES=";
|
||||
const char* CUDA_VISIBLE_DEVICES_STR = "CUDA_VISIBLE_DEVICES=";
|
||||
const char* PATH_SEPERATOR_STR = "\\";
|
||||
#else
|
||||
const char* HIP_VISIBLE_DEVICES_STR = "HIP_VISIBLE_DEVICES";
|
||||
const char* CUDA_VISIBLE_DEVICES_STR = "CUDA_VISIBLE_DEVICES";
|
||||
const char* PATH_SEPERATOR_STR = "/";
|
||||
#endif
|
||||
|
||||
namespace HipTest {
|
||||
|
||||
|
||||
@@ -104,6 +104,7 @@ THE SOFTWARE.
|
||||
#define popen(x,y) _popen(x,y)
|
||||
#define pclose(x) _pclose(x)
|
||||
#define setenv(x,y,z) _putenv_s(x,y)
|
||||
#define unsetenv _putenv
|
||||
#else
|
||||
#define aligned_free(x) free(x)
|
||||
#endif
|
||||
@@ -120,6 +121,9 @@ extern unsigned threadsPerBlock;
|
||||
extern int p_gpuDevice;
|
||||
extern unsigned p_verbose;
|
||||
extern int p_tests;
|
||||
extern const char* HIP_VISIBLE_DEVICES_STR;
|
||||
extern const char* CUDA_VISIBLE_DEVICES_STR;
|
||||
extern const char* PATH_SEPERATOR_STR;
|
||||
|
||||
// ********************* CPP section *********************
|
||||
#ifdef __cplusplus
|
||||
|
||||
@@ -53,13 +53,13 @@ int runTest(int argc, char** argv) {
|
||||
|
||||
hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice);
|
||||
|
||||
struct hipResourceDesc resDesc;
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
|
||||
// Specify texture object parameters
|
||||
struct hipTextureDesc texDesc;
|
||||
hipTextureDesc texDesc;
|
||||
memset(&texDesc, 0, sizeof(texDesc));
|
||||
texDesc.addressMode[0] = hipAddressModeWrap;
|
||||
texDesc.addressMode[1] = hipAddressModeWrap;
|
||||
|
||||
Ссылка в новой задаче
Block a user