Merge branch 'amd-develop' into amd-master

Change-Id: I9f284502dc764b7ad39bd9bc8a40ea7404d0391e


[ROCm/hip commit: 8ac4a51d87]

projects/hip/.vimrc

@@ -0,0 +1,4 @@
+:set tabstop=4
+:set shiftwidth=4
+:set expandtab
+:set smartindent

projects/hip/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};

projects/hip/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

projects/hip/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")

projects/hip/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
 /**
  * @}
  */

projects/hip/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.
+ * @returns #hipSuccess
- * 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.
  *
  * @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 Device which matches hipDeviceProp_t is returned
- */
-
-/**
- * @brief Select compute-device which best matches criteria.
  *
  * @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, #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);
 
@@ -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
+ * hipEventQuery() or hipEventSynchronize() must be used to determine when the event
- * transitions from "recording" (after eventRecord is called) to "recorded"
+ * 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 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
+ * 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
+ * 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,
+ * 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.
+ * 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[out] ptr Pointer to the allocated host pinned memory
- *  @param[in] size Requested memory size
+ *  @param[in]  size Requested memory size
- *  @return Error code
+ *
+ *  @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[out] ptr Pointer to the allocated host pinned memory
- *  @param[in] size Requested memory size
+ *  @param[in]  size Requested memory size
- *  @param[in] flags Type of host memory allocation
+ *  @param[in]  flags Type of host memory allocation
- *  @return Error code
+ *
+ *  @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[out] dstPtr Device Pointer mapped to passed host pointer
- *  @param[in] hstPtr Host Pointer allocated through hipHostAlloc
+ *  @param[in]  hstPtr Host Pointer allocated through hipHostAlloc
- *  @param[in] flags Flags to be passed for extension
+ *  @param[in]  flags Flags to be passed for extension
- *  @return Error code
+ * 
+ *  @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[out] flagsPtr Memory location to store flags
- *  @param[in] hostPtr Host Pointer allocated through hipHostMalloc
+ *  @param[in]  hostPtr Host Pointer allocated through hipHostMalloc
- *  @return Error code
+ *  @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)   ;
 

projects/hip/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,
+    hipErrorMissingConfiguration    = 1001,
-    hipErrorMemoryAllocation        = 2,    ///< Memory allocation error.
+    hipErrorMemoryAllocation        = 1002,    ///< Memory allocation error.
-    hipErrorInitializationError     = 3,    ///< TODO comment from hipErrorInitializationError
+    hipErrorInitializationError     = 1003,    ///< TODO comment from hipErrorInitializationError
-    hipErrorLaunchFailure           = 4,
+    hipErrorLaunchFailure           = 1004,    ///< An exception occurred on the device while executing a kernel.
-    hipErrorPriorLaunchFailure      = 5,
+    hipErrorPriorLaunchFailure      = 1005,
-    hipErrorLaunchTimeOut           = 6,
+    hipErrorLaunchTimeOut           = 1006,
-    hipErrorLaunchOutOfResources    = 7,    ///< Out of resources error.
+    hipErrorLaunchOutOfResources    = 1007,    ///< Out of resources error.
-    hipErrorInvalidDeviceFunction   = 8,
+    hipErrorInvalidDeviceFunction   = 1008,
-    hipErrorInvalidConfiguration    = 9,
+    hipErrorInvalidConfiguration    = 1009,
-    hipErrorInvalidDevice           = 10,   ///< DeviceID must be in range 0...#compute-devices.
+    hipErrorInvalidDevice           = 1010,   ///< 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.
+    hipErrorInvalidValue            = 1011,   ///< One or more of the parameters passed to the API call is NULL or not in an acceptable range.
-    hipErrorInvalidDevicePointer    = 17,   ///< Invalid Device Pointer
+    hipErrorInvalidDevicePointer    = 1017,   ///< Invalid Device Pointer
-    hipErrorInvalidMemcpyDirection  = 21,   ///< Invalid memory copy direction 
+    hipErrorInvalidMemcpyDirection  = 1021,   ///< Invalid memory copy direction 
-    hipErrorUnknown                 = 30,   ///< Unknown error.
+    hipErrorUnknown                 = 1030,   ///< Unknown error.
-    hipErrorInvalidResourceHandle   = 33,   ///< Resource handle (hipEvent_t or hipStream_t) invalid.
+    hipErrorInvalidResourceHandle   = 1033,   ///< 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.
+    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                = 38,   ///< Call to hipGetDeviceCount returned 0 devices
+    hipErrorNoDevice                = 1038,   ///< Call to hipGetDeviceCount returned 0 devices
-    hipErrorPeerAccessAlreadyEnabled = 50,  ///< Peer access was already enabled from the current device.
+    hipErrorPeerAccessAlreadyEnabled = 1050,  ///< Peer access was already enabled from the current device.
 
-    hipErrorPeerAccessNotEnabled    = 51,   ///< Peer access was never enabled from the current device.
+    hipErrorPeerAccessNotEnabled    = 1051,   ///< Peer access was never enabled from the current device.
-    hipErrorRuntimeMemory,                  ///< HSA runtime memory call returned error.  Typically not seen in production systems.
+    hipErrorRuntimeMemory           = 1052,                  ///< 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.
+    hipErrorRuntimeOther            = 1053,                   ///< 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.
+    hipErrorHostMemoryAlreadyRegistered = 1061, ///< Produced when trying to lock a page-locked memory.
-    hipErrorHostMemoryNotRegistered = 62,   ///< Produced when trying to unlock a non-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;
 

projects/hip/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)));
 }

projects/hip/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<N; i+=stride) {
-		C_d[i] = A_d[i] * A_d[i];
+        C_d[i] = A_d[i] * A_d[i];
-	}
+    }
 }
 
 
 int main(int argc, char *argv[])
 {
-	float *A_d, *C_d;
+    float *A_d, *C_d;
-	float *A_h, *C_h;
+    float *A_h, *C_h;
-	size_t N = 1000000;
+    size_t N = 1000000;
-	size_t Nbytes = N * sizeof(float);
+    size_t Nbytes = N * sizeof(float);
 
-	hipDeviceProp_t props;
+    hipDeviceProp_t props;
-	CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/));
+    CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/));
-	printf ("info: running on device %s\n", props.name);
+    printf ("info: running on device %s\n", props.name);
 
-	printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
+    printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
-	A_h = (float*)malloc(Nbytes);
+    A_h = (float*)malloc(Nbytes);
-	CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
+    CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
-	C_h = (float*)malloc(Nbytes);
+    C_h = (float*)malloc(Nbytes);
-	CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
+    CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
 	// Fill with Phi + i
     for (size_t i=0; i<N; i++) 
-	{
+    {
-		A_h[i] = 1.618f + i; 
+        A_h[i] = 1.618f + i; 
-	}
+    }
 
-	printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
+    printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
-	CHECK(hipMalloc(&A_d, Nbytes));
+    CHECK(hipMalloc(&A_d, Nbytes));
-	CHECK(hipMalloc(&C_d, Nbytes));
+    CHECK(hipMalloc(&C_d, Nbytes));
 
-
+    printf ("info: copy Host2Device\n");
-	printf ("info: copy Host2Device\n");
     CHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
 
-	const unsigned blocks = 512;
+    const unsigned blocks = 512;
-	const unsigned threadsPerBlock = 256;
+    const unsigned threadsPerBlock = 256;
 
-	printf ("info: launch 'vector_square' kernel\n");
+    printf ("info: launch 'vector_square' kernel\n");
-	hipLaunchKernel(HIP_KERNEL_NAME(vector_square), dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
+    hipLaunchKernel(HIP_KERNEL_NAME(vector_square), dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
 
-	printf ("info: copy Device2Host\n");
+    printf ("info: copy Device2Host\n");
     CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
 
-	printf ("info: check result\n");
+    printf ("info: check result\n");
     for (size_t i=0; i<N; i++)  {
-		if (C_h[i] != A_h[i] * A_h[i]) {
+        if (C_h[i] != A_h[i] * A_h[i]) {
-			CHECK(hipErrorUnknown);
+            CHECK(hipErrorUnknown);
-		}
+        }
-	}
+    }
-	printf ("PASSED!\n");
+    printf ("PASSED!\n");
 }

projects/hip/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;
         }

projects/hip/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);

projects/hip/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);

projects/hip/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 hipErrorInvalidDeviceFunction      : return "hipErrorInvalidDeviceFunction";
-        case hipErrorInvalidResourceHandle      : return "hipErrorInvalidResourceHandle";
+        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";
     };

projects/hip/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 <typename T>
 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)
 {

projects/hip/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
+# We are going to use HIP source...
-if(NOT DEFINED HIP_PATH)
+get_filename_component(HIP_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/../.. ABSOLUTE)
-        if(NOT DEFINED ENV{HIP_PATH})
+
-        # We are going to use HIP source...
+# ...so we first need to determine the options to cascade to HIP build.
-        get_filename_component(HIP_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/../.. ABSOLUTE)
+if(DEFINED HIP_PLATFORM)
-        execute_process(
+    set(ENV{HIP_PLATFORM} ${HIP_PLATFORM})
-                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()
-else()
-    # We are using HIP_PATH passed to cmake. So just create a fake target.
-    add_custom_target(hip ALL)
 endif()
-MESSAGE("HIP_PATH=" ${HIP_PATH})
+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
+    )
 
 # 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.
+    #Need C++11 for threads in some of the tests.
-	add_definitions(-std=c++11)
+    add_definitions(-std=c++11)
 
-	# NVCC does not not support -rdynamic option
+    # NVCC does not not support -rdynamic option
-	set(CMAKE_SHARED_LIBRARY_LINK_CXX_FLAGS )
+    set(CMAKE_SHARED_LIBRARY_LINK_CXX_FLAGS)
-	set(CMAKE_SHARED_LIBRARY_LINK_C_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)
+# usage : build_hip_executable(exe_name CPP_FILES)
-macro (build_hip_executable exe cpp)
+macro(build_hip_executable exe cpp)
-    add_executable (${exe} ${cpp} ${ARGN} $<TARGET_OBJECTS:test_common> )
+    add_executable(${exe} ${cpp} ${ARGN} $<TARGET_OBJECTS:test_common>)
     add_dependencies(${exe} hip)
 endmacro()
 
 # Make a hip executable, using libc++
-macro (build_hip_executable_libcpp exe cpp)
+macro(build_hip_executable_libcpp exe cpp)
-    build_hip_executable( ${exe} ${cpp} ${ARGN} )
+    build_hip_executable( ${exe} ${cpp} ${ARGN})
-    if (${HIP_PLATFORM} STREQUAL "hcc")
+    if(${HIP_PLATFORM} STREQUAL "hcc")
-        set_source_files_properties (${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --stdlib=libc++ )
+        set_source_files_properties(${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --stdlib=libc++)
     endif()
 endmacro()
 
-function (make_named_test exe testname )
+function(make_named_test exe testname)
-  add_test (NAME ${testname}
+    add_test(NAME ${testname}
-      COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN}
+        COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN}
-            )
+        )
-  set_tests_properties (${testname}
+    set_tests_properties(${testname}
-      PROPERTIES PASS_REGULAR_EXPRESSION "PASSED"
+        PROPERTIES PASS_REGULAR_EXPRESSION "PASSED"
-    )
+        )
 endfunction()
 
-macro (make_test exe )
+macro(make_test exe)
-  string (REPLACE " " "" smush_args ${ARGN})
+    string(REPLACE " " "" smush_args ${ARGN})
-  set (testname ${PROJECT_NAME}/${exe}${smush_args}.tst)
+    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 )
+macro(make_hipify_test sourceFile)
-    #string (REPLACE " " "" smush_args ${ARGN})
+    #string(REPLACE " " "" smush_args ${ARGN})
-  set (testname ${sourceFile}${smush_args}.tst)
+    set(testname ${sourceFile}${smush_args}.tst)
 
-  add_test (NAME ${testname}
+    add_test(NAME ${testname}
-      COMMAND ${HIP_PATH}/bin/hipify ${PROJECT_SOURCE_DIR}/${sourceFile} ${ARGN} 
+        COMMAND ${HIP_PATH}/bin/hipify ${PROJECT_SOURCE_DIR}/${sourceFile} ${ARGN}
-            )
+        )
 endmacro()
 
 
-macro (make_test_matches exe match_string)
+macro(make_test_matches exe match_string)
-  string (REPLACE " " "" smush_args ${ARGN})
+    string(REPLACE " " "" smush_args ${ARGN})
-  set (testname ${exe}${smush_args}.tst)
+    set(testname ${exe}${smush_args}.tst)
-  add_test (NAME ${testname}
+    add_test(NAME ${testname}
-      COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN}
+        COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN}
-            )
+        )
-  set_tests_properties (${testname}
+    set_tests_properties(${testname}
-      PROPERTIES PASS_REGULAR_EXPRESSION ${match_string}
+        PROPERTIES PASS_REGULAR_EXPRESSION ${match_string}
-    )
+        )
 endmacro()
 
-macro (build_hip_executable_sm35 exe cpp)
+macro(build_hip_executable_sm35 exe cpp)
-  build_hip_executable( ${exe} ${cpp} ${ARGN} )
+    build_hip_executable( ${exe} ${cpp} ${ARGN})
-  if (${HIP_PLATFORM} STREQUAL "nvcc")
+    if(${HIP_PLATFORM} STREQUAL "nvcc")
-    set_source_files_properties (${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --gpu-architecture=sm_35 ) 
+        set_source_files_properties(${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --gpu-architecture=sm_35)
-  endif()
+    endif()
 endmacro()
 
 
-build_hip_executable (hipGetDeviceAttribute hipGetDeviceAttribute.cpp) 
+build_hip_executable(hipGetDeviceAttribute hipGetDeviceAttribute.cpp)
-build_hip_executable (hipEnvVar hipEnvVar.cpp)
+build_hip_executable(hipEnvVar hipEnvVar.cpp)
-build_hip_executable (hipEnvVarDriver hipEnvVarDriver.cpp) 
+build_hip_executable(hipEnvVarDriver hipEnvVarDriver.cpp)
-build_hip_executable (hipEventRecord hipEventRecord.cpp)
+build_hip_executable(hipEventRecord hipEventRecord.cpp)
 
-build_hip_executable_libcpp (hipHcc hipHcc.cpp)
+build_hip_executable_libcpp(hipHcc hipHcc.cpp)
-#set_source_files_properties (hipHcc.cpp PROPERTIES COMPILE_FLAGS --stdlib=libc++ )
+#set_source_files_properties(hipHcc.cpp PROPERTIES COMPILE_FLAGS --stdlib=libc++)
 
 # __workweek fix.
-#build_hip_executable_libcpp (hipPointerAttrib hipPointerAttrib.cpp)
+#build_hip_executable_libcpp(hipPointerAttrib hipPointerAttrib.cpp)
-build_hip_executable (hipHostAlloc hipHostAlloc.cpp)
+build_hip_executable(hipHostAlloc hipHostAlloc.cpp)
-build_hip_executable (hipHostGetFlags hipHostGetFlags.cpp)
+build_hip_executable(hipHostGetFlags hipHostGetFlags.cpp)
-build_hip_executable (hipHostRegister hipHostRegister.cpp)
+build_hip_executable(hipHostRegister hipHostRegister.cpp)
-build_hip_executable (hipRandomMemcpyAsync hipRandomMemcpyAsync.cpp)
+build_hip_executable(hipRandomMemcpyAsync hipRandomMemcpyAsync.cpp)
-build_hip_executable (hipFuncSetDeviceFlags hipFuncSetDeviceFlags.cpp)
+build_hip_executable(hipFuncSetDeviceFlags hipFuncSetDeviceFlags.cpp)
-build_hip_executable (hipFuncGetDevice hipFuncGetDevice.cpp)
+build_hip_executable(hipFuncGetDevice hipFuncGetDevice.cpp)
-build_hip_executable (hipFuncSetDevice hipFuncSetDevice.cpp)
+build_hip_executable(hipFuncSetDevice hipFuncSetDevice.cpp)
-build_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp)
+build_hip_executable(hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp)
-build_hip_executable (hipPeerToPeer_simple hipPeerToPeer_simple.cpp)
+build_hip_executable(hipPeerToPeer_simple hipPeerToPeer_simple.cpp)
-build_hip_executable (hipTestMemcpyPin hipTestMemcpyPin.cpp)
+build_hip_executable(hipTestMemcpyPin hipTestMemcpyPin.cpp)
-build_hip_executable (hipDynamicShared hipDynamicShared.cpp)
+build_hip_executable(hipDynamicShared hipDynamicShared.cpp)
-build_hip_executable (hipLaunchParm hipLaunchParm.cpp)
+build_hip_executable(hipLaunchParm hipLaunchParm.cpp)
 
-if (${HIP_PLATFORM} STREQUAL "hcc")
+if(${HIP_PLATFORM} STREQUAL "hcc")
-    build_hip_executable (hipArray hipArray.cpp)
+    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

projects/hip/tests/src/hipDrvMemcpy.cpp

@@ -6,21 +6,39 @@
 #define SIZE LEN<<2
 
 int main(){
-  int *A, *B, *C;
+  int *A, *B;
-  hipDeviceptr Ad, Bd;
+  hipDeviceptr_t Ad, Bd;
   A = new int[LEN];
   B = new int[LEN];
-  C = new int[LEN];
+
   for(int i=0;i<LEN;i++){
     A[i] = i;
   }
+
   hipMalloc((void**)&Ad, SIZE);
   hipMalloc((void**)&Bd, SIZE);
+
   hipMemcpyHtoD(Ad, A, SIZE);
   hipMemcpyDtoD(Bd, Ad, SIZE);
   hipMemcpyDtoH(B, Bd, SIZE);
-  hipMemcpyHtoH(C, B,SIZE);
+
   for(int i=0;i<16;i++){
-    std::cout<<A[i]<<" "<<C[i]<<std::endl;
+    std::cout<<A[i]<<" "<<B[i]<<std::endl;
   }
+
+  int *Ah, *Bh;
+  hipHostMalloc(&Ah, SIZE, 0);
+  hipHostMalloc(&Bh, SIZE, 0);
+  memcpy(Ah, A, SIZE);
+  hipStream_t stream;
+  hipStreamCreate(&stream);
+
+  hipMemcpyHtoDAsync(Ad, Ah, SIZE, stream);
+  hipStreamSynchronize(stream);
+  hipMemcpyDtoDAsync(Bd, Ad, SIZE, stream);
+  hipStreamSynchronize(stream);
+  hipMemcpyDtoHAsync(Bh, Bd, SIZE, stream);
+  hipStreamSynchronize(stream);
+
+  std::cout<<Ah[10]<<" "<<Bh[10]<<std::endl;
 }