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