Merge branch 'amd-master-next' into amd-npi-next

Change-Id: I53f9c054e5df19bd8b130e7610d71b91d525474a

CMakeLists.txt

@@ -8,6 +8,10 @@ set(BUILD_SHARED_LIBS ON  CACHE BOOL "Build shared library (.so) or static lib (
 
 set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
 
+if(NOT ${BUILD_SHARED_LIBS} AND NOT DEFINED ENABLE_HIP_PCH)
+  set(ENABLE_HIP_PCH ON CACHE BOOL "enable/disable pre-compiled hip headers")
+endif()
+
 #############################
 # Options
 #############################

bin/hip_embed_pch.sh

@@ -0,0 +1,58 @@
+#!/bin/bash
+
+#set -x
+
+ROCM_PATH=${ROCM_PATH:-/opt/rocm}
+tmp=/tmp/hip_pch.$$
+mkdir -p $tmp
+
+cat >$tmp/hip_macros.h <<EOF
+#define __device__ __attribute__((device))
+#define __host__ __attribute__((host))
+#define __global__ __attribute__((global))
+#define __constant__ __attribute__((constant))
+#define __shared__ __attribute__((shared))
+
+#define launch_bounds_impl0(requiredMaxThreadsPerBlock)                                            \
+    __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock)))
+#define launch_bounds_impl1(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor)                \
+    __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock),                     \
+                   amdgpu_waves_per_eu(minBlocksPerMultiprocessor)))
+#define select_impl_(_1, _2, impl_, ...) impl_
+#define __launch_bounds__(...)                                                                     \
+    select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__)
+
+// Macro to replace extern __shared__ declarations
+// to local variable definitions
+#define HIP_DYNAMIC_SHARED(type, var) \
+    type* var = (type*)__amdgcn_get_dynamicgroupbaseptr();
+EOF
+
+cat >$tmp/hip_pch.h <<EOF
+#include "hip/hip_runtime.h"
+#include "hip/hip_fp16.h"
+EOF
+
+
+cat >$tmp/hip_pch.mcin <<EOF
+ .type __hip_pch,@object
+  .section .hip_pch,"aMS",@progbits,1
+  .data
+  .globl __hip_pch
+  .globl __hip_pch_size
+  .p2align 3
+__hip_pch:
+  .incbin "$tmp/hip.pch"
+__hip_pch_size:
+  .long __hip_pch_size - __hip_pch
+EOF
+
+$ROCM_PATH/llvm/bin/clang -O3 -c -std=c++17 -isystem /opt/rocm/llvm/lib/clang/11.0.0/include/.. -isystem /opt/rocm/include -nogpulib --cuda-device-only -x hip $tmp/hip_pch.h -E >$tmp/pch.cui
+
+cat $tmp/hip_macros.h >> $tmp/pch.cui
+
+$ROCM_PATH/llvm/bin/clang -cc1 -O3 -emit-pch -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -std=c++17 -fgnuc-version=4.2.1 -o $tmp/hip.pch -x hip-cpp-output - <$tmp/pch.cui
+
+$ROCM_PATH/llvm/bin/llvm-mc -o hip_pch.o $tmp/hip_pch.mcin --filetype=obj
+
+rm -rf $tmp

bin/hip_gen_pch.sh

@@ -0,0 +1,36 @@
+#!/bin/bash
+
+#set -x
+
+cat >/tmp/hip_macros.h <<EOF
+#define __device__ __attribute__((device))
+#define __host__ __attribute__((host))
+#define __global__ __attribute__((global))
+#define __constant__ __attribute__((constant))
+#define __shared__ __attribute__((shared))
+
+#define launch_bounds_impl0(requiredMaxThreadsPerBlock)                                            \
+    __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock)))
+#define launch_bounds_impl1(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor)                \
+    __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock),                     \
+                   amdgpu_waves_per_eu(minBlocksPerMultiprocessor)))
+#define select_impl_(_1, _2, impl_, ...) impl_
+#define __launch_bounds__(...)                                                                     \
+    select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__)
+
+// Macro to replace extern __shared__ declarations
+// to local variable definitions
+#define HIP_DYNAMIC_SHARED(type, var) \
+    type* var = (type*)__amdgcn_get_dynamicgroupbaseptr();
+EOF
+
+cat >/tmp/hip_pch.h <<EOF
+#include "hip/hip_runtime.h"
+#include "hip/hip_fp16.h"
+EOF
+
+/opt/rocm/llvm/bin/clang -O3 -c -std=c++17 -isystem /opt/rocm/llvm/lib/clang/11.0.0/include/.. -isystem /opt/rocm/include -nogpulib --cuda-device-only -x hip /tmp/hip_pch.h -E >/tmp/pch.cui
+
+cat /tmp/hip_macros.h >> /tmp/pch.cui
+
+/opt/rocm/llvm/bin/clang -cc1 -O3 -emit-pch -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -std=c++17 -fgnuc-version=4.2.1 -o /tmp/hip.pch -x hip-cpp-output - </tmp/pch.cui

