Merge branch 'amd-develop' into amd-master

Change-Id: I9f284502dc764b7ad39bd9bc8a40ea7404d0391e
ROCm · Sep 16, 2016 · 8ac4a51 · 8ac4a51
2 parents f0adf7e + 22a7841
commit 8ac4a51
Show file tree

Hide file tree

Showing 16 changed files with 745 additions and 374 deletions.
diff --git a/.vimrc 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
@@ -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
@@ -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
@@ -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
@@ -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
 /**
  * @}
  */