From 2d5140cb27f6beed714f3cac7d88faf6ff2ded59 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 --- include/hcc_detail/hip_runtime_api.h | 6 ++- src/hip_memory.cpp | 79 ++++++++++++++++++++++++++++ tests/src/hipDrvMemcpy.cpp | 28 ++++++++-- 3 files changed, 107 insertions(+), 6 deletions(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index 71e46612de..c782696753 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/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/src/hip_memory.cpp b/src/hip_memory.cpp index 8516d7520d..19b4745b59 100644 --- a/src/hip_memory.cpp +++ b/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/tests/src/hipDrvMemcpy.cpp b/tests/src/hipDrvMemcpy.cpp index 55c9b18818..8322464964 100644 --- a/tests/src/hipDrvMemcpy.cpp +++ b/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 --- include/nvcc_detail/hip_runtime_api.h | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/include/nvcc_detail/hip_runtime_api.h b/include/nvcc_detail/hip_runtime_api.h index 8e9b0d92a7..1ad1af2267 100644 --- a/include/nvcc_detail/hip_runtime_api.h +++ b/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 2cc9873c6b88cb108a4cbd88a69f512b231fb945 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 --- include/hcc_detail/hip_runtime_api.h | 84 ++++++++++++++++++---------- 1 file changed, 53 insertions(+), 31 deletions(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index c782696753..2b1b116c6e 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/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 1b2361ff224cf833246972ce605a3950e64a940b 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 --- include/hcc_detail/hip_runtime_api.h | 6 ++++-- src/hip_error.cpp | 13 +------------ 2 files changed, 5 insertions(+), 14 deletions(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index 2b1b116c6e..a08974f8b5 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/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/src/hip_error.cpp b/src/hip_error.cpp index d9c6dd9aa9..97fcaaf714 100644 --- a/src/hip_error.cpp +++ b/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 f33a0dd4a2dc39eaae6cc352de48f1ece6b62654 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 --- include/hcc_detail/hip_runtime_api.h | 59 +++++++++++++++++++++------- 1 file changed, 45 insertions(+), 14 deletions(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index a08974f8b5..221dfbf768 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/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 67de23d6b9f62f950ecebec92d94e7737bf5f37f 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. --- clang-hipify/src/Cuda2Hip.cpp | 44 ++++++++++++++++++++++++++++++++--- 1 file changed, 41 insertions(+), 3 deletions(-) diff --git a/clang-hipify/src/Cuda2Hip.cpp b/clang-hipify/src/Cuda2Hip.cpp index 77d8b0259e..885626f0b4 100644 --- a/clang-hipify/src/Cuda2Hip.cpp +++ b/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 9749a10abcd232d3fa7e05dcd42e2f82d1684862 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 --- tests/src/CMakeLists.txt | 301 ++++++++++++++++++++++----------------- 1 file changed, 171 insertions(+), 130 deletions(-) diff --git a/tests/src/CMakeLists.txt b/tests/src/CMakeLists.txt index 43a74910b2..2563c31bab 100644 --- a/tests/src/CMakeLists.txt +++ b/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 92fcd74a386748f1b776e3b8ef5f729988d65c5d 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 --- tests/src/CMakeLists.txt | 165 +++++++++++++++++---------------------- 1 file changed, 70 insertions(+), 95 deletions(-) diff --git a/tests/src/CMakeLists.txt b/tests/src/CMakeLists.txt index 2563c31bab..b58b930f7f 100644 --- a/tests/src/CMakeLists.txt +++ b/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 9e15b59dc491fc14d4427fd129fc5c74a83f92c6 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 --- include/hcc_detail/hip_runtime_api.h | 78 +++++++++++++++++++++++++++- 1 file changed, 77 insertions(+), 1 deletion(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index 221dfbf768..6c42d39041 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/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 af8d308745678c92f05c62ed0cc30ff2d39e7449 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 --- include/nvcc_detail/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/nvcc_detail/hip_runtime_api.h b/include/nvcc_detail/hip_runtime_api.h index 1ad1af2267..21ef239894 100644 --- a/include/nvcc_detail/hip_runtime_api.h +++ b/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 8110f562ab59fc1a0baebba377aa2839ff2fbe8d 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 --- src/hip_hcc.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index e2fb952a5b..a1d6e2537f 100644 --- a/src/hip_hcc.cpp +++ b/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 7407cb26006d43ee932fcb1b64e67ad79229745f 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 --- include/hip_runtime_api.h | 46 ++++++++++++++++---------------- src/hip_hcc.cpp | 55 ++++++++++++++++++++++++++++++++++----- 2 files changed, 72 insertions(+), 29 deletions(-) diff --git a/include/hip_runtime_api.h b/include/hip_runtime_api.h index 4cb5bb9bcd..660385972a 100644 --- a/include/hip_runtime_api.h +++ b/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/src/hip_hcc.cpp b/src/hip_hcc.cpp index a1d6e2537f..0f0b03d115 100644 --- a/src/hip_hcc.cpp +++ b/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 49118763711e105e1f52a1eb97aefbf826f89017 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 --- .gitignore | 1 - .vimrc | 4 ++ samples/0_Intro/square/square.hipref.cpp | 66 ++++++++++++------------ 3 files changed, 36 insertions(+), 35 deletions(-) create mode 100644 .vimrc diff --git a/.gitignore b/.gitignore index 450d3bc41b..e298b686bf 100644 --- a/.gitignore +++ b/.gitignore @@ -1,4 +1,3 @@ -.* *.o *.exe *.swp diff --git a/.vimrc b/.vimrc new file mode 100644 index 0000000000..ed64acd347 --- /dev/null +++ b/.vimrc @@ -0,0 +1,4 @@ +:set tabstop=4 +:set shiftwidth=4 +:set expandtab +:set smartindent diff --git a/samples/0_Intro/square/square.hipref.cpp b/samples/0_Intro/square/square.hipref.cpp index aa14077738..7ca3a7500d 100644 --- a/samples/0_Intro/square/square.hipref.cpp +++ b/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 --- cmake/FindHIP.cmake | 2 +- cmake/FindHIP/run_hipcc.cmake | 3 ++- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/cmake/FindHIP.cmake b/cmake/FindHIP.cmake index 2541cf6ec6..1666673beb 100644 --- a/cmake/FindHIP.cmake +++ b/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/cmake/FindHIP/run_hipcc.cmake b/cmake/FindHIP/run_hipcc.cmake index 52ad57f532..6027cb9b0d 100644 --- a/cmake/FindHIP/run_hipcc.cmake +++ b/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 a4c4be2020bb5fb01839ea7348eb602c1f3960f2 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 --- cmake/FindHIP.cmake | 29 +++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) diff --git a/cmake/FindHIP.cmake b/cmake/FindHIP.cmake index 1666673beb..9626e7629f 100644 --- a/cmake/FindHIP.cmake +++ b/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 f3f0c6ae3c36acfd2b6e372f2c003b3eaf8fc36f 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 --- include/nvcc_detail/hip_runtime_api.h | 23 +++++++++++++++++++++++ 1 file changed, 23 insertions(+) diff --git a/include/nvcc_detail/hip_runtime_api.h b/include/nvcc_detail/hip_runtime_api.h index 21ef239894..c90c7cbf29 100644 --- a/include/nvcc_detail/hip_runtime_api.h +++ b/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 13e88c54d7039757d7d9bb00d5b4d6e12b96550f 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 --- src/hip_device.cpp | 43 ++++++++++++++++++++++++++++++++++++++++++- 1 file changed, 42 insertions(+), 1 deletion(-) diff --git a/src/hip_device.cpp b/src/hip_device.cpp index fe7ad6ecd7..6ea4658a4f 100644 --- a/src/hip_device.cpp +++ b/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 71f9efd7583cc25ca285e123c13be8f1aab4f73a 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 --- .gitignore | 1 + 1 file changed, 1 insertion(+) diff --git a/.gitignore b/.gitignore index e298b686bf..450d3bc41b 100644 --- a/.gitignore +++ b/.gitignore @@ -1,3 +1,4 @@ +.* *.o *.exe *.swp From 76949b637400c07fe7950838767745b380b3b221 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 --- include/hcc_detail/hip_runtime_api.h | 5 ++- src/hip_memory.cpp | 60 ---------------------------- 2 files changed, 4 insertions(+), 61 deletions(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index 6c42d39041..d8b7fc4b55 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/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/src/hip_memory.cpp b/src/hip_memory.cpp index 19b4745b59..0478ff26d7 100644 --- a/src/hip_memory.cpp +++ b/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 347aaf7680545a68a8733972070419b57e2ed5ed 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 --- include/hcc_detail/hip_runtime_api.h | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index d8b7fc4b55..be8b513caf 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/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 126b2e0ecac211768f4acc9231ec8bc6d7119aeb 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 --- include/hcc_detail/hip_runtime_api.h | 30 ++++++++++++++-------------- 1 file changed, 15 insertions(+), 15 deletions(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index be8b513caf..db1faefa89 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/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 8ce8e1de76cac8cea3f2414754d994f3430fcba2 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 --- include/hcc_detail/hip_runtime.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/hcc_detail/hip_runtime.h b/include/hcc_detail/hip_runtime.h index 547df405a2..92406b4a75 100644 --- a/include/hcc_detail/hip_runtime.h +++ b/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 3d85cd0ce5b6d9e7fc9f01840907ae59d585a479 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 --- include/hcc_detail/hip_runtime_api.h | 33 ++++++++++++++++++---------- src/hip_event.cpp | 14 ------------ 2 files changed, 21 insertions(+), 26 deletions(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index db1faefa89..7768784e9a 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/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/src/hip_event.cpp b/src/hip_event.cpp index ecc6038b82..77d33cb6c2 100644 --- a/src/hip_event.cpp +++ b/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 8d1ab4c21a85b51596ef34be08238ee1f54165f4 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 --- include/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/hip_runtime_api.h b/include/hip_runtime_api.h index 660385972a..3406bcbbc9 100644 --- a/include/hip_runtime_api.h +++ b/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 22a78410d7015a5d4960da9857d821b1054ab5ca 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 --- include/hcc_detail/hip_runtime_api.h | 24 ++++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index 7768784e9a..3de715cbbc 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/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 */