bin/hipcc

@@ -73,11 +73,11 @@ $HIPCC_LINK_FLAGS_APPEND=$ENV{'HIPCC_LINK_FLAGS_APPEND'};
 # derive HIP_PATH, as dirname $0 could be /opt/rocm/bin or /opt/rocm/hip/bin
 # depending on how it gets invoked.
 # ROCM_PATH which points to <rocm_install_dir> is determined based on whether
-# we find .info/version in the parent of HIP_PATH or not. If it is found,
+# we find bin/rocminfo in the parent of HIP_PATH or not. If it is found,
 # ROCM_PATH is defined relative to HIP_PATH else it is hardcoded to /opt/rocm.
 #
 $HIP_PATH=$ENV{'HIP_PATH'} // dirname(Cwd::abs_path("$0/../")); # use parent directory of hipcc
-if (-e "$HIP_PATH/../.info/version") {
+if (-e "$HIP_PATH/../bin/rocminfo") {
     $ROCM_PATH=$ENV{'ROCM_PATH'} // dirname("$HIP_PATH"); # use parent directory of HIP_PATH
 } else {
     $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";

include/hip/hcc_detail/hip_runtime_api.h

@@ -345,6 +345,12 @@ typedef struct hipLaunchParams_t {
     hipStream_t stream;     ///< Stream identifier
 } hipLaunchParams;
 
+// Pre-Compiled header for online compilation
+#ifdef ENABLE_HIP_PCH
+extern const char* __hip_pch;
+extern unsigned __hip_pch_size;
+void __hipGetPCH(const char** pch, unsigned int*size);
+#endif
 
 // Doxygen end group GlobalDefs
 /**  @} */

include/hip/hcc_detail/math_functions.h

@@ -1533,13 +1533,14 @@ inline _Float16 pow(_Float16 base, int iexp) {
   return __ocml_pown_f16(base, iexp);
 }
 
-#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
-
 #pragma pop_macro("__DEF_FLOAT_FUN")
 #pragma pop_macro("__DEF_FLOAT_FUN2")
 #pragma pop_macro("__DEF_FLOAT_FUN2I")
 #pragma pop_macro("__HIP_OVERLOAD")
 #pragma pop_macro("__HIP_OVERLOAD2")
+
+#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
+
 #pragma pop_macro("__DEVICE__")
 #pragma pop_macro("__RETURN_TYPE")
 

rocclr/CMakeLists.txt

@@ -145,7 +145,15 @@ if(ROCclr_FOUND)
       $<TARGET_PROPERTY:amdrocclr_static,INTERFACE_COMPILE_DEFINITIONS>)
 endif()
 
- # Enable profiling API
+# Short-Term solution for pre-compiled headers for online compilation
+# Enable pre compiled header
+if(${ENABLE_HIP_PCH})
+    execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_gen_pch.sh")
+    execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_embed_pch.sh")
+    add_definitions(-DENABLE_HIP_PCH)
+endif()
+
+# Enable profiling API
 if(USE_PROF_API EQUAL 1)
   find_path(PROF_API_HEADER_DIR prof_protocol.h
     HINTS
@@ -205,17 +213,21 @@ target_link_libraries(host INTERFACE hip::amdhip64)
 
 add_library(device INTERFACE)
 target_link_libraries(device INTERFACE host)
+
+# Short-Term solution for pre-compiled headers for online compilation
+if(${ENABLE_HIP_PCH})
+  target_link_libraries(amdhip64 PRIVATE ${CMAKE_BINARY_DIR}/hip_pch.o)
+endif()
+
 # TODO: we may create host_static and device_static to let app
 # link amdhip64_static
 
 # FIXME: Linux convention is to create static library with same base
 # filename.
-
 if(${BUILD_SHARED_LIBS})
     target_link_libraries(amdhip64 PRIVATE amdrocclr_static Threads::Threads dl hsa-runtime64::hsa-runtime64)
     INSTALL(PROGRAMS $<TARGET_FILE:amdhip64> DESTINATION lib COMPONENT MAIN)
 else()
-
     target_link_libraries(amdhip64 PRIVATE Threads::Threads dl hsa-runtime64::hsa-runtime64 amd_comgr)
     # combine objects of vid and hip into amdhip64_static
     add_custom_target(
@@ -228,9 +240,7 @@ else()
        DEPENDS amdhip64 amdrocclr_static # To make sure this is the last step
        COMMENT "Combining static libs into amdhip64_static"
     )
-
     INSTALL(PROGRAMS $<TARGET_FILE:amdhip64> DESTINATION lib COMPONENT MAIN)
-
 endif()
 
 INSTALL(TARGETS amdhip64 host device EXPORT hip-targets DESTINATION ${LIB_INSTALL_DIR})

rocclr/hip_device_runtime.cpp

@@ -551,29 +551,3 @@ hipError_t hipSetValidDevices ( int* device_arr, int  len ) {
 
   HIP_RETURN(hipErrorNotSupported);
 }
-
-hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount) {
-  HIP_INIT_API(hipExtGetLinkTypeAndHopCount, device1, device2, linktype, hopcount);
-
-  amd::Device* amd_dev_obj1 = nullptr;
-  amd::Device* amd_dev_obj2 = nullptr;
-  const int numDevices = static_cast<int>(g_devices.size());
-
-  if ((device1 < 0) || (device1 >= numDevices) || (device2 < 0) || (device2 >= numDevices)) {
-    HIP_RETURN(hipErrorInvalidDevice);
-  }
-
-  if ((linktype == nullptr) || (hopcount == nullptr)) {
-    HIP_RETURN(hipErrorInvalidValue);
-  }
-
-  amd_dev_obj1 = g_devices[device1]->devices()[0];
-  amd_dev_obj2 = g_devices[device2]->devices()[0];
-
-  if (!amd_dev_obj1->findLinkTypeAndHopCount(amd_dev_obj2, linktype, hopcount)) {
-    HIP_RETURN(hipErrorInvalidHandle);
-  }
-
-  HIP_RETURN(hipSuccess);
-}
-

rocclr/hip_global.cpp

@@ -5,6 +5,13 @@
 #include "hip_code_object.hpp"
 #include "platform/program.hpp"
 
+#ifdef ENABLE_HIP_PCH
+void __hipGetPCH(const char** pch, unsigned int *size) {
+  *pch = __hip_pch;
+  *size = __hip_pch_size;
+}
+#endif
+
 namespace hip {
 
 //Device Vars

rocclr/hip_hcc.map.in

@@ -274,6 +274,7 @@ global:
     hipMemcpyAtoH;
     hipMemcpyHtoA;
     hipMemcpyParam2DAsync;
+    __hipGetPCH;
     };
 local:
     *;

rocclr/hip_memory.cpp

@@ -2029,14 +2029,20 @@ hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsi
 hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr) {
   HIP_INIT_API(hipPointerGetAttributes, attributes, ptr);
 
+  if (attributes == nullptr || ptr == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
   size_t offset = 0;
   amd::Memory* memObj = getMemoryObject(ptr, offset);
   int device = 0;
+  memset(attributes, 0, sizeof(hipPointerAttribute_t));
 
   if (memObj != nullptr) {
     attributes->memoryType = (CL_MEM_SVM_FINE_GRAIN_BUFFER & memObj->getMemFlags())? hipMemoryTypeHost : hipMemoryTypeDevice;
-    attributes->hostPointer = memObj->getSvmPtr();
-    attributes->devicePointer = memObj->getSvmPtr();
+    if (attributes->memoryType == hipMemoryTypeHost) {
+      attributes->hostPointer = static_cast<char*>(memObj->getSvmPtr()) + offset;
+    }
+    attributes->devicePointer = static_cast<char*>(memObj->getSvmPtr()) + offset;
     attributes->isManaged = 0;
     attributes->allocationFlags = memObj->getMemFlags() >> 16;
 

rocclr/hip_peer.cpp

@@ -72,12 +72,48 @@ hipError_t canAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId){
   return hipSuccess;
 }
 
+hipError_t findLinkInfo(int device1, int device2,
+                        std::vector<amd::Device::LinkAttrType>* link_attrs) {
+
+  amd::Device* amd_dev_obj1 = nullptr;
+  amd::Device* amd_dev_obj2 = nullptr;
+  const int numDevices = static_cast<int>(g_devices.size());
+
+  if ((device1 < 0) || (device1 >= numDevices) || (device2 < 0) || (device2 >= numDevices)) {
+    return hipErrorInvalidDevice;
+  }
+
+  amd_dev_obj1 = g_devices[device1]->devices()[0];
+  amd_dev_obj2 = g_devices[device2]->devices()[0];
+
+  if (!amd_dev_obj1->findLinkInfo(*amd_dev_obj2, link_attrs)) {
+    return hipErrorInvalidHandle;
+  }
+
+  return hipSuccess;
+}
+
+hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2,
+                                        uint32_t* linktype, uint32_t* hopcount) {
+  HIP_INIT_API(hipExtGetLinkTypeAndHopCount, device1, device2, linktype, hopcount);
+
+  // Fill out the list of LinkAttributes
+  std::vector<amd::Device::LinkAttrType> link_attrs;
+  link_attrs.push_back(std::make_pair(amd::Device::LinkAttribute::kLinkLinkType, 0));
+  link_attrs.push_back(std::make_pair(amd::Device::LinkAttribute::kLinkHopCount, 0));
+
+  HIP_RETURN_ONFAIL(findLinkInfo(device1, device2, &link_attrs));
+
+  *linktype = static_cast<uint32_t>(link_attrs[0].second);
+  *hopcount = static_cast<uint32_t>(link_attrs[1].second);
+
+  HIP_RETURN(hipSuccess);
+}
+
 hipError_t hipDeviceGetP2PAttribute(int* value, hipDeviceP2PAttr attr,
                                     int srcDevice, int dstDevice) {
   HIP_INIT_API(hipDeviceGetP2PAttribute, value, attr, srcDevice, dstDevice);
 
-  hipError_t hip_error = hipSuccess;
-
   if (value == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
@@ -87,26 +123,50 @@ hipError_t hipDeviceGetP2PAttribute(int* value, hipDeviceP2PAttr attr,
     HIP_RETURN(hipErrorInvalidDevice);
   }
 
+  std::vector<amd::Device::LinkAttrType> link_attrs;
+
   switch (attr) {
-    case hipDevP2PAttrPerformanceRank :
-      assert(0 && "Unimplemented");
+    case hipDevP2PAttrPerformanceRank : {
+      link_attrs.push_back(std::make_pair(amd::Device::LinkAttribute::kLinkLinkType, 0));
       break;
-    case hipDevP2PAttrAccessSupported :
-      hip_error = canAccessPeer(value, srcDevice, dstDevice);
+    }
+    case hipDevP2PAttrAccessSupported : {
+      HIP_RETURN_ONFAIL(canAccessPeer(value, srcDevice, dstDevice));
       break;
-    case hipDevP2PAttrNativeAtomicSupported :
-      assert(0 && "Unimplemented");
+    }
+    case hipDevP2PAttrNativeAtomicSupported : {
+      link_attrs.push_back(std::make_pair(amd::Device::LinkAttribute::kLinkLinkType, 0));
       break;
-    case hipDevP2PAttrHipArrayAccessSupported :
-      assert(0 && "Unimplemented");
+    }
+    case hipDevP2PAttrHipArrayAccessSupported : {
+      hipDeviceProp_t srcDeviceProp;
+      hipDeviceProp_t dstDeviceProp;
+      HIP_RETURN_ONFAIL(hipGetDeviceProperties(&srcDeviceProp, srcDevice));
+      HIP_RETURN_ONFAIL(hipGetDeviceProperties(&dstDeviceProp, dstDevice));
+
+      // Linear layout access is supported if P2P is enabled
+      // Opaque Images are supported only on homogeneous systems
+      // Might have more conditions to check, in future.
+      if (srcDeviceProp.gcnArch == dstDeviceProp.gcnArch) {
+        HIP_RETURN_ONFAIL(canAccessPeer(value, srcDevice, dstDevice));
+      } else {
+        value = 0;
+      }
       break;
-    default :
+    }
+    default : {
       DevLogPrintfError("Invalid attribute attr: %d ", attr);
-      hip_error = hipErrorInvalidValue;
+      HIP_RETURN(hipErrorInvalidValue);
       break;
+    }
   }
 
-  HIP_RETURN(hip_error);
+  if ((attr != hipDevP2PAttrAccessSupported) && (attr != hipDevP2PAttrHipArrayAccessSupported)) {
+    HIP_RETURN_ONFAIL(findLinkInfo(srcDevice, dstDevice, &link_attrs));
+    *value = static_cast<int>(link_attrs[0].second);
+  }
+
+  HIP_RETURN(hipSuccess);
 }
 
 hipError_t hipDeviceCanAccessPeer(int* canAccess, int deviceId, int peerDeviceId) {