From c3ab5a8df8864b47cc11ea6b0251a72166fc186e Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 9 Sep 2016 10:21:52 -0500 Subject: [PATCH 01/25] Added async memcpy driver api Change-Id: I90e8a078d668a408e79f9e1142e7534771467a4f --- hipamd/include/hcc_detail/hip_runtime_api.h | 6 +- hipamd/src/hip_memory.cpp | 79 +++++++++++++++++++++ hipamd/tests/src/hipDrvMemcpy.cpp | 28 ++++++-- 3 files changed, 107 insertions(+), 6 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index 71e46612de..c782696753 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hcc_detail/hip_runtime_api.h @@ -883,7 +883,11 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes); hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes); -hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes); +hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream); + +hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream); + +hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream); /** diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 8516d7520d..19b4745b59 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -576,6 +576,85 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp return ihipLogStatus(e); } +hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream) +{ + HIP_INIT_API(dst, src, sizeBytes, stream); + + hipError_t e = hipSuccess; + + stream = ihipSyncAndResolveStream(stream); + + hipMemcpyKind kind = hipMemcpyHostToDevice; + + if ((dst == NULL) || (src == NULL)) { + e= hipErrorInvalidValue; + } else if (stream) { + try { + stream->copyAsync((void*)dst, src, sizeBytes, kind); + } + catch (ihipException ex) { + e = ex._code; + } + } else { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); +} + +hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream) +{ + HIP_INIT_API(dst, src, sizeBytes, stream); + + hipError_t e = hipSuccess; + + hipMemcpyKind kind = hipMemcpyDeviceToDevice; + + stream = ihipSyncAndResolveStream(stream); + + + if ((dst == NULL) || (src == NULL)) { + e= hipErrorInvalidValue; + } else if (stream) { + try { + stream->copyAsync((void*)dst, (void*)src, sizeBytes, kind); + } + catch (ihipException ex) { + e = ex._code; + } + } else { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); +} + +hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream) +{ + HIP_INIT_API(dst, src, sizeBytes, stream); + + hipError_t e = hipSuccess; + + stream = ihipSyncAndResolveStream(stream); + + hipMemcpyKind kind = hipMemcpyDeviceToHost; + + if ((dst == NULL) || (src == NULL)) { + e= hipErrorInvalidValue; + } else if (stream) { + try { + stream->copyAsync(dst, (void*)src, sizeBytes, kind); + } + catch (ihipException ex) { + e = ex._code; + } + } else { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); +} + // dpitch, spitch, and width in bytes hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { diff --git a/hipamd/tests/src/hipDrvMemcpy.cpp b/hipamd/tests/src/hipDrvMemcpy.cpp index 55c9b18818..8322464964 100644 --- a/hipamd/tests/src/hipDrvMemcpy.cpp +++ b/hipamd/tests/src/hipDrvMemcpy.cpp @@ -6,21 +6,39 @@ #define SIZE LEN<<2 int main(){ - int *A, *B, *C; - hipDeviceptr Ad, Bd; + int *A, *B; + hipDeviceptr_t Ad, Bd; A = new int[LEN]; B = new int[LEN]; - C = new int[LEN]; + for(int i=0;i Date: Fri, 9 Sep 2016 10:25:57 -0500 Subject: [PATCH 02/25] Added async memcpy driver api for nvcc backend Change-Id: If401fc949b834c883cff9de74adb0d532a0fd942 --- hipamd/include/nvcc_detail/hip_runtime_api.h | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/hipamd/include/nvcc_detail/hip_runtime_api.h b/hipamd/include/nvcc_detail/hip_runtime_api.h index 8e9b0d92a7..1ad1af2267 100644 --- a/hipamd/include/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/nvcc_detail/hip_runtime_api.h @@ -232,6 +232,24 @@ inline static hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, return hipCUResultTohipError(cuMemcpyDtoD(dst, src, size)); } +inline static hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, + void* src, size_t size, hipStream_t stream) +{ + return hipCUResultTohipError(cuMemcpyHtoDAsync(dst, src, size, stream)); +} + +inline static hipError_t hipMemcpyDtoHAsync(void* dst, + hipDeviceptr_t src, size_t size, hipStream_t stream) +{ + return hipCUResultTohipError(cuMemcpyDtoH(dst, src, size, stream)); +} + +inline static hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, + hipDeviceptr_t src, size_t size, hipStream_t stream) +{ + return hipCUResultTohipError(cuMemcpyDtoD(dst, src, size, stream)); +} + inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind) { return hipCUDAErrorTohipError(cudaMemcpy(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind))); } From e1d00da7d259ce318bff12a24b9fe1b226fc401c Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Mon, 12 Sep 2016 21:06:54 -0500 Subject: [PATCH 03/25] Cleaned doxygen comments for device management runtime apis Change-Id: I8fbba7d997589a67218cfaa132c598ff750ac544 --- hipamd/include/hcc_detail/hip_runtime_api.h | 84 +++++++++++++-------- 1 file changed, 53 insertions(+), 31 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index c782696753..2b1b116c6e 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hcc_detail/hip_runtime_api.h @@ -189,9 +189,12 @@ typedef enum hipMemcpyKind { */ /** - * @brief Blocks until the default device has completed all preceding requested tasks. + * @brief Waits on all active streams on current device * - * This function waits for all streams on the default device to complete execution, and then returns. + * When this command is invoked, the host thread gets blocked until all the commands associated + * with streams associated with the device. HIP does not support multiple blocking modes (yet!). + * + * @returns hipSuccess * * @see hipSetDevice, hipDeviceReset */ @@ -200,12 +203,12 @@ hipError_t hipDeviceSynchronize(void); /** - * @brief Destroy all resources and reset all state on the default device in the current process. + * @brief The state of current device is discarded and updated to a fresh state. * - * Explicity destroy all memory allocations, events, and queues associated with the default device in the current process. + * Calling this function deletes all streams created, memory allocated, kernels running, events created. + * Make sure that no other thread is using the device or streams, memory, kernels, events associated with the current device. * - * This function will reset the device immmediately, and then return after all resources have been freed. - * The caller must ensure that the device is not being accessed by any other host threads from the active process when this function is called. + * @returns hipSuccess * * @see hipDeviceSynchronize */ @@ -234,6 +237,8 @@ hipError_t hipDeviceReset(void) ; * Thread-pool implementations may inherit the default device of the previous thread. A good practice is to always call hipSetDevice * at the start of HIP coding sequency to establish a known standard device. * + * @returns hipSuccess, hipErrorInvalidDevice, hipErrorDeviceAlreadyInUse + * * @see hipGetDevice, hipGetDeviceCount */ hipError_t hipSetDevice(int deviceId); @@ -250,25 +255,32 @@ hipError_t hipSetDevice(int deviceId); * * @see hipSetDevice, hipGetDevicesizeBytes * - * @returns hipSuccess, hipErrorInvalidDevice + * @returns hipSuccess */ hipError_t hipGetDevice(int *deviceId); /** * @brief Return number of compute-capable devices. + * * @param [output] count Returns number of compute-capable devices. * + * @returns hipSuccess, hipErrorNoDevice + * + * * Returns in @p *count the number of devices that have ability to run compute commands. If there are no such devices, then @ref hipGetDeviceCount will return #hipErrorNoDevice. * If 1 or more devices can be found, then hipGetDeviceCount returns #hipSuccess. */ hipError_t hipGetDeviceCount(int *count); /** - * @brief Query device attribute. + * @brief Query for a specific device attribute. + * * @param [out] pi pointer to value to return * @param [in] attr attribute to query * @param [in] deviceId which device to query for information + * + * @returns hipSuccess, hipErrorInvalidDevice, hipErrorInvalidValue */ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int deviceId); @@ -278,7 +290,7 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int deviceI * @param [out] prop written with device properties * @param [in] deviceId which device to query for information * - * @return #hipSuccess, #hipErrorInvalidDevice + * @return hipSuccess, hipErrorInvalidDevice * @bug HCC always returns 0 for maxThreadsPerMultiProcessor * @bug HCC always returns 0 for regsPerBlock * @bug HCC always returns 0 for l2CacheSize @@ -288,13 +300,12 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int deviceI hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId); - -//Cache partitioning functions: - /** * @brief Set L1/Shared cache partition. * - * @returns #hipSuccess + * @param [in] cacheConfig + * + * @returns hipSuccess, hipErrorInitializationError * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures. * */ @@ -304,7 +315,9 @@ hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig ); /** * @brief Set Cache configuration for a specific function * - * @returns #hipSuccess + * @param [in] cacheConfig + * + * @returns hipSuccess, hipErrorInitializationError * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures. * */ @@ -314,19 +327,21 @@ hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig ); /** * @brief Set Cache configuration for a specific function * - * @returns #hipSuccess + * @param [in] config; + * + * @returns hipSuccess, hipErrorInitializationError * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures. * */ hipError_t hipFuncSetCacheConfig ( hipFuncCache config ); -//--- -//Shared bank config functions: - /** - * @brief Get Shared memory bank configuration. + * @brief Returns bank width of shared memory for current device + * + * @param [out] pConfig + * + * @returns hipSuccess, hipErrorInvalidValue, hipErrorInitializationError * - * @returns #hipSuccess * Note: AMD devices and recent Nvidia GPUS do not support shared cache banking, and the hint is ignored on those architectures. * */ @@ -334,36 +349,43 @@ hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ); /** - * @brief Set Shared memory bank configuration. + * @brief The bank width of shared memory on current device is set + * + * @param [in] config + * + * @returns hipSuccess, hipErrorInvalidValue, hipErrorInitializationError * - * @returns #hipSuccess * Note: AMD devices and recent Nvidia GPUS do not support shared cache banking, and the hint is ignored on those architectures. * */ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ); /** - * @brief Set Device flags + * @brief The current device behavior is changed according the flags passed. + * + * @param [in] flags + * + * @returns hipSuccess, hipErrorInvalidDevice, hipErrorSetOnActiveProcess * - * @returns #hipSuccess * Note: Only hipDeviceScheduleAuto and hipDeviceMapHost are supported * */ hipError_t hipSetDeviceFlags ( unsigned flags); -// end doxygen Device /** - * @} - */ - -/** - * @brief Select compute-device which best matches criteria. + * @brief Device which matches hipDeviceProp_t is returned * * @param [out] device ID * @param [in] device properties pointer * + * @returns hipSuccess, hipErrorInvalidValue + */ +hipError_t hipChooseDevice(int *device, hipDeviceProp_t* prop); + +// end doxygen Device +/** + * @} */ -hipError_t hipChooseDevice(int *device,hipDeviceProp_t* prop); /** *------------------------------------------------------------------------------------------------- From c6808ddc0e1298c3df8001c34ebb2df10f3e0dd2 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 13 Sep 2016 16:54:46 +0530 Subject: [PATCH 04/25] Updated Doxygen related documentation for Error APIs Change-Id: I16873503d8e96df91831e2a0503835a1ff0220e3 --- hipamd/include/hcc_detail/hip_runtime_api.h | 6 ++++-- hipamd/src/hip_error.cpp | 13 +------------ 2 files changed, 5 insertions(+), 14 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index 2b1b116c6e..a08974f8b5 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hcc_detail/hip_runtime_api.h @@ -397,9 +397,12 @@ hipError_t hipChooseDevice(int *device, hipDeviceProp_t* prop); /** * @brief Return last error returned by any HIP runtime API call and resets the stored error code to #hipSuccess * + * @returns return code from last HIP called from the active host thread + * * Returns the last error that has been returned by any of the runtime calls in the same host thread, * and then resets the saved error to #hipSuccess. * + * @see hipGetErrorString, hipGetLastError, hipPeakAtLastError, hipError_t */ hipError_t hipGetLastError(void); @@ -412,8 +415,7 @@ hipError_t hipGetLastError(void); * Returns the last error that has been returned by any of the runtime calls in the same host thread. * Unlike hipGetLastError, this function does not reset the saved error code. * - * - * + * @see hipGetErrorString, hipGetLastError, hipPeakAtLastError, hipError_t */ hipError_t hipPeekAtLastError(void); diff --git a/hipamd/src/hip_error.cpp b/hipamd/src/hip_error.cpp index d9c6dd9aa9..97fcaaf714 100644 --- a/hipamd/src/hip_error.cpp +++ b/hipamd/src/hip_error.cpp @@ -25,9 +25,7 @@ THE SOFTWARE. //------------------------------------------------------------------------------------------------- // Error Handling //--- -/** - * @returns return code from last HIP called from the active host thread. - */ + hipError_t hipGetLastError() { HIP_INIT_API(); @@ -38,8 +36,6 @@ hipError_t hipGetLastError() return e; } - -//--- hipError_t hipPeekAtLastError() { HIP_INIT_API(); @@ -48,7 +44,6 @@ hipError_t hipPeekAtLastError() return ihipLogStatus(tls_lastHipError); } -//--- const char *hipGetErrorName(hipError_t hip_error) { HIP_INIT_API(hip_error); @@ -56,12 +51,6 @@ const char *hipGetErrorName(hipError_t hip_error) return ihipErrorString(hip_error); } - -/** - * @warning : hipGetErrorString returns string from hipGetErrorName - */ - -//--- const char *hipGetErrorString(hipError_t hip_error) { HIP_INIT_API(hip_error); From 9155e1f2fd4a63b2d1050c0ee98634a6776ff0fc Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 13 Sep 2016 17:33:11 +0530 Subject: [PATCH 05/25] Partial- Doxygen related changes for Memory APIs Change-Id: I7ccc143571e0b26c16c70262565d6e69d6e7b6b5 --- hipamd/include/hcc_detail/hip_runtime_api.h | 59 ++++++++++++++++----- 1 file changed, 45 insertions(+), 14 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index a08974f8b5..221dfbf768 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hcc_detail/hip_runtime_api.h @@ -739,6 +739,13 @@ hipError_t hipEventQuery(hipEvent_t event) ; /** * @brief Return attributes for the specified pointer + * + * @param[out] attributes for the specified pointer + * @param[in] pointer to get attributes for + * + * @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue + * + * @see hipGetDeviceCount, hipGetDevice, hipSetDevice, hipChooseDevice */ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr); @@ -747,7 +754,10 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) * * @param[out] ptr Pointer to the allocated memory * @param[in] size Requested memory size + * * @return #hipSuccess + * + * @see hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D, hipMalloc3DArray, hipMallocHost, hipFreeHost, hipHostAlloc */ hipError_t hipMalloc(void** ptr, size_t size) ; @@ -755,19 +765,25 @@ hipError_t hipMalloc(void** ptr, size_t size) ; /** * @brief Allocate pinned host memory * - * @param[out] ptr Pointer to the allocated host pinned memory - * @param[in] size Requested memory size - * @return Error code + * @param[out] ptr Pointer to the allocated host pinned memory + * @param[in] size Requested memory size + * + * @return #hipSuccess, #hipErrorMemoryAllocation + * + * @see hipMalloc, hipMallocPitch, hipMallocArray, hipMalloc3D, hipMalloc3DArray, hipHostAlloc, hipFree, hipFreeArray, hipMallocHost, hipFreeHost, hipHostAlloc */ hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use hipHostMalloc instead"))) ; /** * @brief Allocate device accessible page locked host memory * - * @param[out] ptr Pointer to the allocated host pinned memory - * @param[in] size Requested memory size - * @param[in] flags Type of host memory allocation - * @return Error code + * @param[out] ptr Pointer to the allocated host pinned memory + * @param[in] size Requested memory size + * @param[in] flags Type of host memory allocation + * + * @return #hipSuccess, #hipErrorMemoryAllocation + * + * @see hipSetDeviceFlags, hipMallocHost, hipFreeHost */ hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags) ; hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) __attribute__((deprecated("use hipHostMalloc instead"))) ;; @@ -775,19 +791,24 @@ hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) __attribute /** * @brief Get Device pointer from Host Pointer allocated through hipHostAlloc * - * @param[out] dstPtr Device Pointer mapped to passed host pointer - * @param[in] hstPtr Host Pointer allocated through hipHostAlloc - * @param[in] flags Flags to be passed for extension - * @return Error code + * @param[out] dstPtr Device Pointer mapped to passed host pointer + * @param[in] hstPtr Host Pointer allocated through hipHostAlloc + * @param[in] flags Flags to be passed for extension + * + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryAllocation + * + * @see hipSetDeviceFlags, hipHostAlloc */ hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, unsigned int flags) ; /** * @brief Return flags associated with host pointer * - * @param[out] flagsPtr Memory location to store flags - * @param[in] hostPtr Host Pointer allocated through hipHostMalloc - * @return Error code + * @param[out] flagsPtr Memory location to store flags + * @param[in] hostPtr Host Pointer allocated through hipHostMalloc + * @return #hipSuccess, #hipErrorInvalidValue + * + * @see hipHostAlloc */ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ; @@ -820,6 +841,8 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ; * from the other registered memory region. * * @return #hipSuccess, #hipErrorMemoryAllocation + * + * @see hipHostUnregister, hipHostGetFlags, hipHostGetDevicePointer */ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) ; @@ -828,6 +851,8 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) * * @param[in] hostPtr Host pointer previously registered with #hipHostRegister * @return Error code + * + * @see hipHostRegister */ hipError_t hipHostUnregister(void* hostPtr) ; @@ -842,6 +867,8 @@ hipError_t hipHostUnregister(void* hostPtr) ; * @param[in] width Requested pitched allocation width (in bytes) * @param[in] height Requested pitched allocation height * @return Error code + * + * @see hipMalloc, hipFree, hipMallocArray, hipFreeArray, hipMallocHost, hipFreeHost, hipMalloc3D, hipMalloc3DArray, hipHostAlloc */ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height); @@ -854,6 +881,8 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height * @param[in] ptr Pointer to memory to be freed * @return #hipSuccess * @return #hipErrorInvalidDevicePointer (if pointer is invalid, including host pointers allocated with hipHostMalloc) + * + * @see hipMalloc, hipMallocPitch, hipMallocArray, hipFreeArray, hipMallocHost, hipFreeHost, hipMalloc3D, hipMalloc3DArray, hipHostAlloc */ hipError_t hipFree(void* ptr); @@ -875,6 +904,8 @@ hipError_t hipFreeHost(void* ptr) __attribute__((deprecated("use hipHostFree ins * @param[in] ptr Pointer to memory to be freed * @return #hipSuccess, * #hipErrorInvalidValue (if pointer is invalid, including device pointers allocated with hipMalloc) + * + * @see hipMalloc, hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMallocHost, hipMalloc3D, hipMalloc3DArray, hipHostAlloc */ hipError_t hipHostFree(void* ptr); From 50786a4b705497656d5c95401bd1881517d5b849 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 13 Sep 2016 15:11:21 +0300 Subject: [PATCH 06/25] [HIPIFY] CUDA Driver API porting to HIP : Init and Memory functions support. --- hipamd/clang-hipify/src/Cuda2Hip.cpp | 44 ++++++++++++++++++++++++++-- 1 file changed, 41 insertions(+), 3 deletions(-) diff --git a/hipamd/clang-hipify/src/Cuda2Hip.cpp b/hipamd/clang-hipify/src/Cuda2Hip.cpp index 77d8b0259e..885626f0b4 100644 --- a/hipamd/clang-hipify/src/Cuda2Hip.cpp +++ b/hipamd/clang-hipify/src/Cuda2Hip.cpp @@ -322,6 +322,12 @@ struct cuda2hipMap { cuda2hipRename["CU_STREAM_DEFAULT"] = {"hipStreamDefault", CONV_STREAM, API_DRIVER}; cuda2hipRename["CU_STREAM_NON_BLOCKING"] = {"hipStreamNonBlocking", CONV_STREAM, API_DRIVER}; + // Init + cuda2hipRename["cuInit"] = {"hipInit", CONV_DRIVER, API_DRIVER}; + + // Driver + cuda2hipRename["cuDriverGetVersion"] = {"hipDriverGetVersion", CONV_DRIVER, API_DRIVER}; + // Context cuda2hipRename["cuCtxCreate_v2"] = {"hipCtxCreate", CONV_CONTEXT, API_DRIVER}; cuda2hipRename["cuCtxDestroy_v2"] = {"hipCtxDestroy", CONV_CONTEXT, API_DRIVER}; @@ -356,9 +362,6 @@ struct cuda2hipMap { cuda2hipRename["cuDeviceComputeCapability"] = {"hipDeviceComputeCapability", CONV_DEV, API_DRIVER}; cuda2hipRename["cuDeviceCanAccessPeer"] = {"hipDeviceCanAccessPeer", CONV_DEV, API_DRIVER}; - // Driver - cuda2hipRename["cuDriverGetVersion"] = {"hipDriverGetVersion", CONV_DRIVER, API_DRIVER}; - // Events cuda2hipRename["cuEventCreate"] = {"hipEventCreate", CONV_EVENT, API_DRIVER}; cuda2hipRename["cuEventDestroy_v2"] = {"hipEventDestroy", CONV_EVENT, API_DRIVER}; @@ -387,6 +390,41 @@ struct cuda2hipMap { cuda2hipRename["cuStreamSynchronize"] = {"hipStreamSynchronize", CONV_STREAM, API_DRIVER}; cuda2hipRename["cuStreamWaitEvent"] = {"hipStreamWaitEvent", CONV_STREAM, API_DRIVER}; + // Memory management + cuda2hipRename["cuMemAlloc_v2"] = {"hipMalloc", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemFree_v2"] = {"hipFree", CONV_MEM, API_DRIVER}; + + cuda2hipRename["cuMemHostAlloc"] = {"hipHostMalloc", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemFreeHost"] = {"hipHostFree", CONV_MEM, API_DRIVER}; + + cuda2hipRename["cuMemcpyDtoD_v2"] = {"hipMemcpyDtoD", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemcpyDtoDAsync_v2"] = {"hipMemcpyDtoDAsync", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemcpyDtoH_v2"] = {"hipMemcpyDtoH", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemcpyDtoHAsync_v2"] = {"hipMemcpyDtoHAsync", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemcpyHtoD_v2"] = {"hipMemcpyHtoD", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemcpyHtoDAsync_v2"] = {"hipMemcpyHtoDAsync", CONV_MEM, API_DRIVER}; + + // unsupported yet by HIP + // cuda2hipRename["cuMemsetD8_v2"] = {"hipMemsetD8", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD8Async"] = {"hipMemsetD8Async", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD2D8_v2"] = {"hipMemsetD2D8", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD2D8Async"] = {"hipMemsetD2D8Async", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD16_v2"] = {"hipMemsetD16", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD16Async"] = {"hipMemsetD16Async", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD2D16_v2"] = {"hipMemsetD2D16", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD2D16Async"] = {"hipMemsetD2D16Async", CONV_STREAM, API_DRIVER}; + cuda2hipRename["cuMemsetD32_v2"] = {"hipMemset", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemsetD32Async"] = {"hipMemsetAsync", CONV_MEM, API_DRIVER}; + // unsupported yet by HIP + // cuda2hipRename["cuMemsetD2D32_v2"] = {"hipMemsetD2D32", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD2D32Async"] = {"hipMemsetD2D32Async", CONV_STREAM, API_DRIVER}; + + cuda2hipRename["cuMemGetInfo_v2"] = {"hipMemGetInfo", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemHostRegister_v2"] = {"hipHostRegister", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemHostUnregister"] = {"hipHostUnregister", CONV_MEM, API_DRIVER}; + + + /////////////////////////////// CUDA RT API /////////////////////////////// // Error API cuda2hipRename["cudaGetLastError"] = {"hipGetLastError", CONV_ERR, API_RUNTIME}; From d85012df428c2c16539ed4ed7fa88e631e3c1e52 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 13 Sep 2016 14:03:07 +0530 Subject: [PATCH 07/25] Directed tests: Refactor phase 2 - Restore HIP_BUILD_LOCAL flag. Defaults to 1. - HIP_PATH is used only when HIP_BUILD_LOCAL is 0. - HIP_PLATFORM, HCC_HOME, HIP_LIB_TYPE are passed on to HIP build. - Specifying HCC_HOME automatically sets HIP_DEVELOPER=1 for HIP build. - Always show HIP configure step output. Output is colored magenta. - Fix code indentation. Change-Id: I74f2d9c1cb04bf865313b2db6b55fda567c8e071 --- hipamd/tests/src/CMakeLists.txt | 301 ++++++++++++++++++-------------- 1 file changed, 171 insertions(+), 130 deletions(-) diff --git a/hipamd/tests/src/CMakeLists.txt b/hipamd/tests/src/CMakeLists.txt index 43a74910b2..2563c31bab 100644 --- a/hipamd/tests/src/CMakeLists.txt +++ b/hipamd/tests/src/CMakeLists.txt @@ -1,4 +1,4 @@ -cmake_minimum_required (VERSION 2.6) +cmake_minimum_required(VERSION 2.6) # remove CMAKE_CXX_COMPILER entry from cache since it will be pointing to hipcc unset(CMAKE_CXX_COMPILER CACHE) @@ -9,69 +9,109 @@ project(HIP_Unit_Tests) include(CTest) set(HIPTEST_SOURCE_DIR ${PROJECT_SOURCE_DIR}) +string(ASCII 27 Esc) +set(ColorReset "${Esc}[m") +set(Red "${Esc}[31m") +set(Magenta "${Esc}[35m") + # Enable multi-gpu tests if(NOT DEFINED HIP_MULTI_GPU) set(HIP_MULTI_GPU 0 CACHE BOOL "Run tests requiring more than one GPU") endif() -# Determine HIP_PATH -if(NOT DEFINED HIP_PATH) - if(NOT DEFINED ENV{HIP_PATH}) - # We are going to use HIP source... - get_filename_component(HIP_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/../.. ABSOLUTE) - execute_process( - COMMAND "${CMAKE_COMMAND}" -E remove_directory hip - OUTPUT_QUIET - ERROR_QUIET - ) - execute_process( - COMMAND "${CMAKE_COMMAND}" -E make_directory hip - OUTPUT_QUIET - ERROR_QUIET - ) - message(STATUS "Configuring HIP") - # ...so need to build HIP locally. - execute_process( - COMMAND "${CMAKE_COMMAND}" -DCMAKE_INSTALL_PREFIX=${CMAKE_CURRENT_BINARY_DIR}/hip/localbuild ${HIP_SRC_PATH} - WORKING_DIRECTORY hip - RESULT_VARIABLE hip_build_result - OUTPUT_QUIET - ERROR_QUIET - ) - if(hip_build_result) - message(FATAL_ERROR "Error configuring HIP") - else() - message(STATUS "Configuring HIP - done") - message(STATUS "Building HIP") - endif() - execute_process( - COMMAND "${CMAKE_COMMAND}" --build . --target install - WORKING_DIRECTORY hip - RESULT_VARIABLE hip_build_result - OUTPUT_VARIABLE hip_build_log - ERROR_QUIET - ) - if(hip_build_result) - message(${hip_build_log}) - message(FATAL_ERROR "Error building HIP") - else() - # Building HIP is successful. Point HIP_PATH to this location. - message(STATUS "Building HIP - done") - get_filename_component(HIP_PATH ${CMAKE_CURRENT_BINARY_DIR}/hip/localbuild ABSOLUTE) - endif() - # Add a target to rebuild HIP if HIP source changes. - add_custom_target( - hip ALL - COMMAND "${CMAKE_COMMAND}" --build . --target install - WORKING_DIRECTORY hip - ) +# Determine HIP_BUILD_LOCAL +if(NOT DEFINED HIP_BUILD_LOCAL) + if(NOT DEFINED ENV{HIP_BUILD_LOCAL}) + set(HIP_BUILD_LOCAL 1 CACHE BOOL "Build HIP in local folder") else() - # We are using HIP_PATH from env. So just create a fake target. - set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to installed HIP") - add_custom_target(hip ALL) + set(HIP_BUILD_LOCAL $ENV{HIP_BUILD_LOCAL} CACHE BOOL "Build HIP in local folder") endif() +endif() + +# Determine HIP_PATH +if(HIP_BUILD_LOCAL) + # We are going to use HIP source... + get_filename_component(HIP_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/../.. ABSOLUTE) + + # ...so we first need to determine the options to cascade to HIP build. + if(DEFINED HIP_PLATFORM) + set(ENV{HIP_PLATFORM} ${HIP_PLATFORM}) + endif() + if(DEFINED HCC_HOME) + get_filename_component(HCC_HOME ${HCC_HOME} ABSOLUTE) + set(ENV{HCC_HOME} ${HCC_HOME}) + set(ENV{HIP_DEVELOPER} 1) + endif() + if(DEFINED HIP_LIB_TYPE) + set(ENV{HIP_LIB_TYPE} ${HIP_LIB_TYPE}) + endif() + + # Purge previous HIP installation... + execute_process( + COMMAND "${CMAKE_COMMAND}" -E remove_directory hip + OUTPUT_QUIET + ERROR_QUIET + ) + execute_process( + COMMAND "${CMAKE_COMMAND}" -E make_directory hip + OUTPUT_QUIET + ERROR_QUIET + ) + message(STATUS "Configuring HIP") + + # ...and now build HIP locally. + execute_process( + COMMAND "${CMAKE_COMMAND}" -DCMAKE_INSTALL_PREFIX=${CMAKE_CURRENT_BINARY_DIR}/hip/localbuild ${HIP_SRC_PATH} + WORKING_DIRECTORY hip + RESULT_VARIABLE hip_build_result + OUTPUT_VARIABLE hip_build_log + OUTPUT_STRIP_TRAILING_WHITESPACE + ERROR_QUIET + ) + message("${Magenta}${hip_build_log}${ColorReset}") + + if(hip_build_result) + message(FATAL_ERROR "Error configuring HIP") + else() + message(STATUS "Configuring HIP - done") + message(STATUS "Building HIP") + endif() + execute_process( + COMMAND "${CMAKE_COMMAND}" --build . --target install + WORKING_DIRECTORY hip + RESULT_VARIABLE hip_build_result + OUTPUT_VARIABLE hip_build_log + OUTPUT_STRIP_TRAILING_WHITESPACE + ERROR_QUIET + ) + + # Show HIP build errors if any. + if(hip_build_result) + message("${Red}${hip_build_log}${ColorReset}") + message(FATAL_ERROR "Error building HIP") + else() + # Building HIP is successful. Point HIP_PATH to this location. + message(STATUS "Building HIP - done") + get_filename_component(HIP_PATH ${CMAKE_CURRENT_BINARY_DIR}/hip/localbuild ABSOLUTE) + endif() + + # Add a target to rebuild HIP if HIP source changes. + add_custom_target( + hip ALL + COMMAND "${CMAKE_COMMAND}" --build . --target install + WORKING_DIRECTORY hip + ) else() - # We are using HIP_PATH passed to cmake. So just create a fake target. + if(NOT DEFINED HIP_PATH) + if(NOT DEFINED ENV{HIP_PATH}) + # We set to to default HIP installation path. + set(HIP_PATH /opt/rocm/hip CACHE PATH "Path to installed HIP") + else() + # We are using HIP_PATH from env. + set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to installed HIP") + endif() + endif() + # Create a fake target since we are using an installed HIP. add_custom_target(hip ALL) endif() MESSAGE("HIP_PATH=" ${HIP_PATH}) @@ -80,123 +120,123 @@ MESSAGE("HIP_PATH=" ${HIP_PATH}) execute_process(COMMAND ${HIP_PATH}/bin/hipconfig --platform OUTPUT_VARIABLE HIP_PLATFORM) if(${HIP_PLATFORM} STREQUAL "hcc") - MESSAGE("HIP_PLATFORM=hcc") + MESSAGE("HIP_PLATFORM=hcc") elseif(${HIP_PLATFORM} STREQUAL "nvcc") - MESSAGE("HIP_PLATFORM=nvcc") + MESSAGE("HIP_PLATFORM=nvcc") - #Need C++11 for threads in some of the tests. - add_definitions(-std=c++11) + #Need C++11 for threads in some of the tests. + add_definitions(-std=c++11) - # NVCC does not not support -rdynamic option - set(CMAKE_SHARED_LIBRARY_LINK_CXX_FLAGS ) - set(CMAKE_SHARED_LIBRARY_LINK_C_FLAGS ) + # NVCC does not not support -rdynamic option + set(CMAKE_SHARED_LIBRARY_LINK_CXX_FLAGS) + set(CMAKE_SHARED_LIBRARY_LINK_C_FLAGS) else() - MESSAGE(FATAL_ERROR "UNKNOWN HIP_PLATFORM=" ${HIP_PLATFORM}) + MESSAGE(FATAL_ERROR "UNKNOWN HIP_PLATFORM=" ${HIP_PLATFORM}) endif() set(HIPCC ${HIP_PATH}/bin/hipcc) set(CMAKE_CXX_COMPILER ${HIPCC} CACHE FILEPATH "CXX Compiler" FORCE) -add_library(test_common OBJECT test_common.cpp ) +add_library(test_common OBJECT test_common.cpp) -# usage : build_hip_executable (exe_name CPP_FILES) -macro (build_hip_executable exe cpp) - add_executable (${exe} ${cpp} ${ARGN} $ ) +# usage : build_hip_executable(exe_name CPP_FILES) +macro(build_hip_executable exe cpp) + add_executable(${exe} ${cpp} ${ARGN} $) add_dependencies(${exe} hip) endmacro() # Make a hip executable, using libc++ -macro (build_hip_executable_libcpp exe cpp) - build_hip_executable( ${exe} ${cpp} ${ARGN} ) - if (${HIP_PLATFORM} STREQUAL "hcc") - set_source_files_properties (${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --stdlib=libc++ ) +macro(build_hip_executable_libcpp exe cpp) + build_hip_executable( ${exe} ${cpp} ${ARGN}) + if(${HIP_PLATFORM} STREQUAL "hcc") + set_source_files_properties(${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --stdlib=libc++) endif() endmacro() -function (make_named_test exe testname ) - add_test (NAME ${testname} - COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN} - ) - set_tests_properties (${testname} - PROPERTIES PASS_REGULAR_EXPRESSION "PASSED" - ) +function(make_named_test exe testname) + add_test(NAME ${testname} + COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN} + ) + set_tests_properties(${testname} + PROPERTIES PASS_REGULAR_EXPRESSION "PASSED" + ) endfunction() -macro (make_test exe ) - string (REPLACE " " "" smush_args ${ARGN}) - set (testname ${PROJECT_NAME}/${exe}${smush_args}.tst) +macro(make_test exe) + string(REPLACE " " "" smush_args ${ARGN}) + set(testname ${PROJECT_NAME}/${exe}${smush_args}.tst) - make_named_test(${exe} ${testname} ${ARGN}) + make_named_test(${exe} ${testname} ${ARGN}) endmacro() -macro (make_hipify_test sourceFile ) - #string (REPLACE " " "" smush_args ${ARGN}) - set (testname ${sourceFile}${smush_args}.tst) +macro(make_hipify_test sourceFile) + #string(REPLACE " " "" smush_args ${ARGN}) + set(testname ${sourceFile}${smush_args}.tst) - add_test (NAME ${testname} - COMMAND ${HIP_PATH}/bin/hipify ${PROJECT_SOURCE_DIR}/${sourceFile} ${ARGN} - ) + add_test(NAME ${testname} + COMMAND ${HIP_PATH}/bin/hipify ${PROJECT_SOURCE_DIR}/${sourceFile} ${ARGN} + ) endmacro() -macro (make_test_matches exe match_string) - string (REPLACE " " "" smush_args ${ARGN}) - set (testname ${exe}${smush_args}.tst) - add_test (NAME ${testname} - COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN} - ) - set_tests_properties (${testname} - PROPERTIES PASS_REGULAR_EXPRESSION ${match_string} - ) +macro(make_test_matches exe match_string) + string(REPLACE " " "" smush_args ${ARGN}) + set(testname ${exe}${smush_args}.tst) + add_test(NAME ${testname} + COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN} + ) + set_tests_properties(${testname} + PROPERTIES PASS_REGULAR_EXPRESSION ${match_string} + ) endmacro() -macro (build_hip_executable_sm35 exe cpp) - build_hip_executable( ${exe} ${cpp} ${ARGN} ) - if (${HIP_PLATFORM} STREQUAL "nvcc") - set_source_files_properties (${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --gpu-architecture=sm_35 ) - endif() +macro(build_hip_executable_sm35 exe cpp) + build_hip_executable( ${exe} ${cpp} ${ARGN}) + if(${HIP_PLATFORM} STREQUAL "nvcc") + set_source_files_properties(${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --gpu-architecture=sm_35) + endif() endmacro() -build_hip_executable (hipGetDeviceAttribute hipGetDeviceAttribute.cpp) -build_hip_executable (hipEnvVar hipEnvVar.cpp) -build_hip_executable (hipEnvVarDriver hipEnvVarDriver.cpp) -build_hip_executable (hipEventRecord hipEventRecord.cpp) +build_hip_executable(hipGetDeviceAttribute hipGetDeviceAttribute.cpp) +build_hip_executable(hipEnvVar hipEnvVar.cpp) +build_hip_executable(hipEnvVarDriver hipEnvVarDriver.cpp) +build_hip_executable(hipEventRecord hipEventRecord.cpp) -build_hip_executable_libcpp (hipHcc hipHcc.cpp) -#set_source_files_properties (hipHcc.cpp PROPERTIES COMPILE_FLAGS --stdlib=libc++ ) +build_hip_executable_libcpp(hipHcc hipHcc.cpp) +#set_source_files_properties(hipHcc.cpp PROPERTIES COMPILE_FLAGS --stdlib=libc++) # __workweek fix. -#build_hip_executable_libcpp (hipPointerAttrib hipPointerAttrib.cpp) -build_hip_executable (hipHostAlloc hipHostAlloc.cpp) -build_hip_executable (hipHostGetFlags hipHostGetFlags.cpp) -build_hip_executable (hipHostRegister hipHostRegister.cpp) -build_hip_executable (hipRandomMemcpyAsync hipRandomMemcpyAsync.cpp) -build_hip_executable (hipFuncSetDeviceFlags hipFuncSetDeviceFlags.cpp) -build_hip_executable (hipFuncGetDevice hipFuncGetDevice.cpp) -build_hip_executable (hipFuncSetDevice hipFuncSetDevice.cpp) -build_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp) -build_hip_executable (hipPeerToPeer_simple hipPeerToPeer_simple.cpp) -build_hip_executable (hipTestMemcpyPin hipTestMemcpyPin.cpp) -build_hip_executable (hipDynamicShared hipDynamicShared.cpp) -build_hip_executable (hipLaunchParm hipLaunchParm.cpp) +#build_hip_executable_libcpp(hipPointerAttrib hipPointerAttrib.cpp) +build_hip_executable(hipHostAlloc hipHostAlloc.cpp) +build_hip_executable(hipHostGetFlags hipHostGetFlags.cpp) +build_hip_executable(hipHostRegister hipHostRegister.cpp) +build_hip_executable(hipRandomMemcpyAsync hipRandomMemcpyAsync.cpp) +build_hip_executable(hipFuncSetDeviceFlags hipFuncSetDeviceFlags.cpp) +build_hip_executable(hipFuncGetDevice hipFuncGetDevice.cpp) +build_hip_executable(hipFuncSetDevice hipFuncSetDevice.cpp) +build_hip_executable(hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp) +build_hip_executable(hipPeerToPeer_simple hipPeerToPeer_simple.cpp) +build_hip_executable(hipTestMemcpyPin hipTestMemcpyPin.cpp) +build_hip_executable(hipDynamicShared hipDynamicShared.cpp) +build_hip_executable(hipLaunchParm hipLaunchParm.cpp) -if (${HIP_PLATFORM} STREQUAL "hcc") - build_hip_executable (hipArray hipArray.cpp) +if(${HIP_PLATFORM} STREQUAL "hcc") + build_hip_executable(hipArray hipArray.cpp) endif() make_test(hipEventRecord --iterations 10) -make_test(hipEnvVarDriver " " ) +make_test(hipEnvVarDriver " ") make_test(hipLaunchParm " ") #TODO -reenable -#make_test(hipPointerAttrib " " ) +#make_test(hipPointerAttrib " ") make_test(hipHostAlloc " ") # BS- comment out since test appears broken - asks for device pointer but pointer was never allocated. #make_test(hipHostGetFlags " ") -make_test(hipHcc " " ) +make_test(hipHcc " ") make_test(hipHostRegister " ") make_test(hipRandomMemcpyAsync " ") make_test(hipFuncSetDeviceFlags " ") @@ -204,20 +244,20 @@ make_test(hipFuncGetDevice " ") make_test(hipFuncDeviceSynchronize " ") make_test(hipTestMemcpyPin " ") -if (${HIP_MULTI_GPU}) +if(${HIP_MULTI_GPU}) make_test(hipPeerToPeer_simple " ") # use current device for copy, this fails. make_test(hipPeerToPeer_simple --memcpyWithPeer) make_test(hipPeerToPeer_simple --mirrorPeers) # mirror mapping: test to ensure mirror doesn't destroy orig mapping. endif() -if (${HIP_PLATFORM} STREQUAL "hcc") +if(${HIP_PLATFORM} STREQUAL "hcc") make_test(hipArray " ") make_test(hipFuncSetDevice " ") make_test(hipDynamicShared " ") endif() -make_hipify_test(specialFunc.cu ) +make_hipify_test(specialFunc.cu) # Add subdirs here: @@ -225,3 +265,4 @@ add_subdirectory(context) add_subdirectory(deviceLib) add_subdirectory(runtimeApi) add_subdirectory(kernel) +# vim: ts=4:sw=4:expandtab:smartindent From 33fe29ae9fb8678f7c68d88d4cbf759fc365fead Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 13 Sep 2016 20:40:19 +0530 Subject: [PATCH 08/25] Directed tests: Remove support for HIP_BUILD_LOCAL=0 Change-Id: Ib73845ac8ea0507ca8e476995aa8c237b11d0ccf --- hipamd/tests/src/CMakeLists.txt | 165 ++++++++++++++------------------ 1 file changed, 70 insertions(+), 95 deletions(-) diff --git a/hipamd/tests/src/CMakeLists.txt b/hipamd/tests/src/CMakeLists.txt index 2563c31bab..b58b930f7f 100644 --- a/hipamd/tests/src/CMakeLists.txt +++ b/hipamd/tests/src/CMakeLists.txt @@ -19,102 +19,77 @@ if(NOT DEFINED HIP_MULTI_GPU) set(HIP_MULTI_GPU 0 CACHE BOOL "Run tests requiring more than one GPU") endif() -# Determine HIP_BUILD_LOCAL -if(NOT DEFINED HIP_BUILD_LOCAL) - if(NOT DEFINED ENV{HIP_BUILD_LOCAL}) - set(HIP_BUILD_LOCAL 1 CACHE BOOL "Build HIP in local folder") - else() - set(HIP_BUILD_LOCAL $ENV{HIP_BUILD_LOCAL} CACHE BOOL "Build HIP in local folder") - endif() +# We are going to use HIP source... +get_filename_component(HIP_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/../.. ABSOLUTE) + +# ...so we first need to determine the options to cascade to HIP build. +if(DEFINED HIP_PLATFORM) + set(ENV{HIP_PLATFORM} ${HIP_PLATFORM}) +endif() +if(DEFINED HCC_HOME) + get_filename_component(HCC_HOME ${HCC_HOME} ABSOLUTE) + set(ENV{HCC_HOME} ${HCC_HOME}) + set(ENV{HIP_DEVELOPER} 1) +endif() +if(DEFINED HIP_LIB_TYPE) + set(ENV{HIP_LIB_TYPE} ${HIP_LIB_TYPE}) endif() -# Determine HIP_PATH -if(HIP_BUILD_LOCAL) - # We are going to use HIP source... - get_filename_component(HIP_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/../.. ABSOLUTE) +# Purge previous HIP installation... +execute_process( + COMMAND "${CMAKE_COMMAND}" -E remove_directory hip + OUTPUT_QUIET + ERROR_QUIET + ) +execute_process( + COMMAND "${CMAKE_COMMAND}" -E make_directory hip + OUTPUT_QUIET + ERROR_QUIET + ) +message(STATUS "Configuring HIP") - # ...so we first need to determine the options to cascade to HIP build. - if(DEFINED HIP_PLATFORM) - set(ENV{HIP_PLATFORM} ${HIP_PLATFORM}) - endif() - if(DEFINED HCC_HOME) - get_filename_component(HCC_HOME ${HCC_HOME} ABSOLUTE) - set(ENV{HCC_HOME} ${HCC_HOME}) - set(ENV{HIP_DEVELOPER} 1) - endif() - if(DEFINED HIP_LIB_TYPE) - set(ENV{HIP_LIB_TYPE} ${HIP_LIB_TYPE}) - endif() +# ...and now build HIP locally. +execute_process( + COMMAND "${CMAKE_COMMAND}" -DCMAKE_INSTALL_PREFIX=${CMAKE_CURRENT_BINARY_DIR}/hip/localbuild ${HIP_SRC_PATH} + WORKING_DIRECTORY hip + RESULT_VARIABLE hip_build_result + OUTPUT_VARIABLE hip_build_log + OUTPUT_STRIP_TRAILING_WHITESPACE + ERROR_QUIET + ) +message("${Magenta}${hip_build_log}${ColorReset}") - # Purge previous HIP installation... - execute_process( - COMMAND "${CMAKE_COMMAND}" -E remove_directory hip - OUTPUT_QUIET - ERROR_QUIET - ) - execute_process( - COMMAND "${CMAKE_COMMAND}" -E make_directory hip - OUTPUT_QUIET - ERROR_QUIET - ) - message(STATUS "Configuring HIP") - - # ...and now build HIP locally. - execute_process( - COMMAND "${CMAKE_COMMAND}" -DCMAKE_INSTALL_PREFIX=${CMAKE_CURRENT_BINARY_DIR}/hip/localbuild ${HIP_SRC_PATH} - WORKING_DIRECTORY hip - RESULT_VARIABLE hip_build_result - OUTPUT_VARIABLE hip_build_log - OUTPUT_STRIP_TRAILING_WHITESPACE - ERROR_QUIET - ) - message("${Magenta}${hip_build_log}${ColorReset}") - - if(hip_build_result) - message(FATAL_ERROR "Error configuring HIP") - else() - message(STATUS "Configuring HIP - done") - message(STATUS "Building HIP") - endif() - execute_process( - COMMAND "${CMAKE_COMMAND}" --build . --target install - WORKING_DIRECTORY hip - RESULT_VARIABLE hip_build_result - OUTPUT_VARIABLE hip_build_log - OUTPUT_STRIP_TRAILING_WHITESPACE - ERROR_QUIET - ) - - # Show HIP build errors if any. - if(hip_build_result) - message("${Red}${hip_build_log}${ColorReset}") - message(FATAL_ERROR "Error building HIP") - else() - # Building HIP is successful. Point HIP_PATH to this location. - message(STATUS "Building HIP - done") - get_filename_component(HIP_PATH ${CMAKE_CURRENT_BINARY_DIR}/hip/localbuild ABSOLUTE) - endif() - - # Add a target to rebuild HIP if HIP source changes. - add_custom_target( - hip ALL - COMMAND "${CMAKE_COMMAND}" --build . --target install - WORKING_DIRECTORY hip - ) +if(hip_build_result) + message(FATAL_ERROR "Error configuring HIP") else() - if(NOT DEFINED HIP_PATH) - if(NOT DEFINED ENV{HIP_PATH}) - # We set to to default HIP installation path. - set(HIP_PATH /opt/rocm/hip CACHE PATH "Path to installed HIP") - else() - # We are using HIP_PATH from env. - set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to installed HIP") - endif() - endif() - # Create a fake target since we are using an installed HIP. - add_custom_target(hip ALL) + message(STATUS "Configuring HIP - done") + message(STATUS "Building HIP") endif() -MESSAGE("HIP_PATH=" ${HIP_PATH}) +execute_process( + COMMAND "${CMAKE_COMMAND}" --build . --target install + WORKING_DIRECTORY hip + RESULT_VARIABLE hip_build_result + OUTPUT_VARIABLE hip_build_log + OUTPUT_STRIP_TRAILING_WHITESPACE + ERROR_QUIET + ) + +# Show HIP build errors if any. +if(hip_build_result) + message("${Red}${hip_build_log}${ColorReset}") + message(FATAL_ERROR "Error building HIP") +else() + # Building HIP is successful. Point HIP_PATH to this location. + message(STATUS "Building HIP - done") + get_filename_component(HIP_PATH ${CMAKE_CURRENT_BINARY_DIR}/hip/localbuild ABSOLUTE) +endif() + +# Add a target to rebuild HIP if HIP source changes. +add_custom_target( + hip ALL + COMMAND "${CMAKE_COMMAND}" --build . --target install + WORKING_DIRECTORY hip + ) # Determine HIP_PLATFORM execute_process(COMMAND ${HIP_PATH}/bin/hipconfig --platform OUTPUT_VARIABLE HIP_PLATFORM) @@ -156,10 +131,10 @@ endmacro() function(make_named_test exe testname) add_test(NAME ${testname} COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN} - ) + ) set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION "PASSED" - ) + ) endfunction() macro(make_test exe) @@ -176,7 +151,7 @@ macro(make_hipify_test sourceFile) add_test(NAME ${testname} COMMAND ${HIP_PATH}/bin/hipify ${PROJECT_SOURCE_DIR}/${sourceFile} ${ARGN} - ) + ) endmacro() @@ -185,10 +160,10 @@ macro(make_test_matches exe match_string) set(testname ${exe}${smush_args}.tst) add_test(NAME ${testname} COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN} - ) + ) set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION ${match_string} - ) + ) endmacro() macro(build_hip_executable_sm35 exe cpp) From d3aff68f3705969b836b4decb7607a52dfccc9cd Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 14 Sep 2016 12:53:08 +0530 Subject: [PATCH 09/25] Partial update of memory APIs for doxygen documentation Change-Id: I7865f4591c7f64b2a21702c35f6233b845e7c403 --- hipamd/include/hcc_detail/hip_runtime_api.h | 78 ++++++++++++++++++++- 1 file changed, 77 insertions(+), 1 deletion(-) diff --git a/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index 221dfbf768..6c42d39041 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hcc_detail/hip_runtime_api.h @@ -891,6 +891,10 @@ hipError_t hipFree(void* ptr); /** * @brief Free memory allocated by the hcc hip host memory allocation API. [Deprecated.] * + * @param[in] ptr Pointer to memory to be freed + * @return #hipSuccess, + * #hipErrorInvalidValue (if pointer is invalid, including device pointers allocated with hipMalloc) + * @see hipHostFree */ hipError_t hipFreeHost(void* ptr) __attribute__((deprecated("use hipHostFree instead"))) ; @@ -928,20 +932,88 @@ hipError_t hipHostFree(void* ptr); * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * @param[in] copyType Memory copy type - * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknowni + * + * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer */ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind); +/** + * @brief Copy data from Host to Device + * + * @param[out] dst Data being copy to + * @param[in] src Data being copy from + * @param[in] sizeBytes Data size in bytes + * + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * + * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer + */ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes); +/** + * @brief Copy data from Device to Host + * + * @param[out] dst Data being copy to + * @param[in] src Data being copy from + * @param[in] sizeBytes Data size in bytes + * + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * + * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer + */ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes); +/** + * @brief Copy data from Device to Device + * + * @param[out] dst Data being copy to + * @param[in] src Data being copy from + * @param[in] sizeBytes Data size in bytes + * + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * + * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer + */ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes); +/** + * @brief Copy data from Host to Device asynchronously + * + * @param[out] dst Data being copy to + * @param[in] src Data being copy from + * @param[in] sizeBytes Data size in bytes + * + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * + * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer + */ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream); +/** + * @brief Copy data from Device to Host asynchronously + * + * @param[out] dst Data being copy to + * @param[in] src Data being copy from + * @param[in] sizeBytes Data size in bytes + * + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * + * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer + */ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream); +/** + * @brief Copy data from Device to Device asynchronously + * + * @param[out] dst Data being copy to + * @param[in] src Data being copy from + * @param[in] sizeBytes Data size in bytes + * + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * + * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer + */ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream); @@ -958,6 +1030,8 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz * @param[in] offset - Offset from start of symbol in bytes * @param[in] kind - Type of transfer * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown + * + * @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpy2DFromArray, hipMemcpyArrayToArray, hipMemcpy2DArrayToArray, hipMemcpyFromSymbol, hipMemcpyAsync, hipMemcpy2DAsync, hipMemcpyToArrayAsync, hipMemcpy2DToArrayAsync, hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, hipMemcpyFromSymbolAsync */ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind); @@ -980,6 +1054,8 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz * @param[in] sizeBytes Data size in bytes * @param[in] accelerator_view Accelerator view which the copy is being enqueued * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown + * + * @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpy2DFromArray, hipMemcpyArrayToArray, hipMemcpy2DArrayToArray, hipMemcpyToSymbol, hipMemcpyFromSymbol, hipMemcpy2DAsync, hipMemcpyToArrayAsync, hipMemcpy2DToArrayAsync, hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, hipMemcpyFromSymbolAsync */ #if __cplusplus hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream=0); From 4b1c7eed8ad4defbd2552e790ee746fb820df9aa Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 15 Sep 2016 17:02:37 +0530 Subject: [PATCH 10/25] Fixed broken HIP build on NVCC path Change-Id: Iae9d7ba0c789475681b382c8af18bb9bd6156a52 --- hipamd/include/nvcc_detail/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/include/nvcc_detail/hip_runtime_api.h b/hipamd/include/nvcc_detail/hip_runtime_api.h index 1ad1af2267..21ef239894 100644 --- a/hipamd/include/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/nvcc_detail/hip_runtime_api.h @@ -241,13 +241,13 @@ inline static hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, inline static hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t size, hipStream_t stream) { - return hipCUResultTohipError(cuMemcpyDtoH(dst, src, size, stream)); + return hipCUResultTohipError(cuMemcpyDtoH(dst, src, size)); } inline static hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t size, hipStream_t stream) { - return hipCUResultTohipError(cuMemcpyDtoD(dst, src, size, stream)); + return hipCUResultTohipError(cuMemcpyDtoD(dst, src, size)); } inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind) { From 2c12394d52080bc254a9083c792fab17fac90fe1 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 15 Sep 2016 10:50:26 -0500 Subject: [PATCH 11/25] added new error reporting case Change-Id: I5f0a37dbe396412f5602d04df19d538e451c2696 --- hipamd/src/hip_hcc.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index e2fb952a5b..a1d6e2537f 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -1559,6 +1559,7 @@ const char *ihipErrorString(hipError_t hip_error) { switch (hip_error) { case hipSuccess : return "hipSuccess"; + case hipErrorOutofMemory : return "hipOutofMemory"; case hipErrorMemoryAllocation : return "hipErrorMemoryAllocation"; case hipErrorLaunchOutOfResources : return "hipErrorLaunchOutOfResources"; case hipErrorInvalidValue : return "hipErrorInvalidValue"; From 6cb5a13a2b8f09d069106b759ae3d7d9068367ce Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 15 Sep 2016 11:28:18 -0500 Subject: [PATCH 12/25] added more error codes to hipErrorGetString Change-Id: I80c675905d94813502040fd0caa07985fa8c7dcc --- hipamd/include/hip_runtime_api.h | 46 +++++++++++++------------- hipamd/src/hip_hcc.cpp | 55 ++++++++++++++++++++++++++++---- 2 files changed, 72 insertions(+), 29 deletions(-) diff --git a/hipamd/include/hip_runtime_api.h b/hipamd/include/hip_runtime_api.h index 4cb5bb9bcd..660385972a 100644 --- a/hipamd/include/hip_runtime_api.h +++ b/hipamd/include/hip_runtime_api.h @@ -185,30 +185,30 @@ typedef enum hipError_t { hipErrorIllegalAddress = 700, // Runtime Error Codes start here. - hipErrorMissingConfiguration = 1, - hipErrorMemoryAllocation = 2, ///< Memory allocation error. - hipErrorInitializationError = 3, ///< TODO comment from hipErrorInitializationError - hipErrorLaunchFailure = 4, - hipErrorPriorLaunchFailure = 5, - hipErrorLaunchTimeOut = 6, - hipErrorLaunchOutOfResources = 7, ///< Out of resources error. - hipErrorInvalidDeviceFunction = 8, - hipErrorInvalidConfiguration = 9, - hipErrorInvalidDevice = 10, ///< DeviceID must be in range 0...#compute-devices. - hipErrorInvalidValue = 11, ///< One or more of the parameters passed to the API call is NULL or not in an acceptable range. - hipErrorInvalidDevicePointer = 17, ///< Invalid Device Pointer - hipErrorInvalidMemcpyDirection = 21, ///< Invalid memory copy direction - hipErrorUnknown = 30, ///< Unknown error. - hipErrorInvalidResourceHandle = 33, ///< Resource handle (hipEvent_t or hipStream_t) invalid. - hipErrorNotReady = 34, ///< Indicates that asynchronous operations enqueued earlier are not ready. This is not actually an error, but is used to distinguish from hipSuccess (which indicates completion). APIs that return this error include hipEventQuery and hipStreamQuery. - hipErrorNoDevice = 38, ///< Call to hipGetDeviceCount returned 0 devices - hipErrorPeerAccessAlreadyEnabled = 50, ///< Peer access was already enabled from the current device. + hipErrorMissingConfiguration = 1001, + hipErrorMemoryAllocation = 1002, ///< Memory allocation error. + hipErrorInitializationError = 1003, ///< TODO comment from hipErrorInitializationError + hipErrorLaunchFailure = 1004, + hipErrorPriorLaunchFailure = 1005, + hipErrorLaunchTimeOut = 1006, + hipErrorLaunchOutOfResources = 1007, ///< Out of resources error. + hipErrorInvalidDeviceFunction = 1008, + hipErrorInvalidConfiguration = 1009, + hipErrorInvalidDevice = 1010, ///< DeviceID must be in range 0...#compute-devices. + hipErrorInvalidValue = 1011, ///< One or more of the parameters passed to the API call is NULL or not in an acceptable range. + hipErrorInvalidDevicePointer = 1017, ///< Invalid Device Pointer + hipErrorInvalidMemcpyDirection = 1021, ///< Invalid memory copy direction + hipErrorUnknown = 1030, ///< Unknown error. + hipErrorInvalidResourceHandle = 1033, ///< Resource handle (hipEvent_t or hipStream_t) invalid. + hipErrorNotReady = 1034, ///< Indicates that asynchronous operations enqueued earlier are not ready. This is not actually an error, but is used to distinguish from hipSuccess (which indicates completion). APIs that return this error include hipEventQuery and hipStreamQuery. + hipErrorNoDevice = 1038, ///< Call to hipGetDeviceCount returned 0 devices + hipErrorPeerAccessAlreadyEnabled = 1050, ///< Peer access was already enabled from the current device. - hipErrorPeerAccessNotEnabled = 51, ///< Peer access was never enabled from the current device. - hipErrorRuntimeMemory, ///< HSA runtime memory call returned error. Typically not seen in production systems. - hipErrorRuntimeOther, ///< HSA runtime call other than memory returned error. Typically not seen in production systems. - hipErrorHostMemoryAlreadyRegistered = 61, ///< Produced when trying to lock a page-locked memory. - hipErrorHostMemoryNotRegistered = 62, ///< Produced when trying to unlock a non-page-locked memory. + hipErrorPeerAccessNotEnabled = 1051, ///< Peer access was never enabled from the current device. + hipErrorRuntimeMemory = 1052, ///< HSA runtime memory call returned error. Typically not seen in production systems. + hipErrorRuntimeOther = 1053, ///< HSA runtime call other than memory returned error. Typically not seen in production systems. + hipErrorHostMemoryAlreadyRegistered = 1061, ///< Produced when trying to lock a page-locked memory. + hipErrorHostMemoryNotRegistered = 1062, ///< Produced when trying to unlock a non-page-locked memory. hipErrorTbd ///< Marker that more error codes are needed. } hipError_t; diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index a1d6e2537f..0f0b03d115 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -1559,21 +1559,64 @@ const char *ihipErrorString(hipError_t hip_error) { switch (hip_error) { case hipSuccess : return "hipSuccess"; - case hipErrorOutofMemory : return "hipOutofMemory"; + case hipErrorOutOfMemory : return "hipErrorOutOfMemory"; + case hipErrorNotInitialized : return "hipErrorNotInitialized"; + case hipErrorDeinitialized : return "hipErrorDeinitialized"; + case hipErrorProfilerDisabled : return "hipErrorProfilerDisabled"; + case hipErrorProfilerNotInitialized : return "hipErrorProfilerNotInitialized"; + case hipErrorProfilerAlreadyStarted : return "hipErrorProfilerAlreadyStarted"; + case hipErrorProfilerAlreadyStopped : return "hipErrorProfilerAlreadyStopped"; + case hipErrorInvalidImage : return "hipErrorInvalidImage"; + case hipErrorInvalidContext : return "hipErrorInvalidContext"; + case hipErrorContextAlreadyCurrent : return "hipErrorContextAlreadyCurrent"; + case hipErrorMapFailed : return "hipErrorMapFailed"; + case hipErrorUnmapFailed : return "hipErrorUnmapFailed"; + case hipErrorArrayIsMapped : return "hipErrorArrayIsMapped"; + case hipErrorAlreadyMapped : return "hipErrorAlreadyMapped"; + case hipErrorNoBinaryForGpu : return "hipErrorNoBinaryForGpu"; + case hipErrorAlreadyAcquired : return "hipErrorAlreadyAcquired"; + case hipErrorNotMapped : return "hipErrorNotMapped"; + case hipErrorNotMappedAsArray : return "hipErrorNotMappedAsArray"; + case hipErrorNotMappedAsPointer : return "hipErrorNotMappedAsPointer"; + case hipErrorECCNotCorrectable : return "hipErrorECCNotCorrectable"; + case hipErrorUnsupportedLimit : return "hipErrorUnsupportedLimit"; + case hipErrorContextAlreadyInUse : return "hipErrorContextAlreadyInUse"; + case hipErrorPeerAccessUnsupported : return "hipErrorPeerAccessUnsupported"; + case hipErrorInvalidKernelFile : return "hipErrorInvalidKernelFile"; + case hipErrorInvalidGraphicsContext : return "hipErrorInvalidGraphicsContext"; + case hipErrorInvalidSource : return "hipErrorInvalidSource"; + case hipErrorFileNotFound : return "hipErrorFileNotFound"; + case hipErrorSharedObjectSymbolNotFound : return "hipErrorSharedObjectSymbolNotFound"; + case hipErrorSharedObjectInitFailed : return "hipErrorSharedObjectInitFailed"; + case hipErrorOperatingSystem : return "hipErrorOperatingSystem"; + case hipErrorInvalidHandle : return "hipErrorInvalidHandle"; + case hipErrorNotFound : return "hipErrorNotFound"; + case hipErrorIllegalAddress : return "hipErrorIllegalAddress"; + + case hipErrorMissingConfiguration : return "hipErrorMissingConfiguration"; case hipErrorMemoryAllocation : return "hipErrorMemoryAllocation"; + case hipErrorInitializationError : return "hipErrorInitializationError"; + case hipErrorLaunchFailure : return "hipErrorLaunchFailure"; + case hipErrorPriorLaunchFailure : return "hipErrorPriorLaunchFailure"; + case hipErrorLaunchTimeOut : return "hipErrorLaunchTimeOut"; case hipErrorLaunchOutOfResources : return "hipErrorLaunchOutOfResources"; - case hipErrorInvalidValue : return "hipErrorInvalidValue"; - case hipErrorInvalidResourceHandle : return "hipErrorInvalidResourceHandle"; + case hipErrorInvalidDeviceFunction : return "hipErrorInvalidDeviceFunction"; + case hipErrorInvalidConfiguration : return "hipErrorInvalidConfiguration"; case hipErrorInvalidDevice : return "hipErrorInvalidDevice"; + case hipErrorInvalidValue : return "hipErrorInvalidValue"; + case hipErrorInvalidDevicePointer : return "hipErrorInvalidDevicePointer"; case hipErrorInvalidMemcpyDirection : return "hipErrorInvalidMemcpyDirection"; - case hipErrorNoDevice : return "hipErrorNoDevice"; + case hipErrorUnknown : return "hipErrorUnknown"; + case hipErrorInvalidResourceHandle : return "hipErrorInvalidResourceHandle"; case hipErrorNotReady : return "hipErrorNotReady"; - case hipErrorPeerAccessNotEnabled : return "hipErrorPeerAccessNotEnabled"; + case hipErrorNoDevice : return "hipErrorNoDevice"; case hipErrorPeerAccessAlreadyEnabled : return "hipErrorPeerAccessAlreadyEnabled"; + case hipErrorPeerAccessNotEnabled : return "hipErrorPeerAccessNotEnabled"; case hipErrorRuntimeMemory : return "hipErrorRuntimeMemory"; case hipErrorRuntimeOther : return "hipErrorRuntimeOther"; - case hipErrorUnknown : return "hipErrorUnknown"; + case hipErrorHostMemoryAlreadyRegistered : return "hipErrorHostMemoryAlreadyRegistered"; + case hipErrorHostMemoryNotRegistered : return "hipErrorHostMemoryNotRegistered"; case hipErrorTbd : return "hipErrorTbd"; default : return "hipErrorUnknown"; }; From 8faf2ed7e3dff0b55a1a0f2f3233c08290792e93 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 15 Sep 2016 11:40:17 -0500 Subject: [PATCH 13/25] added vimrc for current project 1. Added vimrc config file for HIP 2. Corrected square sample indent Change-Id: I3e1d92403571148fe6825db6ad63ad925ae69519 --- hipamd/.gitignore | 1 - hipamd/.vimrc | 4 ++ .../samples/0_Intro/square/square.hipref.cpp | 66 +++++++++---------- 3 files changed, 36 insertions(+), 35 deletions(-) create mode 100644 hipamd/.vimrc diff --git a/hipamd/.gitignore b/hipamd/.gitignore index 450d3bc41b..e298b686bf 100644 --- a/hipamd/.gitignore +++ b/hipamd/.gitignore @@ -1,4 +1,3 @@ -.* *.o *.exe *.swp diff --git a/hipamd/.vimrc b/hipamd/.vimrc new file mode 100644 index 0000000000..ed64acd347 --- /dev/null +++ b/hipamd/.vimrc @@ -0,0 +1,4 @@ +:set tabstop=4 +:set shiftwidth=4 +:set expandtab +:set smartindent diff --git a/hipamd/samples/0_Intro/square/square.hipref.cpp b/hipamd/samples/0_Intro/square/square.hipref.cpp index aa14077738..7ca3a7500d 100644 --- a/hipamd/samples/0_Intro/square/square.hipref.cpp +++ b/hipamd/samples/0_Intro/square/square.hipref.cpp @@ -31,7 +31,6 @@ THE SOFTWARE. }\ } - /* * Square each element in the array A and write to array C. */ @@ -43,55 +42,54 @@ vector_square(hipLaunchParm lp, T *C_d, const T *A_d, size_t N) size_t stride = hipBlockDim_x * hipGridDim_x ; for (size_t i=offset; i Date: Fri, 16 Sep 2016 10:16:34 +0530 Subject: [PATCH 14/25] FindHIP: determine HIP_PLATFORM at build time instead of compile time Change-Id: Icb99d23f2874c4c15771641d7e3e91bb5c918f27 --- hipamd/cmake/FindHIP.cmake | 2 +- hipamd/cmake/FindHIP/run_hipcc.cmake | 3 ++- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/hipamd/cmake/FindHIP.cmake b/hipamd/cmake/FindHIP.cmake index 2541cf6ec6..1666673beb 100644 --- a/hipamd/cmake/FindHIP.cmake +++ b/hipamd/cmake/FindHIP.cmake @@ -416,7 +416,7 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files) # Create up the comment string file(RELATIVE_PATH generated_file_relative_path "${CMAKE_BINARY_DIR}" "${generated_file}") - set(hip_build_comment_string "Building HIPCC (using ${HIP_PLATFORM}) object ${generated_file_relative_path}") + set(hip_build_comment_string "Building HIPCC object ${generated_file_relative_path}") # Build the generated file and dependency file add_custom_command( diff --git a/hipamd/cmake/FindHIP/run_hipcc.cmake b/hipamd/cmake/FindHIP/run_hipcc.cmake index 52ad57f532..6027cb9b0d 100644 --- a/hipamd/cmake/FindHIP/run_hipcc.cmake +++ b/hipamd/cmake/FindHIP/run_hipcc.cmake @@ -22,8 +22,8 @@ endif() # Set these up as variables to make reading the generated file easier set(HIP_HIPCC_EXECUTABLE "@HIP_HIPCC_EXECUTABLE@") # path +set(HIP_HIPCONFIG_EXECUTABLE "@HIP_HIPCONFIG_EXECUTABLE@") #path set(HIP_HOST_COMPILER "@HIP_HOST_COMPILER@") # path -set(HIP_PLATFORM "@HIP_PLATFORM@") #string set(CMAKE_COMMAND "@CMAKE_COMMAND@") # path set(HIP_run_make2cmake "@HIP_run_make2cmake@") # path @@ -38,6 +38,7 @@ set(source_file "@source_file@") # path set(host_flag "@host_flag@") # bool # Determine compiler and compiler flags +execute_process(COMMAND ${HIP_HIPCONFIG_EXECUTABLE} --platform OUTPUT_VARIABLE HIP_PLATFORM OUTPUT_STRIP_TRAILING_WHITESPACE) if(NOT host_flag) set(__CC ${HIP_HIPCC_EXECUTABLE}) if(HIP_PLATFORM STREQUAL "hcc") From 730268e08794156e6e861b2eedb9648eb741bacf Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 16 Sep 2016 10:17:27 +0530 Subject: [PATCH 15/25] FindHIP: Initial support for HIP_ADD_LIBRARY Change-Id: Ic8b93c2a9daff2dddb3a94b62981d3a1f03afd5c --- hipamd/cmake/FindHIP.cmake | 29 +++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) diff --git a/hipamd/cmake/FindHIP.cmake b/hipamd/cmake/FindHIP.cmake index 1666673beb..9626e7629f 100644 --- a/hipamd/cmake/FindHIP.cmake +++ b/hipamd/cmake/FindHIP.cmake @@ -342,6 +342,24 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files) HIP_PARSE_HIPCC_OPTIONS(HIP_HCC_FLAGS ${_hcc_options}) HIP_PARSE_HIPCC_OPTIONS(HIP_NVCC_FLAGS ${_nvcc_options}) + # Check if we are building shared library. + set(_hip_build_shared_libs FALSE) + list(FIND _hip_cmake_options SHARED _hip_found_SHARED) + list(FIND _hip_cmake_options MODULE _hip_found_MODULE) + if(_hip_found_SHARED GREATER -1 OR _hip_found_MODULE GREATER -1) + set(_hip_build_shared_libs TRUE) + endif() + list(FIND _hip_cmake_options STATIC _hip_found_STATIC) + if(_hip_found_STATIC GREATER -1) + set(_hip_build_shared_libs FALSE) + endif() + + # If we are building a shared library, add extra flags to HIP_HIPCC_FLAGS + if(_hip_build_shared_libs) + list(APPEND HIP_HCC_FLAGS "-fPIC") + list(APPEND HIP_NVCC_FLAGS "--shared -Xcompiler '-fPIC'") + endif() + # Set host compiler set(HIP_HOST_COMPILER "${CMAKE_${HIP_C_OR_CXX}_COMPILER}") @@ -458,4 +476,15 @@ macro(HIP_ADD_EXECUTABLE hip_target) set_target_properties(${hip_target} PROPERTIES LINKER_LANGUAGE HIP) endmacro() +############################################################################### +# HIP_ADD_LIBRARY +############################################################################### +macro(HIP_ADD_LIBRARY hip_target) + # Separate the sources from the options + HIP_GET_SOURCES_AND_OPTIONS(_sources _cmake_options _hipcc_options _hcc_options _nvcc_options ${ARGN}) + HIP_PREPARE_TARGET_COMMANDS(${hip_target} OBJ _generated_files ${_sources} ${_cmake_options} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options}) + add_library(${hip_target} ${_cmake_options} ${_generated_files} ${_sources}) + set_target_properties(${hip_target} PROPERTIES LINKER_LANGUAGE ${HIP_C_OR_CXX}) +endmacro() + # vim: ts=4:sw=4:expandtab:smartindent From 2c1bd4bd108f331ec7fd7b996b3d412d2852682d Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 16 Sep 2016 11:15:03 +0530 Subject: [PATCH 16/25] HIP/NVCC path for hipChooseDevice Change-Id: I64690d7df1f5493e067c7ff5e3b39a10fcf323bb --- hipamd/include/nvcc_detail/hip_runtime_api.h | 23 ++++++++++++++++++++ 1 file changed, 23 insertions(+) diff --git a/hipamd/include/nvcc_detail/hip_runtime_api.h b/hipamd/include/nvcc_detail/hip_runtime_api.h index 21ef239894..c90c7cbf29 100644 --- a/hipamd/include/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/nvcc_detail/hip_runtime_api.h @@ -214,6 +214,29 @@ inline static hipError_t hipSetDevice(int device) { return hipCUDAErrorTohipError(cudaSetDevice(device)); } +inline static hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop ) +{ + cudaDeviceProp cdprop; + memset(&cdprop,0x0,sizeof(cudaDeviceProp)); + cdprop.major= prop->major; + cdprop.minor = prop->minor; + cdprop.totalGlobalMem = prop->totalGlobalMem ; + cdprop.sharedMemPerBlock = prop->sharedMemPerBlock; + cdprop.regsPerBlock = prop->regsPerBlock; + cdprop.warpSize = prop->warpSize ; + cdprop.maxThreadsPerBlock = prop->maxThreadsPerBlock ; + cdprop.clockRate = prop->clockRate; + cdprop.totalConstMem = prop->totalConstMem ; + cdprop.multiProcessorCount = prop->multiProcessorCount ; + cdprop.l2CacheSize = prop->l2CacheSize ; + cdprop.maxThreadsPerMultiProcessor = prop->maxThreadsPerMultiProcessor ; + cdprop.computeMode = prop->computeMode ; + cdprop.canMapHostMemory = prop->canMapHostMemory; + cdprop.memoryClockRate = prop->memoryClockRate; + cdprop.memoryBusWidth = prop->memoryBusWidth; + return hipCUDAErrorTohipError(cudaChooseDevice(device,&cdprop)); +} + inline static hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t size) { From 6139d5f587092c6c73bfe02d9bf8ed6b71bccb03 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 16 Sep 2016 11:18:34 +0530 Subject: [PATCH 17/25] Added further properties in hipChooseDevice Change-Id: I8f7f5800e8a95873318a534f09e75e037ad31b06 --- hipamd/src/hip_device.cpp | 43 ++++++++++++++++++++++++++++++++++++++- 1 file changed, 42 insertions(+), 1 deletion(-) diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index fe7ad6ecd7..6ea4658a4f 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -305,7 +305,48 @@ hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop ) } } } - + if(prop->totalGlobalMem != 0) { + inPropCount++; + if(tempProp.totalGlobalMem >= prop->totalGlobalMem) { + matchedPropCount++; + } + } + if(prop->sharedMemPerBlock != 0) { + inPropCount++; + if(tempProp.sharedMemPerBlock >= prop->sharedMemPerBlock) { + matchedPropCount++; + } + } + if(prop->maxThreadsPerBlock != 0) { + inPropCount++; + if(tempProp.maxThreadsPerBlock >= prop->maxThreadsPerBlock ) { + matchedPropCount++; + } + } + if(prop->totalConstMem != 0) { + inPropCount++; + if(tempProp.totalConstMem >= prop->totalConstMem ) { + matchedPropCount++; + } + } + if(prop->multiProcessorCount != 0) { + inPropCount++; + if(tempProp.multiProcessorCount >= prop->multiProcessorCount ) { + matchedPropCount++; + } + } + if(prop->maxThreadsPerMultiProcessor != 0) { + inPropCount++; + if(tempProp.maxThreadsPerMultiProcessor >= prop->maxThreadsPerMultiProcessor ) { + matchedPropCount++; + } + } + if(prop->memoryClockRate != 0) { + inPropCount++; + if(tempProp.memoryClockRate >= prop->memoryClockRate ) { + matchedPropCount++; + } + } if(inPropCount == matchedPropCount) { *device = i; } From 4ceb73baa1104018c19a60d42cf07a518544bcd9 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 16 Sep 2016 12:39:33 +0530 Subject: [PATCH 18/25] Ignore .* files again Change-Id: Idf430bea898537df2cdef11796b12e52fa4cd9c6 --- hipamd/.gitignore | 1 + 1 file changed, 1 insertion(+) diff --git a/hipamd/.gitignore b/hipamd/.gitignore index e298b686bf..450d3bc41b 100644 --- a/hipamd/.gitignore +++ b/hipamd/.gitignore @@ -1,3 +1,4 @@ +.* *.o *.exe *.swp From cb3bff0a0dfad7096cc63a5a852fea9a69588a10 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 16 Sep 2016 12:59:41 +0530 Subject: [PATCH 19/25] Partial updates of error codes for Memory APIs Change-Id: If924867cc62f56dc9eaf81f5499ba66f417c393a --- hipamd/include/hcc_detail/hip_runtime_api.h | 5 +- hipamd/src/hip_memory.cpp | 60 --------------------- 2 files changed, 4 insertions(+), 61 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index 6c42d39041..d8b7fc4b55 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hcc_detail/hip_runtime_api.h @@ -1042,6 +1042,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz * @warning If host or dest are not pinned, the memory copy will be performed synchronously. For best performance, use hipHostMalloc to * allocate host memory that is transferred asynchronously. * + * @warning on HCC hipMemcpyAsync does not support overlapped H2D and D2H copies. * For hipMemcpy, the copy is always performed by the device associated with the specified stream. * * For multi-gpu or peer-to-peer configurations, it is recommended to use a stream which is a attached to the device where the src data is physically located. @@ -1101,7 +1102,9 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st * @brief Query memory info. * Return snapshot of free memory, and total allocatable memory on the device. * - * Returns in *free a snapshot of the current free memory o + * Returns in *free a snapshot of the current free memory. + * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue (if free != NULL due to bugs) + * @warning On HCC, the free memory only accounts for memory allocated by this process and may be optimistic. **/ hipError_t hipMemGetInfo (size_t * free, size_t * total) ; diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 19b4745b59..0478ff26d7 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -28,12 +28,6 @@ // Memory // // -// - -//--- -/** - * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice - */ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) { HIP_INIT_API(attributes, ptr); @@ -77,12 +71,6 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) return ihipLogStatus(e); } - -/** - * @returns #hipSuccess, - * @returns #hipErrorInvalidValue if flags are not 0 - * @returns #hipErrorMemoryAllocation if hostPointer is not a tracked allocation. - */ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsigned flags) { HIP_INIT_API(devicePointer, hostPointer, flags); @@ -107,13 +95,6 @@ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsi return ihipLogStatus(e); } - - - -//--- -/** - * @returns #hipSuccess #hipErrorMemoryAllocation - */ hipError_t hipMalloc(void** ptr, size_t sizeBytes) { HIP_INIT_API(ptr, sizeBytes); @@ -153,8 +134,6 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) return ihipLogStatus(hip_status); } - - hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_INIT_API(ptr, sizeBytes, flags); @@ -193,7 +172,6 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) return ihipLogStatus(hip_status); } - //--- // TODO - remove me, this is deprecated. hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) @@ -201,7 +179,6 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) return hipHostMalloc(ptr, sizeBytes, flags); }; - //--- // TODO - remove me, this is deprecated. hipError_t hipMallocHost(void** ptr, size_t sizeBytes) @@ -253,7 +230,6 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height return ihipLogStatus(hip_status); } - hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { hipChannelFormatDesc cd; @@ -262,7 +238,6 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel return cd; } - hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) { @@ -324,8 +299,6 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, return ihipLogStatus(hip_status); } - -//--- hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) { HIP_INIT_API(flagsPtr, hostPtr); @@ -350,8 +323,6 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) return ihipLogStatus(hip_status); } - -//--- hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) { HIP_INIT_API(hostPtr, sizeBytes, flags); @@ -395,7 +366,6 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) return ihipLogStatus(hip_status); } -//--- hipError_t hipHostUnregister(void *hostPtr) { HIP_INIT_API(hostPtr); @@ -413,8 +383,6 @@ hipError_t hipHostUnregister(void *hostPtr) return ihipLogStatus(hip_status); } - -//--- hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind) { HIP_INIT_API(symbolName, src, count, offset, kind); @@ -478,7 +446,6 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) return ihipLogStatus(e); } - hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { HIP_INIT_API(dst, src, sizeBytes); @@ -542,15 +509,6 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) return ihipLogStatus(e); } - - -/** - * @result #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidMemcpyDirection, - * @result #hipErrorInvalidValue : If dst==NULL or src==NULL, or other bad argument. - * @warning on HCC hipMemcpyAsync does not support overlapped H2D and D2H copies. - * @warning on HCC hipMemcpyAsync requires that any host pointers are pinned (ie via the hipMallocHost call). - */ -//--- hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_API(dst, src, sizeBytes, kind, stream); @@ -655,7 +613,6 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h return ihipLogStatus(e); } -// dpitch, spitch, and width in bytes hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { @@ -682,7 +639,6 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, return ihipLogStatus(e); } -// wOffset, width, and spitch in bytes hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { @@ -757,8 +713,6 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, return ihipLogStatus(e); } - - // TODO - make member function of stream? template hc::completion_future @@ -797,11 +751,7 @@ ihipMemsetKernel(hipStream_t stream, return cf; } - - // TODO-sync: function is async unless target is pinned host memory - then these are fully sync. -/** @return #hipErrorInvalidValue -*/ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream ) { HIP_INIT_API(dst, value, sizeBytes, stream); @@ -851,7 +801,6 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s return ihipLogStatus(e); }; - hipError_t hipMemset(void* dst, int value, size_t sizeBytes ) { hipStream_t stream = hipStreamNull; @@ -903,11 +852,6 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes ) return ihipLogStatus(e); } - -/* - * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue (if free != NULL due to bug)S - * @warning On HCC, the free memory only accounts for memory allocated by this process and may be optimistic. - */ hipError_t hipMemGetInfo (size_t *free, size_t *total) { HIP_INIT_API(free, total); @@ -936,8 +880,6 @@ hipError_t hipMemGetInfo (size_t *free, size_t *total) return ihipLogStatus(e); } - -//--- hipError_t hipFree(void* ptr) { HIP_INIT_API(ptr); @@ -965,7 +907,6 @@ hipError_t hipFree(void* ptr) return ihipLogStatus(hipStatus); } - hipError_t hipHostFree(void* ptr) { HIP_INIT_API(ptr); @@ -993,7 +934,6 @@ hipError_t hipHostFree(void* ptr) return ihipLogStatus(hipStatus); }; - // TODO - deprecated function. hipError_t hipFreeHost(void* ptr) { From 58a828ede57ffe7ccb1a5f9ae2a8bed2bf022aea Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 16 Sep 2016 14:14:45 +0530 Subject: [PATCH 20/25] Corrected wrong error code name in memory APIs Change-Id: I16a8ebf6d530f2e7d266a1a63a76ba7b9710c8c1 --- hipamd/include/hcc_detail/hip_runtime_api.h | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index d8b7fc4b55..be8b513caf 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hcc_detail/hip_runtime_api.h @@ -945,7 +945,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer */ @@ -958,7 +958,7 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes); * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer */ @@ -971,7 +971,7 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes); * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer */ @@ -984,7 +984,7 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer */ @@ -997,7 +997,7 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, h * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer */ @@ -1010,7 +1010,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValuei + * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, hipMemAllocPitch, hipMemcpy2D, hipMemcpy2DAsync, hipMemcpy2DUnaligned, hipMemcpyAtoA, hipMemcpyAtoD, hipMemcpyAtoH, hipMemcpyAtoHAsync, hipMemcpyDtoA, hipMemcpyDtoD, hipMemcpyDtoDAsync, hipMemcpyDtoH, hipMemcpyDtoHAsync, hipMemcpyHtoA, hipMemcpyHtoAAsync, hipMemcpyHtoDAsync, hipMemFree, hipMemFreeHost, hipMemGetAddressRange, hipMemGetInfo, hipMemHostAlloc, hipMemHostGetDevicePointer */ From 5b0a4eeaba38b71da8675e71138f781e3fa43450 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 16 Sep 2016 14:59:08 +0530 Subject: [PATCH 21/25] Updated error codes in Memory APIs Change-Id: I328090db929867fd3500fc471b69f86c6f595912 --- hipamd/include/hcc_detail/hip_runtime_api.h | 30 ++++++++++----------- 1 file changed, 15 insertions(+), 15 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index be8b513caf..db1faefa89 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hcc_detail/hip_runtime_api.h @@ -194,7 +194,7 @@ typedef enum hipMemcpyKind { * When this command is invoked, the host thread gets blocked until all the commands associated * with streams associated with the device. HIP does not support multiple blocking modes (yet!). * - * @returns hipSuccess + * @returns #hipSuccess * * @see hipSetDevice, hipDeviceReset */ @@ -208,7 +208,7 @@ hipError_t hipDeviceSynchronize(void); * Calling this function deletes all streams created, memory allocated, kernels running, events created. * Make sure that no other thread is using the device or streams, memory, kernels, events associated with the current device. * - * @returns hipSuccess + * @returns #hipSuccess * * @see hipDeviceSynchronize */ @@ -237,7 +237,7 @@ hipError_t hipDeviceReset(void) ; * Thread-pool implementations may inherit the default device of the previous thread. A good practice is to always call hipSetDevice * at the start of HIP coding sequency to establish a known standard device. * - * @returns hipSuccess, hipErrorInvalidDevice, hipErrorDeviceAlreadyInUse + * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorDeviceAlreadyInUse * * @see hipGetDevice, hipGetDeviceCount */ @@ -253,9 +253,9 @@ 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. * - * @see hipSetDevice, hipGetDevicesizeBytes + * @returns #hipSuccess * - * @returns hipSuccess + * @see hipSetDevice, hipGetDevicesizeBytes */ hipError_t hipGetDevice(int *deviceId); @@ -265,7 +265,7 @@ hipError_t hipGetDevice(int *deviceId); * * @param [output] count Returns number of compute-capable devices. * - * @returns hipSuccess, hipErrorNoDevice + * @returns #hipSuccess, #hipErrorNoDevice * * * Returns in @p *count the number of devices that have ability to run compute commands. If there are no such devices, then @ref hipGetDeviceCount will return #hipErrorNoDevice. @@ -280,7 +280,7 @@ hipError_t hipGetDeviceCount(int *count); * @param [in] attr attribute to query * @param [in] deviceId which device to query for information * - * @returns hipSuccess, hipErrorInvalidDevice, hipErrorInvalidValue + * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue */ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int deviceId); @@ -290,7 +290,7 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int deviceI * @param [out] prop written with device properties * @param [in] deviceId which device to query for information * - * @return hipSuccess, hipErrorInvalidDevice + * @return #hipSuccess, #hipErrorInvalidDevice * @bug HCC always returns 0 for maxThreadsPerMultiProcessor * @bug HCC always returns 0 for regsPerBlock * @bug HCC always returns 0 for l2CacheSize @@ -305,7 +305,7 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId); * * @param [in] cacheConfig * - * @returns hipSuccess, hipErrorInitializationError + * @returns #hipSuccess, #hipErrorInitializationError * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures. * */ @@ -317,7 +317,7 @@ hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig ); * * @param [in] cacheConfig * - * @returns hipSuccess, hipErrorInitializationError + * @returns #hipSuccess, #hipErrorInitializationError * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures. * */ @@ -329,7 +329,7 @@ hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig ); * * @param [in] config; * - * @returns hipSuccess, hipErrorInitializationError + * @returns #hipSuccess, #hipErrorInitializationError * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures. * */ @@ -340,7 +340,7 @@ hipError_t hipFuncSetCacheConfig ( hipFuncCache config ); * * @param [out] pConfig * - * @returns hipSuccess, hipErrorInvalidValue, hipErrorInitializationError + * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError * * Note: AMD devices and recent Nvidia GPUS do not support shared cache banking, and the hint is ignored on those architectures. * @@ -353,7 +353,7 @@ hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ); * * @param [in] config * - * @returns hipSuccess, hipErrorInvalidValue, hipErrorInitializationError + * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError * * Note: AMD devices and recent Nvidia GPUS do not support shared cache banking, and the hint is ignored on those architectures. * @@ -365,7 +365,7 @@ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ); * * @param [in] flags * - * @returns hipSuccess, hipErrorInvalidDevice, hipErrorSetOnActiveProcess + * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorSetOnActiveProcess * * Note: Only hipDeviceScheduleAuto and hipDeviceMapHost are supported * @@ -378,7 +378,7 @@ hipError_t hipSetDeviceFlags ( unsigned flags); * @param [out] device ID * @param [in] device properties pointer * - * @returns hipSuccess, hipErrorInvalidValue + * @returns #hipSuccess, #hipErrorInvalidValue */ hipError_t hipChooseDevice(int *device, hipDeviceProp_t* prop); From f51b5dda0c3f061372087fdfd648f5c10934735c Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 16 Sep 2016 15:04:52 +0530 Subject: [PATCH 22/25] Fix fence functions doxygen group Change-Id: I222eecb0eb0174b9f19e15b16c1ee06b7ba1e67e --- hipamd/include/hcc_detail/hip_runtime.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_runtime.h b/hipamd/include/hcc_detail/hip_runtime.h index 547df405a2..92406b4a75 100644 --- a/hipamd/include/hcc_detail/hip_runtime.h +++ b/hipamd/include/hcc_detail/hip_runtime.h @@ -500,7 +500,7 @@ __device__ float __dsqrt_rz(double x); /** *------------------------------------------------------------------------------------------------- *------------------------------------------------------------------------------------------------- - * @defgroup Memory Fence Functions + * @defgroup Fence Fence Functions * @{ * * @@ -545,7 +545,7 @@ __device__ void __threadfence(void) __attribute__((deprecated("Provided for com __device__ void __threadfence_system(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional"))); -// doxygen end Memory Fence +// doxygen end Fence Fence /** * @} */ From b14d2c8c8e0b0cfb50a32f47fe087927253b9a58 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 16 Sep 2016 15:05:48 +0530 Subject: [PATCH 23/25] Updates to doxygen documentation for HIP Event APIs Change-Id: Ic2e4a5725aebca14ac824da43436aacd0fe0a45c --- hipamd/include/hcc_detail/hip_runtime_api.h | 33 +++++++++++++-------- hipamd/src/hip_event.cpp | 14 --------- 2 files changed, 21 insertions(+), 26 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index db1faefa89..7768784e9a 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hcc_detail/hip_runtime_api.h @@ -596,11 +596,13 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags); * @brief Create an event with the specified flags * * @param[in,out] event Returns the newly created event. - * @param[in] flags Flags to control event behavior. #hipEventDefault, #hipEventBlockingSync, #hipEventDisableTiming, #hipEventInterprocess + * @param[in] flags Flags to control event behavior. Valid values are #hipEventDefault, #hipEventBlockingSync, #hipEventDisableTiming, #hipEventInterprocess * - * @warning On HCC platform, #hipEventInterprocess is not supported. + * @warning On HCC platform, flags must be #hipEventDefault. * - * @returns #cudaSuccess + * @returns #hipSuccess, #hipErrorInitializationError, #hipErrorInvalidValue, #hipErrorLaunchFailure, #hipErrorMemoryAllocation + * + * @see hipEventCreate, hipEventSynchronize, hipEventDestroy, hipEventElapsedTime */ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags); @@ -610,6 +612,9 @@ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags); * * @param[in,out] event Returns the newly created event. * + * @returns #hipSuccess, #hipErrorInitializationError, #hipErrorInvalidValue, #hipErrorLaunchFailure, #hipErrorMemoryAllocation + * + * @see hipEventCreateWithFlags, hipEventRecord, hipEventQuery, hipEventSynchronize, hipEventDestroy, hipEventElapsedTime */ hipError_t hipEventCreate(hipEvent_t* event); @@ -619,7 +624,7 @@ hipError_t hipEventCreate(hipEvent_t* event); * * @param[in] event event to record. * @param[in] stream stream in which to record event. - * @returns #hipSuccess, #hipErrorInvalidResourceHandle + * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError, #hipErrorInvalidResourceHandle, #hipErrorLaunchFailure * * hipEventQuery or hipEventSynchronize must be used to determine when the event * transitions from "recording" (after eventRecord is called) to "recorded" @@ -635,7 +640,7 @@ hipError_t hipEventCreate(hipEvent_t* event); * If this function is called on a 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. This shoul be avoided. * - * @see hipEventElapsedTime + * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventSynchronize, hipEventDestroy, hipEventElapsedTime * */ #ifdef __cplusplus @@ -648,17 +653,18 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream); * @brief Destroy the specified event. * * @param[in] event Event to destroy. - * @return : #hipSuccess, + * @returns #hipSuccess, #hipErrorInitializationError, #hipErrorInvalidValue, #hipErrorLaunchFailure * * Releases memory associated with the event. If the event is recording but has not completed recording when hipEventDestroy is called, * the function will return immediately and the completion_future resources will be released later, when the hipDevice is synchronized. * + * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventSynchronize, hipEventRecord, hipEventElapsedTime */ hipError_t hipEventDestroy(hipEvent_t event); /** - * @brief: Wait for an event to complete. + * @brief Wait for an event to complete. * * This function will block until the event is ready, waiting for all previous work in the stream specified when event was recorded with hipEventRecord. * @@ -667,8 +673,9 @@ hipError_t hipEventDestroy(hipEvent_t event); * TODO-hcc - This function needs to support hipEventBlockingSync parameter. * * @param[in] event Event on which to wait. - * @return #hipSuccess, #hipErrorInvalidResourceHandle, + * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError, #hipErrorInvalidResourceHandle, #hipErrorLaunchFailure * + * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventDestroy, hipEventRecord, hipEventElapsedTime */ hipError_t hipEventSynchronize(hipEvent_t event); @@ -676,10 +683,10 @@ hipError_t hipEventSynchronize(hipEvent_t event); /** * @brief Return the elapsed time between two events. * - * @param[out]] ms : Return time between start and stop in ms. + * @param[out] ms : Return time between start and stop in ms. * @param[in] start : Start event. * @param[in] stop : Stop event. - * @return : #hipSuccess, #hipErrorInvalidResourceHandle, #hipErrorNotReady, + * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotReady, #hipErrorInvalidResourceHandle, #hipErrorInitializationError, #hipErrorLaunchFailure * * Computes the elapsed time between two events. Time is computed in ms, with * a resolution of approximately 1 us. @@ -696,6 +703,8 @@ hipError_t hipEventSynchronize(hipEvent_t event); * If hipEventRecord has been called on both events, but the timestamp has not yet been recorded on one or * both events (that is, hipEventQuery would return #hipErrorNotReady on at least one of the events), then * #hipErrorNotReady is returned. + * + * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventDestroy, hipEventRecord, hipEventSynchronize */ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop); @@ -704,13 +713,13 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop); * @brief Query event status * * @param[in] event Event to query. - * @returns #hipSuccess, hipEventNotReady + * @returns #hipSuccess, #hipErrorNotReady, #hipErrorInvalidResourceHandle, #hipErrorInvalidValue, #hipErrorInitializationError, #hipErrorLaunchFailure * * Query the status of the specified event. This function will return #hipErrorNotReady if all commands * in the appropriate stream (specified to hipEventRecord) have completed. If that work has not completed, * or if hipEventRecord was not called on the event, then hipSuccess is returned. * - * + * @see hipEventCreate, hipEventCreateWithFlags, hipEventRecord, hipEventDestroy, hipEventSynchronize, hipEventElapsedTime */ hipError_t hipEventQuery(hipEvent_t event) ; diff --git a/hipamd/src/hip_event.cpp b/hipamd/src/hip_event.cpp index ecc6038b82..77d33cb6c2 100644 --- a/hipamd/src/hip_event.cpp +++ b/hipamd/src/hip_event.cpp @@ -48,9 +48,6 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) return e; } -/** - * @warning : flags must be 0. - */ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) { HIP_INIT_API(event, flags); @@ -58,7 +55,6 @@ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) return ihipLogStatus(ihipEventCreate(event, flags)); } - hipError_t hipEventCreate(hipEvent_t* event) { HIP_INIT_API(event); @@ -66,8 +62,6 @@ hipError_t hipEventCreate(hipEvent_t* event) return ihipLogStatus(ihipEventCreate(event, 0)); } - -//--- hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { HIP_INIT_API(event, stream); @@ -100,8 +94,6 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) } } - -//--- hipError_t hipEventDestroy(hipEvent_t event) { HIP_INIT_API(event); @@ -115,8 +107,6 @@ hipError_t hipEventDestroy(hipEvent_t event) return ihipLogStatus(hipSuccess); } - -//--- hipError_t hipEventSynchronize(hipEvent_t event) { HIP_INIT_API(event); @@ -142,8 +132,6 @@ hipError_t hipEventSynchronize(hipEvent_t event) } } - -//--- hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) { HIP_INIT_API(ms, start, stop); @@ -187,8 +175,6 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) return ihipLogStatus(status); } - -//--- hipError_t hipEventQuery(hipEvent_t event) { HIP_INIT_API(event); From 1dc1c056d4be0a28b1c08d7debe9bd8f290b6ba1 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 16 Sep 2016 15:18:17 +0530 Subject: [PATCH 24/25] Fix typo in comment Change-Id: I217012e10d8e5322aa0dc7dd5915be10df35fb53 --- hipamd/include/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/include/hip_runtime_api.h b/hipamd/include/hip_runtime_api.h index 660385972a..3406bcbbc9 100644 --- a/hipamd/include/hip_runtime_api.h +++ b/hipamd/include/hip_runtime_api.h @@ -141,7 +141,7 @@ typedef struct hipPointerAttribute_t { /* - asdasd* @brief hipError_t + * @brief hipError_t * @enum * @ingroup Enumerations */ @@ -188,7 +188,7 @@ typedef enum hipError_t { hipErrorMissingConfiguration = 1001, hipErrorMemoryAllocation = 1002, ///< Memory allocation error. hipErrorInitializationError = 1003, ///< TODO comment from hipErrorInitializationError - hipErrorLaunchFailure = 1004, + hipErrorLaunchFailure = 1004, ///< An exception occurred on the device while executing a kernel. hipErrorPriorLaunchFailure = 1005, hipErrorLaunchTimeOut = 1006, hipErrorLaunchOutOfResources = 1007, ///< Out of resources error. From 70a1a498f68ecaf0f3c3fae34719230aed1fb7e7 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 16 Sep 2016 15:26:58 +0530 Subject: [PATCH 25/25] Some more event API documentation fixes Change-Id: I159706e6771bec5270dda81800c1c6fd5f13d254 --- hipamd/include/hcc_detail/hip_runtime_api.h | 24 ++++++++++----------- 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index 7768784e9a..3de715cbbc 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hcc_detail/hip_runtime_api.h @@ -626,8 +626,8 @@ hipError_t hipEventCreate(hipEvent_t* event); * @param[in] stream stream in which to record event. * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError, #hipErrorInvalidResourceHandle, #hipErrorLaunchFailure * - * hipEventQuery or hipEventSynchronize must be used to determine when the event - * transitions from "recording" (after eventRecord is called) to "recorded" + * hipEventQuery() or hipEventSynchronize() must be used to determine when the event + * transitions from "recording" (after hipEventRecord() is called) to "recorded" * (when timestamps are set, if requested). * * Events which are recorded in a non-NULL stream will transition to @@ -635,7 +635,7 @@ hipError_t hipEventCreate(hipEvent_t* event); * the specified stream, after all previous * commands in that stream have completed executing. * - * If hipEventRecord has been previously called aon event, then this call will overwrite any existing state in event. + * If hipEventRecord() has been previously called aon 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 - either * outstanding recording may save state into the event, and the order is not guaranteed. This shoul be avoided. @@ -655,7 +655,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream); * @param[in] event Event to destroy. * @returns #hipSuccess, #hipErrorInitializationError, #hipErrorInvalidValue, #hipErrorLaunchFailure * - * Releases memory associated with the event. If the event is recording but has not completed recording when hipEventDestroy is called, + * Releases memory associated with the event. If the event is recording but has not completed recording when hipEventDestroy() is called, * the function will return immediately and the completion_future resources will be released later, when the hipDevice is synchronized. * * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventSynchronize, hipEventRecord, hipEventElapsedTime @@ -666,9 +666,9 @@ hipError_t hipEventDestroy(hipEvent_t event); /** * @brief Wait for an event to complete. * - * This function will block until the event is ready, waiting for all previous work in the stream specified when event was recorded with hipEventRecord. + * This function will block until the event is ready, waiting for all previous work in the stream specified when event was recorded with hipEventRecord(). * - * If hipEventRecord has not been called on @p event, this function returns immediately. + * If hipEventRecord() has not been called on @p event, this function returns immediately. * * TODO-hcc - This function needs to support hipEventBlockingSync parameter. * @@ -697,11 +697,11 @@ hipError_t hipEventSynchronize(hipEvent_t event); * Events which are recorded in a non-NULL stream will record their timestamp * when they reach the head of the specified stream, after all previous * commands in that stream have completed executing. Thus the time that - * the event recorded may be significantly after the host calls hipEventRecord. + * the event recorded may be significantly after the host calls hipEventRecord(). * - * If hipEventRecord has not been called on either event, then #hipErrorInvalidResourceHandle is returned. - * If hipEventRecord has been called on both events, but the timestamp has not yet been recorded on one or - * both events (that is, hipEventQuery would return #hipErrorNotReady on at least one of the events), then + * If hipEventRecord() has not been called on either event, then #hipErrorInvalidResourceHandle is returned. + * If hipEventRecord() has been called on both events, but the timestamp has not yet been recorded on one or + * both events (that is, hipEventQuery() would return #hipErrorNotReady on at least one of the events), then * #hipErrorNotReady is returned. * * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventDestroy, hipEventRecord, hipEventSynchronize @@ -716,8 +716,8 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop); * @returns #hipSuccess, #hipErrorNotReady, #hipErrorInvalidResourceHandle, #hipErrorInvalidValue, #hipErrorInitializationError, #hipErrorLaunchFailure * * Query the status of the specified event. This function will return #hipErrorNotReady if all commands - * in the appropriate stream (specified to hipEventRecord) have completed. If that work has not completed, - * or if hipEventRecord was not called on the event, then hipSuccess is returned. + * in the appropriate stream (specified to hipEventRecord()) have completed. If that work has not completed, + * or if hipEventRecord() was not called on the event, then #hipSuccess is returned. * * @see hipEventCreate, hipEventCreateWithFlags, hipEventRecord, hipEventDestroy, hipEventSynchronize, hipEventElapsedTime */