diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 7b286d638d..b804d8d62d 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -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; diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 21401b0626..c4fbd1db6d 100644 --- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -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`* | diff --git a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index cbe4afb009..4fa033e21c 100644 --- a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -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`* | diff --git a/projects/hip/hipify-clang/src/ArgParse.cpp b/projects/hip/hipify-clang/src/ArgParse.cpp index 4f648c996f..cbf2f941e8 100644 --- a/projects/hip/hipify-clang/src/ArgParse.cpp +++ b/projects/hip/hipify-clang/src/ArgParse.cpp @@ -30,7 +30,7 @@ cl::opt OutputFilename("o", cl::cat(ToolTemplateCategory)); cl::opt OutputDir("o-dir", - cl::desc("Output direcory"), + cl::desc("Output directory"), cl::value_desc("directory"), cl::cat(ToolTemplateCategory)); @@ -45,17 +45,17 @@ cl::opt GeneratePython("python", cl::cat(ToolTemplateCategory)); cl::opt 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 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 TemporaryDir("temp-dir", - cl::desc("Temporary direcory"), + cl::desc("Temporary directory"), cl::value_desc("directory"), cl::cat(ToolTemplateCategory)); diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp index f88086f2cb..4603af69d0 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -1105,161 +1105,195 @@ const std::map 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 diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp index 7b34d97ab4..d8acb9882c 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp @@ -595,42 +595,51 @@ const std::map 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 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 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 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 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 diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index a6f9cc0826..68e3277270 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/include/hip/hcc_detail/device_functions.h @@ -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; diff --git a/projects/hip/include/hip/hcc_detail/device_library_decls.h b/projects/hip/include/hip/hcc_detail/device_library_decls.h index 8cfb020caa..ac35823cd2 100644 --- a/projects/hip/include/hip/hcc_detail/device_library_decls.h +++ b/projects/hip/include/hip/hcc_detail/device_library_decls.h @@ -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); diff --git a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp index c493eec933..c9953c27d4 100644 --- a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp @@ -176,5 +176,4 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, numBlocks, dimBlocks, sharedMemBytes, stream, &config[0]); } - #pragma GCC visibility pop diff --git a/projects/hip/include/hip/hcc_detail/helpers.hpp b/projects/hip/include/hip/hcc_detail/helpers.hpp index cdae0e0fc4..b94b126994 100644 --- a/projects/hip/include/hip/hcc_detail/helpers.hpp +++ b/projects/hip/include/hip/hcc_detail/helpers.hpp @@ -27,6 +27,26 @@ THE SOFTWARE. // std::false_type, std result_of and std::true_type. #include // For std::declval. +#ifdef __has_include // Check if __has_include is present +# if __has_include() // Check for version header +# include +# 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 using void_t_ = void; -#if (__cplusplus < 201402L) -template -struct is_callable_impl : is_callable_impl {}; +#if HIP_HAS_INVOCABLE +template +struct is_callable_impl; -// Pointer to member function, call through non-pointer. -template -struct is_callable_impl< - F(C, Ts...), 0u, - void_t_().*std::declval())(std::declval()...))> > - : std::true_type {}; - -// Pointer to member function, call through pointer. -template -struct is_callable_impl< - F(C, Ts...), 1u, - void_t_()).*std::declval())(std::declval()...))> > - : std::true_type {}; - -// Pointer to member data, call through non-pointer, no args. -template -struct is_callable_impl().*std::declval())> > - : std::true_type {}; - -// Pointer to member data, call through pointer, no args. -template -struct is_callable_impl().*std::declval())> > - : std::true_type {}; - -// General call, n args. template -struct is_callable_impl()(std::declval()...))> > - : std::true_type {}; - -// Not callable. -template -struct is_callable_impl : std::false_type {}; -#elif (__cplusplus < 201703L) +struct is_callable_impl : std::is_invocable {}; +#elif HIP_HAS_RESULT_OF_SFINAE template struct is_callable_impl : std::false_type {}; template struct is_callable_impl::type > > : std::true_type {}; #else +template +auto simple_invoke(T Base::*pmd, Derived&& ref) +-> decltype(static_cast(ref).*pmd); + +template +auto simple_invoke(PMD&& pmd, Pointer&& ptr) +-> decltype((*static_cast(ptr)).*static_cast(pmd)); -// C++17 +template +auto simple_invoke(T Base::*pmd, const std::reference_wrapper& ref) +-> decltype(ref.get().*pmd); + +template +auto simple_invoke(T Base::*pmf, Derived&& ref, Args&&... args) +-> decltype((static_cast(ref).*pmf)(static_cast(args)...)); + +template +auto simple_invoke(PMF&& pmf, Pointer&& ptr, Args&&... args) +-> decltype(((*static_cast(ptr)).*static_cast(pmf))(static_cast(args)...)); + +template +auto simple_invoke(T Base::*pmf, const std::reference_wrapper& ref, Args&&... args) +-> decltype((ref.get().*pmf)(static_cast(args)...)); + +template +auto simple_invoke(F&& f, Ts&&... xs) +-> decltype(f(static_cast(xs)...)); template struct is_callable_impl : std::false_type {}; template -struct is_callable_impl > > : std::true_type {}; +struct is_callable_impl(), std::declval()...))> > + : std::true_type {}; + #endif + template struct is_callable : is_callable_impl {}; diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 4acee9e235..11648ce123 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -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; } diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index ba7f2f4ade..48cec81da4 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -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 diff --git a/projects/hip/include/hip/hcc_detail/hip_vector_types.h b/projects/hip/include/hip/hcc_detail/hip_vector_types.h index 3a10a06372..582a359fbd 100644 --- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/hip/include/hip/hcc_detail/hip_vector_types.h @@ -40,6 +40,7 @@ THE SOFTWARE. #endif #if defined(__cplusplus) && defined(__clang__) + #include #include 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{}>::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{}>::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{}>::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{}>::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{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::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{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::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{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::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{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::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{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::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{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::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 x; }; + + using value_type = T; }; template @@ -209,6 +264,8 @@ THE SOFTWARE. hip_impl::Scalar_accessor x; hip_impl::Scalar_accessor y; }; + + using value_type = T; }; template @@ -367,6 +424,8 @@ THE SOFTWARE. T z; }; }; + + using value_type = T; }; template @@ -380,6 +439,8 @@ THE SOFTWARE. hip_impl::Scalar_accessor z; hip_impl::Scalar_accessor w; }; + + using value_type = T; }; template diff --git a/projects/hip/include/hip/hip_ext.h b/projects/hip/include/hip/hip_ext.h index 3df911ee7a..a618462d6e 100644 --- a/projects/hip/include/hip/hip_ext.h +++ b/projects/hip/include/hip/hip_ext.h @@ -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 +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{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(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 diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index d2f81c8b7b..f3ce66c94e 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.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 -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 -inline static hipError_t hipBindTexture(size_t* offset, const struct texture& tex, - const void* devPtr, size_t size = UINT_MAX) { - return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size)); -} - -template -inline static hipError_t hipBindTexture(size_t* offset, struct texture& 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 -inline static hipError_t hipUnbindTexture(struct texture* tex) { - return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); -} - -template -inline static hipError_t hipUnbindTexture(struct texture &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 -inline static hipError_t hipBindTextureToArray(struct texture& tex, - hipArray_const_t array, - const struct hipChannelFormatDesc& desc) { - return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); -} - -template -inline static hipError_t hipBindTextureToArray(struct texture *tex, - hipArray_const_t array, - const struct hipChannelFormatDesc* desc) { - return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); -} - -template -inline static hipError_t hipBindTextureToArray(struct texture& tex, - hipArray_const_t array) { - return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array)); -} - -template -inline static hipChannelFormatDesc hipCreateChannelDesc() { - return cudaCreateChannelDesc(); -} - 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 +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 +inline static hipError_t hipBindTexture(size_t* offset, const struct texture& tex, + const void* devPtr, size_t size = UINT_MAX) { + return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size)); +} + +template +inline static hipError_t hipBindTexture(size_t* offset, struct texture& tex, + const void* devPtr, const hipChannelFormatDesc& desc, + size_t size = UINT_MAX) { + return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size)); +} + +template +inline static hipError_t hipUnbindTexture(struct texture* tex) { + return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); +} + +template +inline static hipError_t hipUnbindTexture(struct texture &tex) { + return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); +} + +template +inline static hipError_t hipBindTextureToArray(struct texture& tex, + hipArray_const_t array, + const hipChannelFormatDesc& desc) { + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); +} + +template +inline static hipError_t hipBindTextureToArray(struct texture *tex, + hipArray_const_t array, + const hipChannelFormatDesc* desc) { + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); +} + +template +inline static hipError_t hipBindTextureToArray(struct texture& tex, + hipArray_const_t array) { + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array)); +} + +template +inline static hipChannelFormatDesc hipCreateChannelDesc() { + return cudaCreateChannelDesc(); +} #endif //__CUDACC__ #endif // HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H diff --git a/projects/hip/src/hip_device.cpp b/projects/hip/src/hip_device.cpp index 6406b48c2d..403194483a 100644 --- a/projects/hip/src/hip_device.cpp +++ b/projects/hip/src/hip_device.cpp @@ -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); diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp index dcab059774..3d1a366463 100644 --- a/projects/hip/src/hip_event.cpp +++ b/projects/hip/src/hip_event.cpp @@ -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) { diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 377ce0aad4..e93490a08e 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -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 (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; } //--- diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index 6337edd370..8ee47eba4a 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -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) { diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 266f9b51d6..00a7af59f8 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -27,6 +27,7 @@ THE SOFTWARE. #include "hip_hcc_internal.h" #include "trace_helper.h" +#include #include __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 + 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(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 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(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(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 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(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(di.agentBaseAddress) + + (static_cast(dst) - + static_cast(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 lck{ + dst, [](void* p) { hsa_amd_memory_unlock(p); }}; + + throwing_result_check(hsa_amd_memory_lock(dst, n, &si.agentOwner, 1, + const_cast(&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(src))}; + + if (si.type == HSA_EXT_POINTER_TYPE_LOCKED) { + src = static_cast(si.agentBaseAddress) + + (static_cast(src) - + static_cast(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 lck{ + const_cast(src), [](void* p) { hsa_amd_memory_unlock(p); }}; + + throwing_result_check(hsa_amd_memory_lock(const_cast(src), n, + &di.agentOwner, 1, + const_cast(&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 lck0{ + nullptr, [](void* p) { hsa_amd_memory_unlock(p); }}; + std::unique_ptr 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(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(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(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, diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index bf3289bb49..a9f30197d5 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -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 (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); } diff --git a/projects/hip/tests/src/cppstd/hipInvocable11.cpp b/projects/hip/tests/src/cppstd/hipInvocable11.cpp new file mode 100644 index 0000000000..6fa7e53622 --- /dev/null +++ b/projects/hip/tests/src/cppstd/hipInvocable11.cpp @@ -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 + +int main() { passed(); } diff --git a/projects/hip/tests/src/cppstd/hipInvocable11Fallback.cpp b/projects/hip/tests/src/cppstd/hipInvocable11Fallback.cpp new file mode 100644 index 0000000000..071dce8997 --- /dev/null +++ b/projects/hip/tests/src/cppstd/hipInvocable11Fallback.cpp @@ -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 + +int main() { passed(); } diff --git a/projects/hip/tests/src/cppstd/hipInvocable14.cpp b/projects/hip/tests/src/cppstd/hipInvocable14.cpp new file mode 100644 index 0000000000..41b58bd7dd --- /dev/null +++ b/projects/hip/tests/src/cppstd/hipInvocable14.cpp @@ -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 + +int main() { passed(); } diff --git a/projects/hip/tests/src/cppstd/hipInvocable17.cpp b/projects/hip/tests/src/cppstd/hipInvocable17.cpp new file mode 100644 index 0000000000..aa37aac5fb --- /dev/null +++ b/projects/hip/tests/src/cppstd/hipInvocable17.cpp @@ -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 + +int main() { passed(); } diff --git a/projects/hip/tests/src/cppstd/is_callable_test.hpp b/projects/hip/tests/src/cppstd/is_callable_test.hpp new file mode 100644 index 0000000000..87e4e35b8b --- /dev/null +++ b/projects/hip/tests/src/cppstd/is_callable_test.hpp @@ -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 +#include + +using hip_impl::is_callable; + +template +struct callable_rank : callable_rank +{}; + +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::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); + + typedef void (*is_callable_function_pointer)(int); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); + + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable, callable_test_param)>::value, "callable failed"); + static_assert(not is_callable, callable_test_param)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); +}; + +struct test2 +{ + typedef int(callable_rank<0>::*fn)(int); + + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(not is_callable&)>::value, "Failed"); + static_assert(not is_callable 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> ST; + + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(not is_callable::value, "Failed"); + +}; + +struct test4 +{ + typedef int(callable_rank<0>::*fn); + + static_assert(not is_callable::value, "Failed"); +}; + +struct test5 +{ + typedef int(callable_rank<0>::*fn); + + static_assert(is_callable&)>::value, "Failed"); + static_assert(is_callable&&)>::value, "Failed"); + static_assert(is_callable&)>::value, "Failed"); + static_assert(is_callable&)>::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> ST; + + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + +}; + +struct test7 +{ + typedef void(*fp)(callable_rank<0>&, int); + + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(not is_callable&, int)>::value, "Failed"); + static_assert(not is_callable::value, "Failed"); + static_assert(not is_callable&)>::value, "Failed"); +}; + +struct test8 +{ + typedef void(&fp)(callable_rank<0>&, int); + + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(not is_callable&, int)>::value, "Failed"); + static_assert(not is_callable::value, "Failed"); + static_assert(not is_callable&)>::value, "Failed"); +}; diff --git a/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp b/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp index 14479881ff..70c8320073 100644 --- a/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp +++ b/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp @@ -34,6 +34,7 @@ THE SOFTWARE. #include #include #include +#include #include 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; } diff --git a/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp b/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp index ba23931ee5..4bf5d2c87d 100644 --- a/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp +++ b/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp @@ -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; } diff --git a/projects/hip/tests/src/gcc/hipMalloc.c b/projects/hip/tests/src/gcc/hipMalloc.c index e3c7a66919..19e90d5222 100644 --- a/projects/hip/tests/src/gcc/hipMalloc.c +++ b/projects/hip/tests/src/gcc/hipMalloc.c @@ -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 */ diff --git a/projects/hip/tests/src/hipEnvVar.cpp b/projects/hip/tests/src/hipEnvVar.cpp index a0b9503b41..7cf38fc77e 100644 --- a/projects/hip/tests/src/hipEnvVar.cpp +++ b/projects/hip/tests/src/hipEnvVar.cpp @@ -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 */ diff --git a/projects/hip/tests/src/hipEnvVarDriver.cpp b/projects/hip/tests/src/hipEnvVarDriver.cpp index 4817d157a4..07379f0878 100644 --- a/projects/hip/tests/src/hipEnvVarDriver.cpp +++ b/projects/hip/tests/src/hipEnvVarDriver.cpp @@ -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 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); } diff --git a/projects/hip/tests/src/hipHcc.cpp b/projects/hip/tests/src/hipHcc.cpp index 2d4523ecd6..9bcc999b34 100644 --- a/projects/hip/tests/src/hipHcc.cpp +++ b/projects/hip/tests/src/hipHcc.cpp @@ -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 */ diff --git a/projects/hip/tests/src/hiprtc/saxpy.cpp b/projects/hip/tests/src/hiprtc/saxpy.cpp old mode 100644 new mode 100755 index 191e48b846..156a44afe4 --- a/projects/hip/tests/src/hiprtc/saxpy.cpp +++ b/projects/hip/tests/src/hiprtc/saxpy.cpp @@ -35,6 +35,7 @@ THE SOFTWARE. #include #include #include +#include static constexpr auto NUM_THREADS{128}; static constexpr auto NUM_BLOCKS{32}; diff --git a/projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp b/projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp new file mode 100644 index 0000000000..4b73352bb7 --- /dev/null +++ b/projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp @@ -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 + +extern "C" __device__ HIP_vector_base::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(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(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; +} diff --git a/projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp b/projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp new file mode 100644 index 0000000000..742798c0b9 --- /dev/null +++ b/projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp @@ -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 + +// 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::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(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(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(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; +} diff --git a/projects/hip/tests/src/kernel/hipExtLaunchKernelGGL.cpp b/projects/hip/tests/src/kernel/hipExtLaunchKernelGGL.cpp new file mode 100644 index 0000000000..4164a87f1e --- /dev/null +++ b/projects/hip/tests/src/kernel/hipExtLaunchKernelGGL.cpp @@ -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(A_d), static_cast(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(); +} diff --git a/projects/hip/tests/src/runtimeApi/module/hipLaunchMultiKernelMultiDevice.cpp b/projects/hip/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp similarity index 100% rename from projects/hip/tests/src/runtimeApi/module/hipLaunchMultiKernelMultiDevice.cpp rename to projects/hip/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp index f98d5e4ec4..5896794e90 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp @@ -32,7 +32,6 @@ THE SOFTWARE. #include #include #include -#include #define LEN 64 #define SIZE LEN * sizeof(float) diff --git a/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp b/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp deleted file mode 100644 index e6492c7ce2..0000000000 --- a/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp +++ /dev/null @@ -1,145 +0,0 @@ -#include -#include -#include -#include "test_common.h" -#include - -/* 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 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(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; -} diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp deleted file mode 100644 index b2fb168fbc..0000000000 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp +++ /dev/null @@ -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 -#include -#include - - -//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;iresult[i] != data->copyOfOriginalData[i]+1) - { - std::cout<<"Error value : "<result[i]<<"| Expected value :"<copyOfOriginalData[i]+1<result = hResultData; - inputParam->copyOfOriginalData = hData; - - HIPCHECK(hipStreamCreate(&stream)); - HIPCHECK(hipStreamAddCallback(stream,callback,inputParam,0)); - - // Wait for stream add callback to complete - std::unique_lock l(gMutx); - - while(!callbackCompleted) - condVar.wait(l); - - // Will destroy device memory hence explicite hipFree is not needed - HIPCHECK(hipDeviceReset()); - - free(hData); - free(hResultData); - passed(); -} diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp deleted file mode 100644 index 3babf5869e..0000000000 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp +++ /dev/null @@ -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 -#include -#include -#include -#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 :"<copyOf_hData[i]< 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;ihData[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 l(completionSignal.mu); - while(completionSignal.completedThreads != numWorkloads) - { - completionSignal.cv.wait(l); - } - - //clean-up - hipDeviceReset(); - free(workloads); - - passed(); - return 0; -} diff --git a/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp b/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp index b6c016ecc0..4580220d1d 100644 --- a/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp +++ b/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp @@ -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; diff --git a/projects/hip/tests/src/test_common.cpp b/projects/hip/tests/src/test_common.cpp index 0205310351..e7a2622662 100644 --- a/projects/hip/tests/src/test_common.cpp +++ b/projects/hip/tests/src/test_common.cpp @@ -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 { diff --git a/projects/hip/tests/src/test_common.h b/projects/hip/tests/src/test_common.h index 92528e38e3..426ea846b1 100644 --- a/projects/hip/tests/src/test_common.h +++ b/projects/hip/tests/src/test_common.h @@ -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 diff --git a/projects/hip/tests/src/texture/hipTextureObj2D.cpp b/projects/hip/tests/src/texture/hipTextureObj2D.cpp index 148c966e80..649c748fd2 100644 --- a/projects/hip/tests/src/texture/hipTextureObj2D.cpp +++ b/projects/hip/tests/src/texture/hipTextureObj2D.cpp @@ -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;