diff --git a/.vimrc b/.vimrc new file mode 100644 index 0000000000..ed64acd347 --- /dev/null +++ b/.vimrc @@ -0,0 +1,4 @@ +:set tabstop=4 +:set shiftwidth=4 +:set expandtab +:set smartindent diff --git a/clang-hipify/src/Cuda2Hip.cpp b/clang-hipify/src/Cuda2Hip.cpp index 77d8b0259e..885626f0b4 100644 --- a/clang-hipify/src/Cuda2Hip.cpp +++ b/clang-hipify/src/Cuda2Hip.cpp @@ -322,6 +322,12 @@ struct cuda2hipMap { cuda2hipRename["CU_STREAM_DEFAULT"] = {"hipStreamDefault", CONV_STREAM, API_DRIVER}; cuda2hipRename["CU_STREAM_NON_BLOCKING"] = {"hipStreamNonBlocking", CONV_STREAM, API_DRIVER}; + // Init + cuda2hipRename["cuInit"] = {"hipInit", CONV_DRIVER, API_DRIVER}; + + // Driver + cuda2hipRename["cuDriverGetVersion"] = {"hipDriverGetVersion", CONV_DRIVER, API_DRIVER}; + // Context cuda2hipRename["cuCtxCreate_v2"] = {"hipCtxCreate", CONV_CONTEXT, API_DRIVER}; cuda2hipRename["cuCtxDestroy_v2"] = {"hipCtxDestroy", CONV_CONTEXT, API_DRIVER}; @@ -356,9 +362,6 @@ struct cuda2hipMap { cuda2hipRename["cuDeviceComputeCapability"] = {"hipDeviceComputeCapability", CONV_DEV, API_DRIVER}; cuda2hipRename["cuDeviceCanAccessPeer"] = {"hipDeviceCanAccessPeer", CONV_DEV, API_DRIVER}; - // Driver - cuda2hipRename["cuDriverGetVersion"] = {"hipDriverGetVersion", CONV_DRIVER, API_DRIVER}; - // Events cuda2hipRename["cuEventCreate"] = {"hipEventCreate", CONV_EVENT, API_DRIVER}; cuda2hipRename["cuEventDestroy_v2"] = {"hipEventDestroy", CONV_EVENT, API_DRIVER}; @@ -387,6 +390,41 @@ struct cuda2hipMap { cuda2hipRename["cuStreamSynchronize"] = {"hipStreamSynchronize", CONV_STREAM, API_DRIVER}; cuda2hipRename["cuStreamWaitEvent"] = {"hipStreamWaitEvent", CONV_STREAM, API_DRIVER}; + // Memory management + cuda2hipRename["cuMemAlloc_v2"] = {"hipMalloc", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemFree_v2"] = {"hipFree", CONV_MEM, API_DRIVER}; + + cuda2hipRename["cuMemHostAlloc"] = {"hipHostMalloc", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemFreeHost"] = {"hipHostFree", CONV_MEM, API_DRIVER}; + + cuda2hipRename["cuMemcpyDtoD_v2"] = {"hipMemcpyDtoD", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemcpyDtoDAsync_v2"] = {"hipMemcpyDtoDAsync", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemcpyDtoH_v2"] = {"hipMemcpyDtoH", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemcpyDtoHAsync_v2"] = {"hipMemcpyDtoHAsync", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemcpyHtoD_v2"] = {"hipMemcpyHtoD", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemcpyHtoDAsync_v2"] = {"hipMemcpyHtoDAsync", CONV_MEM, API_DRIVER}; + + // unsupported yet by HIP + // cuda2hipRename["cuMemsetD8_v2"] = {"hipMemsetD8", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD8Async"] = {"hipMemsetD8Async", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD2D8_v2"] = {"hipMemsetD2D8", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD2D8Async"] = {"hipMemsetD2D8Async", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD16_v2"] = {"hipMemsetD16", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD16Async"] = {"hipMemsetD16Async", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD2D16_v2"] = {"hipMemsetD2D16", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD2D16Async"] = {"hipMemsetD2D16Async", CONV_STREAM, API_DRIVER}; + cuda2hipRename["cuMemsetD32_v2"] = {"hipMemset", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemsetD32Async"] = {"hipMemsetAsync", CONV_MEM, API_DRIVER}; + // unsupported yet by HIP + // cuda2hipRename["cuMemsetD2D32_v2"] = {"hipMemsetD2D32", CONV_STREAM, API_DRIVER}; + // cuda2hipRename["cuMemsetD2D32Async"] = {"hipMemsetD2D32Async", CONV_STREAM, API_DRIVER}; + + cuda2hipRename["cuMemGetInfo_v2"] = {"hipMemGetInfo", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemHostRegister_v2"] = {"hipHostRegister", CONV_MEM, API_DRIVER}; + cuda2hipRename["cuMemHostUnregister"] = {"hipHostUnregister", CONV_MEM, API_DRIVER}; + + + /////////////////////////////// CUDA RT API /////////////////////////////// // Error API cuda2hipRename["cudaGetLastError"] = {"hipGetLastError", CONV_ERR, API_RUNTIME}; diff --git a/cmake/FindHIP.cmake b/cmake/FindHIP.cmake index 2541cf6ec6..9626e7629f 100644 --- a/cmake/FindHIP.cmake +++ b/cmake/FindHIP.cmake @@ -342,6 +342,24 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files) HIP_PARSE_HIPCC_OPTIONS(HIP_HCC_FLAGS ${_hcc_options}) HIP_PARSE_HIPCC_OPTIONS(HIP_NVCC_FLAGS ${_nvcc_options}) + # Check if we are building shared library. + set(_hip_build_shared_libs FALSE) + list(FIND _hip_cmake_options SHARED _hip_found_SHARED) + list(FIND _hip_cmake_options MODULE _hip_found_MODULE) + if(_hip_found_SHARED GREATER -1 OR _hip_found_MODULE GREATER -1) + set(_hip_build_shared_libs TRUE) + endif() + list(FIND _hip_cmake_options STATIC _hip_found_STATIC) + if(_hip_found_STATIC GREATER -1) + set(_hip_build_shared_libs FALSE) + endif() + + # If we are building a shared library, add extra flags to HIP_HIPCC_FLAGS + if(_hip_build_shared_libs) + list(APPEND HIP_HCC_FLAGS "-fPIC") + list(APPEND HIP_NVCC_FLAGS "--shared -Xcompiler '-fPIC'") + endif() + # Set host compiler set(HIP_HOST_COMPILER "${CMAKE_${HIP_C_OR_CXX}_COMPILER}") @@ -416,7 +434,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( @@ -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 diff --git a/cmake/FindHIP/run_hipcc.cmake b/cmake/FindHIP/run_hipcc.cmake index 52ad57f532..6027cb9b0d 100644 --- a/cmake/FindHIP/run_hipcc.cmake +++ b/cmake/FindHIP/run_hipcc.cmake @@ -22,8 +22,8 @@ endif() # Set these up as variables to make reading the generated file easier set(HIP_HIPCC_EXECUTABLE "@HIP_HIPCC_EXECUTABLE@") # path +set(HIP_HIPCONFIG_EXECUTABLE "@HIP_HIPCONFIG_EXECUTABLE@") #path set(HIP_HOST_COMPILER "@HIP_HOST_COMPILER@") # path -set(HIP_PLATFORM "@HIP_PLATFORM@") #string set(CMAKE_COMMAND "@CMAKE_COMMAND@") # path set(HIP_run_make2cmake "@HIP_run_make2cmake@") # path @@ -38,6 +38,7 @@ set(source_file "@source_file@") # path set(host_flag "@host_flag@") # bool # Determine compiler and compiler flags +execute_process(COMMAND ${HIP_HIPCONFIG_EXECUTABLE} --platform OUTPUT_VARIABLE HIP_PLATFORM OUTPUT_STRIP_TRAILING_WHITESPACE) if(NOT host_flag) set(__CC ${HIP_HIPCC_EXECUTABLE}) if(HIP_PLATFORM STREQUAL "hcc") diff --git a/include/hcc_detail/hip_runtime.h b/include/hcc_detail/hip_runtime.h index 547df405a2..92406b4a75 100644 --- a/include/hcc_detail/hip_runtime.h +++ b/include/hcc_detail/hip_runtime.h @@ -500,7 +500,7 @@ __device__ float __dsqrt_rz(double x); /** *------------------------------------------------------------------------------------------------- *------------------------------------------------------------------------------------------------- - * @defgroup Memory Fence Functions + * @defgroup Fence Fence Functions * @{ * * @@ -545,7 +545,7 @@ __device__ void __threadfence(void) __attribute__((deprecated("Provided for com __device__ void __threadfence_system(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional"))); -// doxygen end Memory Fence +// doxygen end Fence Fence /** * @} */ diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index 71e46612de..3de715cbbc 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/include/hcc_detail/hip_runtime_api.h @@ -189,9 +189,12 @@ typedef enum hipMemcpyKind { */ /** - * @brief Blocks until the default device has completed all preceding requested tasks. + * @brief Waits on all active streams on current device * - * This function waits for all streams on the default device to complete execution, and then returns. + * When this command is invoked, the host thread gets blocked until all the commands associated + * with streams associated with the device. HIP does not support multiple blocking modes (yet!). + * + * @returns #hipSuccess * * @see hipSetDevice, hipDeviceReset */ @@ -200,12 +203,12 @@ hipError_t hipDeviceSynchronize(void); /** - * @brief Destroy all resources and reset all state on the default device in the current process. + * @brief The state of current device is discarded and updated to a fresh state. * - * Explicity destroy all memory allocations, events, and queues associated with the default device in the current process. + * Calling this function deletes all streams created, memory allocated, kernels running, events created. + * Make sure that no other thread is using the device or streams, memory, kernels, events associated with the current device. * - * This function will reset the device immmediately, and then return after all resources have been freed. - * The caller must ensure that the device is not being accessed by any other host threads from the active process when this function is called. + * @returns #hipSuccess * * @see hipDeviceSynchronize */ @@ -234,6 +237,8 @@ hipError_t hipDeviceReset(void) ; * Thread-pool implementations may inherit the default device of the previous thread. A good practice is to always call hipSetDevice * at the start of HIP coding sequency to establish a known standard device. * + * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorDeviceAlreadyInUse + * * @see hipGetDevice, hipGetDeviceCount */ hipError_t hipSetDevice(int deviceId); @@ -248,27 +253,34 @@ 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, hipErrorInvalidDevice + * @see hipSetDevice, hipGetDevicesizeBytes */ 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); @@ -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); /** *------------------------------------------------------------------------------------------------- @@ -375,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); @@ -390,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); @@ -572,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, flags must be #hipEventDefault. * - * @warning On HCC platform, #hipEventInterprocess is not supported. + * @returns #hipSuccess, #hipErrorInitializationError, #hipErrorInvalidValue, #hipErrorLaunchFailure, #hipErrorMemoryAllocation * - * @returns #cudaSuccess + * @see hipEventCreate, hipEventSynchronize, hipEventDestroy, hipEventElapsedTime */ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags); @@ -586,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); @@ -595,10 +624,10 @@ 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" + * 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 @@ -606,12 +635,12 @@ 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. * - * @see hipEventElapsedTime + * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventSynchronize, hipEventDestroy, hipEventElapsedTime * */ #ifdef __cplusplus @@ -624,27 +653,29 @@ 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, + * 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. + * 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. * * @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); @@ -652,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. @@ -666,12 +697,14 @@ 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 */ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop); @@ -680,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. - * + * 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) ; @@ -715,6 +748,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); @@ -723,7 +763,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) ; @@ -731,19 +774,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"))) ;; @@ -751,19 +800,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) ; @@ -796,6 +850,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) ; @@ -804,6 +860,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) ; @@ -818,6 +876,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); @@ -830,6 +890,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); @@ -838,6 +900,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"))) ; @@ -851,6 +917,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); @@ -873,17 +941,89 @@ 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, #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 + */ 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, #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 + */ 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, #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 + */ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes); -hipError_t hipMemcpyHtoH(void* dst, void* 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, #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 + */ +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, #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 + */ +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, #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 + */ +hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream); /** @@ -899,6 +1039,8 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes); * @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); @@ -909,6 +1051,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. @@ -921,6 +1064,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); @@ -966,7 +1111,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/include/hip_runtime_api.h b/include/hip_runtime_api.h index 4cb5bb9bcd..3406bcbbc9 100644 --- a/include/hip_runtime_api.h +++ b/include/hip_runtime_api.h @@ -141,7 +141,7 @@ typedef struct hipPointerAttribute_t { /* - asdasd* @brief hipError_t + * @brief hipError_t * @enum * @ingroup Enumerations */ @@ -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. - - 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. + hipErrorMissingConfiguration = 1001, + hipErrorMemoryAllocation = 1002, ///< Memory allocation error. + hipErrorInitializationError = 1003, ///< TODO comment from hipErrorInitializationError + hipErrorLaunchFailure = 1004, ///< An exception occurred on the device while executing a kernel. + 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 = 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/include/nvcc_detail/hip_runtime_api.h b/include/nvcc_detail/hip_runtime_api.h index 8e9b0d92a7..c90c7cbf29 100644 --- a/include/nvcc_detail/hip_runtime_api.h +++ b/include/nvcc_detail/hip_runtime_api.h @@ -214,6 +214,29 @@ inline static hipError_t hipSetDevice(int device) { return hipCUDAErrorTohipError(cudaSetDevice(device)); } +inline static hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop ) +{ + cudaDeviceProp cdprop; + memset(&cdprop,0x0,sizeof(cudaDeviceProp)); + cdprop.major= prop->major; + cdprop.minor = prop->minor; + cdprop.totalGlobalMem = prop->totalGlobalMem ; + cdprop.sharedMemPerBlock = prop->sharedMemPerBlock; + cdprop.regsPerBlock = prop->regsPerBlock; + cdprop.warpSize = prop->warpSize ; + cdprop.maxThreadsPerBlock = prop->maxThreadsPerBlock ; + cdprop.clockRate = prop->clockRate; + cdprop.totalConstMem = prop->totalConstMem ; + cdprop.multiProcessorCount = prop->multiProcessorCount ; + cdprop.l2CacheSize = prop->l2CacheSize ; + cdprop.maxThreadsPerMultiProcessor = prop->maxThreadsPerMultiProcessor ; + cdprop.computeMode = prop->computeMode ; + cdprop.canMapHostMemory = prop->canMapHostMemory; + cdprop.memoryClockRate = prop->memoryClockRate; + cdprop.memoryBusWidth = prop->memoryBusWidth; + return hipCUDAErrorTohipError(cudaChooseDevice(device,&cdprop)); +} + inline static hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t size) { @@ -232,6 +255,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)); +} + +inline static hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, + hipDeviceptr_t src, size_t size, hipStream_t stream) +{ + return hipCUResultTohipError(cuMemcpyDtoD(dst, src, size)); +} + inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind) { return hipCUDAErrorTohipError(cudaMemcpy(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind))); } diff --git a/samples/0_Intro/square/square.hipref.cpp b/samples/0_Intro/square/square.hipref.cpp index aa14077738..7ca3a7500d 100644 --- a/samples/0_Intro/square/square.hipref.cpp +++ b/samples/0_Intro/square/square.hipref.cpp @@ -31,7 +31,6 @@ THE SOFTWARE. }\ } - /* * Square each element in the array A and write to array C. */ @@ -43,55 +42,54 @@ vector_square(hipLaunchParm lp, T *C_d, const T *A_d, size_t N) size_t stride = hipBlockDim_x * hipGridDim_x ; for (size_t i=offset; itotalGlobalMem != 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; } diff --git a/src/hip_error.cpp b/src/hip_error.cpp index d9c6dd9aa9..97fcaaf714 100644 --- a/src/hip_error.cpp +++ b/src/hip_error.cpp @@ -25,9 +25,7 @@ THE SOFTWARE. //------------------------------------------------------------------------------------------------- // Error Handling //--- -/** - * @returns return code from last HIP called from the active host thread. - */ + hipError_t hipGetLastError() { HIP_INIT_API(); @@ -38,8 +36,6 @@ hipError_t hipGetLastError() return e; } - -//--- hipError_t hipPeekAtLastError() { HIP_INIT_API(); @@ -48,7 +44,6 @@ hipError_t hipPeekAtLastError() return ihipLogStatus(tls_lastHipError); } -//--- const char *hipGetErrorName(hipError_t hip_error) { HIP_INIT_API(hip_error); @@ -56,12 +51,6 @@ const char *hipGetErrorName(hipError_t hip_error) return ihipErrorString(hip_error); } - -/** - * @warning : hipGetErrorString returns string from hipGetErrorName - */ - -//--- const char *hipGetErrorString(hipError_t hip_error) { HIP_INIT_API(hip_error); diff --git a/src/hip_event.cpp b/src/hip_event.cpp index ecc6038b82..77d33cb6c2 100644 --- a/src/hip_event.cpp +++ b/src/hip_event.cpp @@ -48,9 +48,6 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) return e; } -/** - * @warning : flags must be 0. - */ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) { HIP_INIT_API(event, flags); @@ -58,7 +55,6 @@ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) return ihipLogStatus(ihipEventCreate(event, flags)); } - hipError_t hipEventCreate(hipEvent_t* event) { HIP_INIT_API(event); @@ -66,8 +62,6 @@ hipError_t hipEventCreate(hipEvent_t* event) return ihipLogStatus(ihipEventCreate(event, 0)); } - -//--- hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { HIP_INIT_API(event, stream); @@ -100,8 +94,6 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) } } - -//--- hipError_t hipEventDestroy(hipEvent_t event) { HIP_INIT_API(event); @@ -115,8 +107,6 @@ hipError_t hipEventDestroy(hipEvent_t event) return ihipLogStatus(hipSuccess); } - -//--- hipError_t hipEventSynchronize(hipEvent_t event) { HIP_INIT_API(event); @@ -142,8 +132,6 @@ hipError_t hipEventSynchronize(hipEvent_t event) } } - -//--- hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) { HIP_INIT_API(ms, start, stop); @@ -187,8 +175,6 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) return ihipLogStatus(status); } - -//--- hipError_t hipEventQuery(hipEvent_t event) { HIP_INIT_API(event); diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index e2fb952a5b..0f0b03d115 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -1559,20 +1559,64 @@ const char *ihipErrorString(hipError_t hip_error) { switch (hip_error) { case hipSuccess : return "hipSuccess"; + 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"; }; diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 8516d7520d..0478ff26d7 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -28,12 +28,6 @@ // Memory // // -// - -//--- -/** - * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice - */ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) { HIP_INIT_API(attributes, ptr); @@ -77,12 +71,6 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) return ihipLogStatus(e); } - -/** - * @returns #hipSuccess, - * @returns #hipErrorInvalidValue if flags are not 0 - * @returns #hipErrorMemoryAllocation if hostPointer is not a tracked allocation. - */ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsigned flags) { HIP_INIT_API(devicePointer, hostPointer, flags); @@ -107,13 +95,6 @@ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsi return ihipLogStatus(e); } - - - -//--- -/** - * @returns #hipSuccess #hipErrorMemoryAllocation - */ hipError_t hipMalloc(void** ptr, size_t sizeBytes) { HIP_INIT_API(ptr, sizeBytes); @@ -153,8 +134,6 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) return ihipLogStatus(hip_status); } - - hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_INIT_API(ptr, sizeBytes, flags); @@ -193,7 +172,6 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) return ihipLogStatus(hip_status); } - //--- // TODO - remove me, this is deprecated. hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) @@ -201,7 +179,6 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) return hipHostMalloc(ptr, sizeBytes, flags); }; - //--- // TODO - remove me, this is deprecated. hipError_t hipMallocHost(void** ptr, size_t sizeBytes) @@ -253,7 +230,6 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height return ihipLogStatus(hip_status); } - hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { hipChannelFormatDesc cd; @@ -262,7 +238,6 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel return cd; } - hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) { @@ -324,8 +299,6 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, return ihipLogStatus(hip_status); } - -//--- hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) { HIP_INIT_API(flagsPtr, hostPtr); @@ -350,8 +323,6 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) return ihipLogStatus(hip_status); } - -//--- hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) { HIP_INIT_API(hostPtr, sizeBytes, flags); @@ -395,7 +366,6 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) return ihipLogStatus(hip_status); } -//--- hipError_t hipHostUnregister(void *hostPtr) { HIP_INIT_API(hostPtr); @@ -413,8 +383,6 @@ hipError_t hipHostUnregister(void *hostPtr) return ihipLogStatus(hip_status); } - -//--- hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind) { HIP_INIT_API(symbolName, src, count, offset, kind); @@ -478,7 +446,6 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) return ihipLogStatus(e); } - hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { HIP_INIT_API(dst, src, sizeBytes); @@ -542,15 +509,6 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) return ihipLogStatus(e); } - - -/** - * @result #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidMemcpyDirection, - * @result #hipErrorInvalidValue : If dst==NULL or src==NULL, or other bad argument. - * @warning on HCC hipMemcpyAsync does not support overlapped H2D and D2H copies. - * @warning on HCC hipMemcpyAsync requires that any host pointers are pinned (ie via the hipMallocHost call). - */ -//--- hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_API(dst, src, sizeBytes, kind, stream); @@ -576,7 +534,85 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp return ihipLogStatus(e); } -// dpitch, spitch, and width in bytes +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); +} + hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { @@ -603,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) { @@ -678,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 @@ -718,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); @@ -772,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; @@ -824,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); @@ -857,8 +880,6 @@ hipError_t hipMemGetInfo (size_t *free, size_t *total) return ihipLogStatus(e); } - -//--- hipError_t hipFree(void* ptr) { HIP_INIT_API(ptr); @@ -886,7 +907,6 @@ hipError_t hipFree(void* ptr) return ihipLogStatus(hipStatus); } - hipError_t hipHostFree(void* ptr) { HIP_INIT_API(ptr); @@ -914,7 +934,6 @@ hipError_t hipHostFree(void* ptr) return ihipLogStatus(hipStatus); }; - // TODO - deprecated function. hipError_t hipFreeHost(void* ptr) { diff --git a/tests/src/CMakeLists.txt b/tests/src/CMakeLists.txt index 43a74910b2..b58b930f7f 100644 --- a/tests/src/CMakeLists.txt +++ b/tests/src/CMakeLists.txt @@ -1,4 +1,4 @@ -cmake_minimum_required (VERSION 2.6) +cmake_minimum_required(VERSION 2.6) # remove CMAKE_CXX_COMPILER entry from cache since it will be pointing to hipcc unset(CMAKE_CXX_COMPILER CACHE) @@ -9,194 +9,209 @@ 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 - ) - 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) - 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() + +# 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() - # We are using HIP_PATH passed to cmake. So just create a fake target. - 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) 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) - -if (${HIP_PLATFORM} STREQUAL "hcc") - build_hip_executable (hipArray hipArray.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) 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 +219,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 +240,4 @@ add_subdirectory(context) add_subdirectory(deviceLib) add_subdirectory(runtimeApi) add_subdirectory(kernel) +# vim: ts=4:sw=4:expandtab:smartindent diff --git a/tests/src/hipDrvMemcpy.cpp b/tests/src/hipDrvMemcpy.cpp index 55c9b18818..8322464964 100644 --- a/tests/src/hipDrvMemcpy.cpp +++ b/tests/src/hipDrvMemcpy.cpp @@ -6,21 +6,39 @@ #define SIZE LEN<<2 int main(){ - int *A, *B, *C; - hipDeviceptr Ad, Bd; + int *A, *B; + hipDeviceptr_t Ad, Bd; A = new int[LEN]; B = new int[LEN]; - C = new int[LEN]; + for(int i=0;i