From 78a36b4fb81d9eb2fda5861d05ccf02a134e400a Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Thu, 30 Jan 2020 15:05:53 -0500 Subject: [PATCH 1/5] [HIP][HIPIFY] Add some missing flags for cooperative launch and occupancy APIs [ROCm/hip commit: 6e867eacb6a8d970993b27c99853e47b3841da68] --- projects/hip/bin/hipify-perl | 5 +++++ .../CUDA_Driver_API_functions_supported_by_HIP.md | 5 +++-- .../CUDA_Runtime_API_functions_supported_by_HIP.md | 4 ++-- .../hipify-clang/src/CUDA2HIP_Driver_API_types.cpp | 2 ++ .../hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp | 2 +- .../hip/include/hip/hcc_detail/hip_runtime_api.h | 11 ++++++++--- projects/hip/include/hip/hip_runtime_api.h | 4 ++++ .../hip/include/hip/nvcc_detail/hip_runtime_api.h | 12 ++++++++++++ projects/hip/src/hip_hcc.cpp | 2 ++ 9 files changed, 39 insertions(+), 8 deletions(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 4fe9b80349..d18e227c3d 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -1186,6 +1186,7 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_PTX\b/hipErrorInvalidKernelFile/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_SOURCE\b/hipErrorInvalidSource/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_VALUE\b/hipErrorInvalidValue/g; + $ft{'numeric_literal'} += s/\bCUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE\b/hipErrorCooperativeLaunchTooLarge/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_LAUNCH_FAILED\b/hipErrorLaunchFailure/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_LAUNCH_OUT_OF_RESOURCES\b/hipErrorLaunchOutOfResources/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_LAUNCH_TIMEOUT\b/hipErrorLaunchTimeOut/g; @@ -1659,6 +1660,7 @@ sub simpleSubstitutions { $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/\bcudaErrorCooperativeLaunchTooLarge\b/hipErrorCooperativeLaunchTooLarge/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; @@ -1801,6 +1803,9 @@ sub simpleSubstitutions { $ft{'define'} += s/\bcudaTextureType3D\b/hipTextureType3D/g; $ft{'define'} += s/\bcudaTextureTypeCubemap\b/hipTextureTypeCubemap/g; $ft{'define'} += s/\bcudaTextureTypeCubemapLayered\b/hipTextureTypeCubemapLayered/g; + $ft{'define'} += s/\bcudaOccupancyDefault\b/hipOccupancyDefault/g; + $ft{'define'} += s/\bcudaCooperativeLaunchMultiDeviceNoPreSync\b/hipCooperativeLaunchMultiDeviceNoPreSync/g; + $ft{'define'} += s/\bcudaCooperativeLaunchMultiDeviceNoPostSync\b/hipCooperativeLaunchMultiDeviceNoPostSync/g; } # CUDA extern __shared__ syntax replace with HIP_DYNAMIC_SHARED() macro 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 ff14cf9bc1..5be6a4da8d 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 @@ -447,6 +447,7 @@ | 717 |*`CUDA_ERROR_INVALID_ADDRESS_SPACE`* | | | 718 |*`CUDA_ERROR_INVALID_PC`* | | | 719 |*`CUDA_ERROR_LAUNCH_FAILED`* |*`hipErrorLaunchFailure`* | +| 720 |*`CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE`* |*`hipErrorCooperativeLaunchTooLarge`* | | 800 |*`CUDA_ERROR_NOT_PERMITTED`* | | | 801 |*`CUDA_ERROR_NOT_SUPPORTED`* |*`hipErrorNotSupported`* | | 802 |*`CUDA_ERROR_SYSTEM_NOT_READY`* | | 10.0 | @@ -1063,8 +1064,8 @@ | `cuFuncSetSharedMemConfig` | | | `cuLaunchKernel` | `hipModuleLaunchKernel` | | `cuLaunchHostFunc` | | 10.0 | -| `cuLaunchCooperativeKernel` | | 9.0 | -| `cuLaunchCooperativeKernelMultiDevice` | | 9.0 | +| `cuLaunchCooperativeKernel` | `hipLaunchCooperativeKernel` | 9.0 | +| `cuLaunchCooperativeKernelMultiDevice` | `hipLaunchCooperativeKernelMultiDevice` | 9.0 | ## **19. Execution Control [DEPRECATED]** 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 201d2aad0f..54e0c89e06 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 @@ -800,7 +800,7 @@ | 717 |*`cudaErrorInvalidAddressSpace`* | | | | 718 |*`cudaErrorInvalidPc`* | | | | 719 |*`cudaErrorLaunchFailure`* | |*`hipErrorLaunchFailure`* | -| 720 |*`cudaErrorCooperativeLaunchTooLarge`* | 9.0 | | +| 720 |*`cudaErrorCooperativeLaunchTooLarge`* | 9.0 |*`hipErrorCooperativeLaunchTooLarge`* | | 800 |*`cudaErrorNotPermitted`* | | | | 801 |*`cudaErrorNotSupported`* | |*`hipErrorNotSupported`* | | 802 |*`cudaErrorSystemNotReady`* | 10.0 | | @@ -1045,7 +1045,7 @@ | define |`cudaMemAttachGlobal` | |`hipMemAttachGlobal` | | define |`cudaMemAttachHost` | |`hipMemAttachHost` | | define |`cudaMemAttachSingle` | | | -| define |`cudaOccupancyDefault` | | | +| define |`cudaOccupancyDefault` | |`hipOccupancyDefault` | | define |`cudaOccupancyDisableCachingOverride` | | | | define |`cudaPeerAccessDefault` | | | | define |`cudaStreamDefault` | |`hipStreamDefault` | 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 1c3c2634f7..a4d419384d 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -1284,6 +1284,8 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CUDA_ERROR_INVALID_PC", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 718 // cudaErrorLaunchFailure {"CUDA_ERROR_LAUNCH_FAILED", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 719 + // cudaErrorCooperativeLaunchTooLarge + {"CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 720 // cudaErrorNotPermitted {"CUDA_ERROR_NOT_PERMITTED", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 800 // cudaErrorNotSupported 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 5993f5d770..ade8b105f1 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp @@ -787,7 +787,7 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { {"cudaErrorInvalidPc", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 718 // CUDA_ERROR_LAUNCH_FAILED {"cudaErrorLaunchFailure", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 719 - // no analogue + // CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE {"cudaErrorCooperativeLaunchTooLarge", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 720 // CUDA_ERROR_NOT_PERMITTED {"cudaErrorNotPermitted", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 800 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 659a6c3c3a..8ee44be6ac 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -212,6 +212,11 @@ enum hipLimit_t { #define hipArrayCubemap 0x04 #define hipArrayTextureGather 0x08 +#define hipOccupancyDefault 0x00 + +#define hipCooperativeLaunchMultiDeviceNoPreSync 0x01 +#define hipCooperativeLaunchMultiDeviceNoPostSync 0x02 + /* * @brief hipJitOption * @enum @@ -2903,7 +2908,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigne * @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th * default stream is used with associated synchronization rules. * - * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue + * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue, hipErrorCooperativeLaunchTooLarge */ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDimX, void** kernelParams, unsigned int sharedMemBytes, @@ -2917,7 +2922,7 @@ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDim * @param [in] numDevices Size of the launchParamsList array. * @param [in] flags Flags to control launch behavior. * - * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue + * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue, hipErrorCooperativeLaunchTooLarge */ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags); @@ -2960,7 +2965,7 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( * @param [in] flags Extra flags for occupancy calculation (currently ignored) */ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( - int* numBlocks, const void* f, int blockSize, size_t dynSharedMemPerBlk, unsigned int flags); + int* numBlocks, const void* f, int blockSize, size_t dynSharedMemPerBlk, unsigned int flags __dparm(hipOccupancyDefault)); #if __HIP_VDI__ && !defined(__HCC__) /** diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index 64b2a85d8a..025688e98c 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -255,6 +255,10 @@ typedef enum __HIP_NODISCARD hipError_t { 713, ///< Produced when trying to unlock a non-page-locked memory. hipErrorLaunchFailure = 719, ///< An exception occurred on the device while executing a kernel. + hipErrorCooperativeLaunchTooLarge = + 720, ///< This error indicates that the number of blocks launched per grid for a kernel + ///< that was launched via cooperative launch APIs exceeds the maximum number of + ///< allowed blocks for the current device hipErrorNotSupported = 801, ///< Produced when the hip API is not supported/implemented hipErrorUnknown = 999, //< Unknown error. // HSA Runtime Error Codes start here. 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 05a162478a..1ee84b5057 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -140,6 +140,12 @@ typedef enum cudaChannelFormatKind hipChannelFormatKind; #define hipLimitMallocHeapSize cudaLimitMallocHeapSize #define hipIpcMemLazyEnablePeerAccess cudaIpcMemLazyEnablePeerAccess +#define hipOccupancyDefault cudaOccupancyDefault + +#define hipCooperativeLaunchMultiDeviceNoPreSync cudaCooperativeLaunchMultiDeviceNoPreSync +#define hipCooperativeLaunchMultiDeviceNoPostSync cudaCooperativeLaunchMultiDeviceNoPostSync + + // enum CUjit_option redefines #define hipJitOptionMaxRegisters CU_JIT_MAX_REGISTERS #define hipJitOptionThreadsPerBlock CU_JIT_THREADS_PER_BLOCK @@ -267,6 +273,8 @@ inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) { return hipErrorNotInitialized; case cudaErrorLaunchFailure: return hipErrorLaunchFailure; + case cudaErrorCooperativeLaunchTooLarge: + return hipErrorCooperativeLaunchTooLarge; case cudaErrorPriorLaunchFailure: return hipErrorPriorLaunchFailure; case cudaErrorLaunchOutOfResources: @@ -445,6 +453,8 @@ inline static hipError_t hipCUResultTohipError(CUresult cuError) { return hipErrorHostMemoryNotRegistered; case CUDA_ERROR_LAUNCH_FAILED: return hipErrorLaunchFailure; + case CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE: + return hipErrorCooperativeLaunchTooLarge; case CUDA_ERROR_NOT_SUPPORTED: return hipErrorNotSupported; case CUDA_ERROR_UNKNOWN: @@ -601,6 +611,8 @@ inline static cudaError_t hipErrorToCudaError(hipError_t hError) { return cudaErrorSetOnActiveProcess; case hipErrorLaunchFailure: return cudaErrorLaunchFailure; + case hipErrorCooperativeLaunchTooLarge: + return cudaErrorCooperativeLaunchTooLarge; case hipErrorNotSupported: return cudaErrorNotSupported; // HSA: does not exist in CUDA diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 175d301ee1..85c08943bf 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -1796,6 +1796,8 @@ const char* ihipErrorString(hipError_t hip_error) { return "hipErrorMissingConfiguration"; case hipErrorLaunchFailure: return "hipErrorLaunchFailure"; + case hipErrorCooperativeLaunchTooLarge: + return "hipErrorCooperativeLaunchTooLarge"; case hipErrorPriorLaunchFailure: return "hipErrorPriorLaunchFailure"; case hipErrorLaunchTimeOut: From adc9134318fb7f5f43b0613b2346a40714d21660 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Fri, 31 Jan 2020 13:08:25 -0500 Subject: [PATCH 2/5] code clean up [ROCm/hip commit: c8137263d62af5f0016fda855de74f335f3c43a2] --- projects/hip/bin/hipify-perl | 5 ----- .../CUDA_Driver_API_functions_supported_by_HIP.md | 4 ++-- .../hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp | 8 ++++---- .../hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp | 8 ++++---- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 6 ++++-- 5 files changed, 14 insertions(+), 17 deletions(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index d18e227c3d..4fe9b80349 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -1186,7 +1186,6 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_PTX\b/hipErrorInvalidKernelFile/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_SOURCE\b/hipErrorInvalidSource/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_VALUE\b/hipErrorInvalidValue/g; - $ft{'numeric_literal'} += s/\bCUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE\b/hipErrorCooperativeLaunchTooLarge/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_LAUNCH_FAILED\b/hipErrorLaunchFailure/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_LAUNCH_OUT_OF_RESOURCES\b/hipErrorLaunchOutOfResources/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_LAUNCH_TIMEOUT\b/hipErrorLaunchTimeOut/g; @@ -1660,7 +1659,6 @@ sub simpleSubstitutions { $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/\bcudaErrorCooperativeLaunchTooLarge\b/hipErrorCooperativeLaunchTooLarge/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; @@ -1803,9 +1801,6 @@ sub simpleSubstitutions { $ft{'define'} += s/\bcudaTextureType3D\b/hipTextureType3D/g; $ft{'define'} += s/\bcudaTextureTypeCubemap\b/hipTextureTypeCubemap/g; $ft{'define'} += s/\bcudaTextureTypeCubemapLayered\b/hipTextureTypeCubemapLayered/g; - $ft{'define'} += s/\bcudaOccupancyDefault\b/hipOccupancyDefault/g; - $ft{'define'} += s/\bcudaCooperativeLaunchMultiDeviceNoPreSync\b/hipCooperativeLaunchMultiDeviceNoPreSync/g; - $ft{'define'} += s/\bcudaCooperativeLaunchMultiDeviceNoPostSync\b/hipCooperativeLaunchMultiDeviceNoPostSync/g; } # CUDA extern __shared__ syntax replace with HIP_DYNAMIC_SHARED() macro 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 5be6a4da8d..8956533ed0 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 @@ -1064,8 +1064,8 @@ | `cuFuncSetSharedMemConfig` | | | `cuLaunchKernel` | `hipModuleLaunchKernel` | | `cuLaunchHostFunc` | | 10.0 | -| `cuLaunchCooperativeKernel` | `hipLaunchCooperativeKernel` | 9.0 | -| `cuLaunchCooperativeKernelMultiDevice` | `hipLaunchCooperativeKernelMultiDevice` | 9.0 | +| `cuLaunchCooperativeKernel` | | 9.0 | +| `cuLaunchCooperativeKernelMultiDevice` | | 9.0 | ## **19. Execution Control [DEPRECATED]** 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 a4d419384d..054de19800 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -1047,7 +1047,7 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CUoccupancy_flags_enum", {"hipOccupancyFlags", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // CUoccupancy_flags enum values // cudaOccupancyDefault - {"CU_OCCUPANCY_DEFAULT", {"hipOccupancyDefault", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x00 + {"CU_OCCUPANCY_DEFAULT", {"hipOccupancyDefault", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x00 // cudaOccupancyDisableCachingOverride {"CU_OCCUPANCY_DISABLE_CACHING_OVERRIDE", {"hipOccupancyDisableCachingOverride", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01 @@ -1285,7 +1285,7 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ // cudaErrorLaunchFailure {"CUDA_ERROR_LAUNCH_FAILED", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 719 // cudaErrorCooperativeLaunchTooLarge - {"CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 720 + {"CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 720 // cudaErrorNotPermitted {"CUDA_ERROR_NOT_PERMITTED", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 800 // cudaErrorNotSupported @@ -1600,9 +1600,9 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ // cudaArrayColorAttachment {"CUDA_ARRAY3D_COLOR_ATTACHMENT", {"hipArrayColorAttachment", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x20 // cudaCooperativeLaunchMultiDeviceNoPreSync - {"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC", {"hipCooperativeLaunchMultiDeviceNoPreSync", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01 + {"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC", {"hipCooperativeLaunchMultiDeviceNoPreSync", "", CONV_DEFINE, API_DRIVER}}, // 0x01 // cudaCooperativeLaunchMultiDeviceNoPostSync - {"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC", {"hipCooperativeLaunchMultiDeviceNoPostSync", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02 + {"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC", {"hipCooperativeLaunchMultiDeviceNoPostSync", "", CONV_DEFINE, API_DRIVER}}, // 0x02 // cudaExternalMemoryDedicated {"CUDA_EXTERNAL_MEMORY_DEDICATED", {"hipExternalMemoryDedicated", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1 // cudaExternalSemaphoreSignalSkipNvSciBufMemSync 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 ade8b105f1..6eb9bfb2be 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp @@ -788,7 +788,7 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { // CUDA_ERROR_LAUNCH_FAILED {"cudaErrorLaunchFailure", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 719 // CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE - {"cudaErrorCooperativeLaunchTooLarge", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 720 + {"cudaErrorCooperativeLaunchTooLarge", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 720 // CUDA_ERROR_NOT_PERMITTED {"cudaErrorNotPermitted", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 800 // CUDA_ERROR_NOT_SUPPORTED @@ -1327,9 +1327,9 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { // CUDA_ARRAY3D_COLOR_ATTACHMENT {"cudaArrayColorAttachment", {"hipArrayColorAttachment", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x20 // CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC - {"cudaCooperativeLaunchMultiDeviceNoPreSync", {"hipCooperativeLaunchMultiDeviceNoPreSync", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x01 + {"cudaCooperativeLaunchMultiDeviceNoPreSync", {"hipCooperativeLaunchMultiDeviceNoPreSync", "", CONV_DEFINE, API_RUNTIME}}, // 0x01 // CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC - {"cudaCooperativeLaunchMultiDeviceNoPostSync", {"hipCooperativeLaunchMultiDeviceNoPostSync", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x02 + {"cudaCooperativeLaunchMultiDeviceNoPostSync", {"hipCooperativeLaunchMultiDeviceNoPostSync", "", CONV_DEFINE, API_RUNTIME}}, // 0x02 // CU_DEVICE_CPU ((CUdevice)-1) {"cudaCpuDeviceId", {"hipCpuDeviceId", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // ((int)-1) // CU_DEVICE_INVALID ((CUdevice)-2) @@ -1412,7 +1412,7 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { // no analogue {"cudaTextureTypeCubemapLayered", {"hipTextureTypeCubemapLayered", "", CONV_DEFINE, API_RUNTIME}}, // 0xFC // CU_OCCUPANCY_DEFAULT - {"cudaOccupancyDefault", {"hipOccupancyDefault", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x00 + {"cudaOccupancyDefault", {"hipOccupancyDefault", "", CONV_DEFINE, API_RUNTIME}}, // 0x00 // CU_OCCUPANCY_DISABLE_CACHING_OVERRIDE {"cudaOccupancyDisableCachingOverride", {"hipOccupancyDisableCachingOverride", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x01 // CU_STREAM_DEFAULT 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 1ee84b5057..bcddce58b2 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -142,8 +142,10 @@ typedef enum cudaChannelFormatKind hipChannelFormatKind; #define hipOccupancyDefault cudaOccupancyDefault -#define hipCooperativeLaunchMultiDeviceNoPreSync cudaCooperativeLaunchMultiDeviceNoPreSync -#define hipCooperativeLaunchMultiDeviceNoPostSync cudaCooperativeLaunchMultiDeviceNoPostSync +#define hipCooperativeLaunchMultiDeviceNoPreSync + cudaCooperativeLaunchMultiDeviceNoPreSync +#define hipCooperativeLaunchMultiDeviceNoPostSync + cudaCooperativeLaunchMultiDeviceNoPostSync // enum CUjit_option redefines From d5ea43964d4e97db3dae673dfc1810030c0b66a5 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Fri, 31 Jan 2020 14:36:37 -0500 Subject: [PATCH 3/5] update the hipify-perl [ROCm/hip commit: 5c03a003c3ba6115ad06527bd509e7d3125cce5b] --- projects/hip/bin/hipify-perl | 8 ++++++++ 1 file changed, 8 insertions(+) mode change 100755 => 100644 projects/hip/bin/hipify-perl diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl old mode 100755 new mode 100644 index 4fe9b80349..f6de5abae4 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -1172,6 +1172,7 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bCUDA_ERROR_ASSERT\b/hipErrorAssert/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_CONTEXT_ALREADY_CURRENT\b/hipErrorContextAlreadyCurrent/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_CONTEXT_ALREADY_IN_USE\b/hipErrorContextAlreadyInUse/g; + $ft{'numeric_literal'} += s/\bCUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE\b/hipErrorCooperativeLaunchTooLarge/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_DEINITIALIZED\b/hipErrorDeinitialized/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_ECC_UNCORRECTABLE\b/hipErrorECCNotCorrectable/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_FILE_NOT_FOUND\b/hipErrorFileNotFound/g; @@ -1528,6 +1529,7 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bCU_MEMORYTYPE_UNIFIED\b/hipMemoryTypeUnified/g; $ft{'numeric_literal'} += s/\bCU_MEM_ATTACH_GLOBAL\b/hipMemAttachGlobal/g; $ft{'numeric_literal'} += s/\bCU_MEM_ATTACH_HOST\b/hipMemAttachHost/g; + $ft{'numeric_literal'} += s/\bCU_OCCUPANCY_DEFAULT\b/hipOccupancyDefault/g; $ft{'numeric_literal'} += s/\bCU_RESOURCE_TYPE_ARRAY\b/hipResourceTypeArray/g; $ft{'numeric_literal'} += s/\bCU_RESOURCE_TYPE_LINEAR\b/hipResourceTypeLinear/g; $ft{'numeric_literal'} += s/\bCU_RESOURCE_TYPE_MIPMAPPED_ARRAY\b/hipResourceTypeMipmappedArray/g; @@ -1636,6 +1638,7 @@ sub simpleSubstitutions { $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/\bcudaErrorCooperativeLaunchTooLarge\b/hipErrorCooperativeLaunchTooLarge/g; $ft{'numeric_literal'} += s/\bcudaErrorCudartUnloading\b/hipErrorDeinitialized/g; $ft{'numeric_literal'} += s/\bcudaErrorDeviceAlreadyInUse\b/hipErrorContextAlreadyInUse/g; $ft{'numeric_literal'} += s/\bcudaErrorDeviceUninitialized\b/hipErrorInvalidContext/g; @@ -1749,6 +1752,8 @@ sub simpleSubstitutions { $ft{'define'} += s/\bCUDA_ARRAY3D_LAYERED\b/hipArrayLayered/g; $ft{'define'} += s/\bCUDA_ARRAY3D_SURFACE_LDST\b/hipArraySurfaceLoadStore/g; $ft{'define'} += s/\bCUDA_ARRAY3D_TEXTURE_GATHER\b/hipArrayTextureGather/g; + $ft{'define'} += s/\bCUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC\b/hipCooperativeLaunchMultiDeviceNoPostSync/g; + $ft{'define'} += s/\bCUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC\b/hipCooperativeLaunchMultiDeviceNoPreSync/g; $ft{'define'} += s/\bCU_LAUNCH_PARAM_BUFFER_POINTER\b/HIP_LAUNCH_PARAM_BUFFER_POINTER/g; $ft{'define'} += s/\bCU_LAUNCH_PARAM_BUFFER_SIZE\b/HIP_LAUNCH_PARAM_BUFFER_SIZE/g; $ft{'define'} += s/\bCU_LAUNCH_PARAM_END\b/HIP_LAUNCH_PARAM_END/g; @@ -1769,6 +1774,8 @@ sub simpleSubstitutions { $ft{'define'} += s/\bcudaArrayLayered\b/hipArrayLayered/g; $ft{'define'} += s/\bcudaArraySurfaceLoadStore\b/hipArraySurfaceLoadStore/g; $ft{'define'} += s/\bcudaArrayTextureGather\b/hipArrayTextureGather/g; + $ft{'define'} += s/\bcudaCooperativeLaunchMultiDeviceNoPostSync\b/hipCooperativeLaunchMultiDeviceNoPostSync/g; + $ft{'define'} += s/\bcudaCooperativeLaunchMultiDeviceNoPreSync\b/hipCooperativeLaunchMultiDeviceNoPreSync/g; $ft{'define'} += s/\bcudaDeviceBlockingSync\b/hipDeviceScheduleBlockingSync/g; $ft{'define'} += s/\bcudaDeviceLmemResizeToMax\b/hipDeviceLmemResizeToMax/g; $ft{'define'} += s/\bcudaDeviceMapHost\b/hipDeviceMapHost/g; @@ -1792,6 +1799,7 @@ sub simpleSubstitutions { $ft{'define'} += s/\bcudaIpcMemLazyEnablePeerAccess\b/hipIpcMemLazyEnablePeerAccess/g; $ft{'define'} += s/\bcudaMemAttachGlobal\b/hipMemAttachGlobal/g; $ft{'define'} += s/\bcudaMemAttachHost\b/hipMemAttachHost/g; + $ft{'define'} += s/\bcudaOccupancyDefault\b/hipOccupancyDefault/g; $ft{'define'} += s/\bcudaStreamDefault\b/hipStreamDefault/g; $ft{'define'} += s/\bcudaStreamNonBlocking\b/hipStreamNonBlocking/g; $ft{'define'} += s/\bcudaTextureType1D\b/hipTextureType1D/g; From 5db60dab96eaf32b80a15a8a3deeb800f308bb44 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Mon, 10 Feb 2020 11:44:40 -0500 Subject: [PATCH 4/5] Fix a typo causing a build error [ROCm/hip commit: 5a29f27455c15634ac4a95f4a09455b8824d4a0d] --- projects/hip/include/hip/hcc_detail/hip_runtime_api.h | 1 + 1 file changed, 1 insertion(+) 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 1636a782a0..76209ef6a7 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -2964,6 +2964,7 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( * @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block * @param [in] flags Extra flags for occupancy calculation (currently ignored) */ +hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags __dparm(hipOccupancyDefault)); #if __HIP_VDI__ && !defined(__HCC__) From 07c3e4f89841095360f60206b3fedb5b6e1072eb Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Tue, 11 Feb 2020 12:16:51 -0500 Subject: [PATCH 5/5] fix build error in nvcc path [ROCm/hip commit: 959f1b0f0e6dd1da478920940bdbd590d5c06a4b] --- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) 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 bcddce58b2..6e0d02d0c0 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -142,9 +142,9 @@ typedef enum cudaChannelFormatKind hipChannelFormatKind; #define hipOccupancyDefault cudaOccupancyDefault -#define hipCooperativeLaunchMultiDeviceNoPreSync +#define hipCooperativeLaunchMultiDeviceNoPreSync \ cudaCooperativeLaunchMultiDeviceNoPreSync -#define hipCooperativeLaunchMultiDeviceNoPostSync +#define hipCooperativeLaunchMultiDeviceNoPostSync \ cudaCooperativeLaunchMultiDeviceNoPostSync