Merge 'master' into 'amd-master'

Change-Id: I4f54ac4fd1993ab9dfca9aa34a153e9eb86f45ae

hipify-clang/CMakeLists.txt

@@ -74,12 +74,12 @@ if ((LLVM_PACKAGE_VERSION VERSION_EQUAL "7") OR (LLVM_PACKAGE_VERSION VERSION_GR
 endif()
 
 set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}")
-set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}")
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS} -std=c++14")
 if(MSVC)
     target_compile_options(hipify-clang PRIVATE "/Od /GR- /EHs- /EHc-")
     set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS")
 else()
-    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden")
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -pthread -fno-rtti -fvisibility-inlines-hidden")
 endif()
 
 set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"")

hipify-clang/src/main.cpp

@@ -426,6 +426,7 @@ int main(int argc, const char **argv) {
     if (Tool.runAndSave(&actionFactory)) {
       currentStat.hasErrors = true;
       LLVM_DEBUG(llvm::dbgs() << "Skipped some replacements.\n");
+      Result = 1;
     }
     // Copy the tmpfile to the output
     if (!NoOutput && !currentStat.hasErrors) {

include/hip/hcc_detail/code_object_bundle.hpp

@@ -91,6 +91,10 @@ struct Bundled_code {
 
 #define magic_string_  "__CLANG_OFFLOAD_BUNDLE__"
 
+#ifdef __GNUC__
+#pragma GCC visibility push (default)
+#endif
+
 class Bundled_code_header {
     // DATA - STATICS
     static constexpr auto magic_string_sz_ = sizeof(magic_string_) - 1;
@@ -172,6 +176,10 @@ class Bundled_code_header {
     size_t bundled_code_size = 0;
 };
 
+#ifdef __GNUC__
+#pragma GCC visibility pop
+#endif
+
 // CREATORS
 template <typename RandomAccessIterator>
 Bundled_code_header::Bundled_code_header(RandomAccessIterator f, RandomAccessIterator l)

include/hip/hcc_detail/device_library_decls.h

@@ -59,6 +59,10 @@ extern "C" __device__ __attribute__((const)) float __ocml_trunc_f32(float);
 extern "C" __device__ __attribute__((const)) float __ocml_fmin_f32(float, float);
 extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);
 
+extern "C" __device__ __attribute__((convergent)) void __ockl_gws_init(uint nwm1, uint rid);
+extern "C" __device__ __attribute__((convergent)) void __ockl_gws_barrier(uint nwm1, uint rid);
+
+
 // Introduce local address space
 #define __local __attribute__((address_space(3)))
 

include/hip/hcc_detail/helpers.hpp

@@ -89,7 +89,7 @@ template <typename, typename = void>
 struct is_callable_impl : std::false_type {};
 
 template <FunctionalProcedure F, typename... Ts>
-struct is_callable_impl<F(Ts...), void_t_<std::result_of_t<F(Ts...)> > > : std::true_type {};
+struct is_callable_impl<F(Ts...), void_t_<typename> std::result_of<F>::type(Ts...) > > : std::true_type {};
 #else
 
 // C++17

samples/1_Utils/hipInfo/hipInfo.cpp

@@ -83,6 +83,7 @@ void printDeviceProp(int deviceId) {
     cout << setw(w1) << "Name: " << props.name << endl;
     cout << setw(w1) << "pciBusID: " << props.pciBusID << endl;
     cout << setw(w1) << "pciDeviceID: " << props.pciDeviceID << endl;
+    cout << setw(w1) << "pciDomainID: " << props.pciDomainID << endl;
     cout << setw(w1) << "multiProcessorCount: " << props.multiProcessorCount << endl;
     cout << setw(w1) << "maxThreadsPerMultiProcessor: " << props.maxThreadsPerMultiProcessor
          << endl;
@@ -100,6 +101,7 @@ void printDeviceProp(int deviceId) {
     cout << setw(w1) << "totalConstMem: " << props.totalConstMem << endl;
     cout << setw(w1) << "sharedMemPerBlock: " << (float)props.sharedMemPerBlock / 1024.0 << " KB"
          << endl;
+    cout << setw(w1) << "canMapHostMemory: " << props.canMapHostMemory << endl;
     cout << setw(w1) << "regsPerBlock: " << props.regsPerBlock << endl;
     cout << setw(w1) << "warpSize: " << props.warpSize << endl;
     cout << setw(w1) << "l2CacheSize: " << props.l2CacheSize << endl;
@@ -114,6 +116,8 @@ void printDeviceProp(int deviceId) {
     cout << setw(w1) << "major: " << props.major << endl;
     cout << setw(w1) << "minor: " << props.minor << endl;
     cout << setw(w1) << "concurrentKernels: " << props.concurrentKernels << endl;
+    cout << setw(w1) << "cooperativeLaunch: " << props.cooperativeLaunch << endl;
+    cout << setw(w1) << "cooperativeMultiDeviceLaunch: " << props.cooperativeMultiDeviceLaunch << endl;
     cout << setw(w1) << "arch.hasGlobalInt32Atomics: " << props.arch.hasGlobalInt32Atomics << endl;
     cout << setw(w1) << "arch.hasGlobalFloatAtomicExch: " << props.arch.hasGlobalFloatAtomicExch
          << endl;
@@ -134,7 +138,14 @@ void printDeviceProp(int deviceId) {
     cout << setw(w1) << "arch.has3dGrid: " << props.arch.has3dGrid << endl;
     cout << setw(w1) << "arch.hasDynamicParallelism: " << props.arch.hasDynamicParallelism << endl;
     cout << setw(w1) << "gcnArch: " << props.gcnArch << endl;
-
+    cout << setw(w1) << "isIntegrated: " << props.integrated << endl;
+    cout << setw(w1) << "maxTexture1D: " << props.maxTexture1D << endl;
+    cout << setw(w1) << "maxTexture2D.width: " << props.maxTexture2D[0] << endl;
+    cout << setw(w1) << "maxTexture2D.height: " << props.maxTexture2D[1] << endl;
+    cout << setw(w1) << "maxTexture3D.width: " << props.maxTexture3D[0] << endl;
+    cout << setw(w1) << "maxTexture3D.height: " << props.maxTexture3D[1] << endl;
+    cout << setw(w1) << "maxTexture3D.depth: " << props.maxTexture3D[2] << endl;
+    
     int deviceCnt;
     hipGetDeviceCount(&deviceCnt);
     cout << setw(w1) << "peers: ";
@@ -192,6 +203,7 @@ int main(int argc, char* argv[]) {
     HIPCHECK(hipGetDeviceCount(&deviceCnt));
 
     for (int i = 0; i < deviceCnt; i++) {
+        hipSetDevice(i);
         printDeviceProp(i);
     }
 

src/hip_memory.cpp

@@ -2156,46 +2156,64 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) {
 hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags) {
     HIP_INIT_API(hipIpcOpenMemHandle, devPtr, &handle, flags);
     hipError_t hipStatus = hipSuccess;
-    if (devPtr == NULL) {
-        hipStatus = hipErrorInvalidValue;
-    } else {
-#if USE_IPC
-        // Get the current device agent.
-        hc::accelerator acc;
-        hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
-        if (!agent) return hipErrorInvalidResourceHandle;
+    if (devPtr == NULL)
+        return ihipLogStatus(hipErrorInvalidValue);
 
-        ihipIpcMemHandle_t* iHandle = (ihipIpcMemHandle_t*)&handle;
-        // Attach ipc memory
-        auto ctx = ihipGetTlsDefaultCtx();
-        {
-            LockedAccessor_CtxCrit_t crit(ctx->criticalData());
-            // the peerCnt always stores self so make sure the trace actually
-            hsa_status_t hsa_status = hsa_amd_ipc_memory_attach(
-                (hsa_amd_ipc_memory_t*)&(iHandle->ipc_handle), iHandle->psize, crit->peerCnt(),
-                crit->peerAgents(), devPtr);
-            if (hsa_status != HSA_STATUS_SUCCESS) hipStatus = hipErrorMapBufferObjectFailed;
-        }
+#if USE_IPC
+    // Get the current device agent.
+    hc::accelerator acc;
+    hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
+    if (!agent)
+        return ihipLogStatus(hipErrorInvalidResourceHandle);
+
+    ihipIpcMemHandle_t* iHandle = (ihipIpcMemHandle_t*)&handle;
+    // Attach ipc memory
+    auto ctx = ihipGetTlsDefaultCtx();
+    {
+        LockedAccessor_CtxCrit_t crit(ctx->criticalData());
+        auto device = ctx->getWriteableDevice();
+        // the peerCnt always stores self so make sure the trace actually
+        if(hsa_amd_ipc_memory_attach(
+            (hsa_amd_ipc_memory_t*)&(iHandle->ipc_handle), iHandle->psize, crit->peerCnt(),
+             crit->peerAgents(), devPtr) != HSA_STATUS_SUCCESS)
+            return ihipLogStatus(hipErrorRuntimeOther);
+
+        hc::AmPointerInfo ampi(NULL, *devPtr, *devPtr, sizeof(*devPtr), acc, true, true);
+        am_status_t am_status = hc::am_memtracker_add(*devPtr,ampi);
+        if (am_status != AM_SUCCESS)
+            return ihipLogStatus(hipErrorMapBufferObjectFailed);
+
+#if USE_APP_PTR_FOR_CTX
+        am_status = hc::am_memtracker_update(*devPtr, device->_deviceId, 0, ctx);
 #else
-        hipStatus = hipErrorRuntimeOther;
+        am_status = hc::am_memtracker_update(*devPtr, device->_deviceId, 0);
 #endif
+        if(am_status != AM_SUCCESS)
+            return ihipLogStatus(hipErrorMapBufferObjectFailed);
     }
+#else
+    hipStatus = hipErrorRuntimeOther;
+#endif
+
     return ihipLogStatus(hipStatus);
 }
 
 hipError_t hipIpcCloseMemHandle(void* devPtr) {
     HIP_INIT_API(hipIpcCloseMemHandle, devPtr);
     hipError_t hipStatus = hipSuccess;
-    if (devPtr == NULL) {
-        hipStatus = hipErrorInvalidValue;
-    } else {
+    if (devPtr == NULL)
+        return ihipLogStatus(hipErrorInvalidValue);
+
 #if USE_IPC
-        hsa_status_t hsa_status = hsa_amd_ipc_memory_detach(devPtr);
-        if (hsa_status != HSA_STATUS_SUCCESS) return hipErrorInvalidResourceHandle;
+    if(hc::am_memtracker_remove(devPtr) != AM_SUCCESS)
+        return ihipLogStatus(hipErrorInvalidValue);
+
+    if (hsa_amd_ipc_memory_detach(devPtr) != HSA_STATUS_SUCCESS)
+        return ihipLogStatus(hipErrorInvalidResourceHandle);
 #else
-        hipStatus = hipErrorRuntimeOther;
+    hipStatus = hipErrorRuntimeOther;
 #endif
-    }
+
     return ihipLogStatus(hipStatus);
 }
 

src/hip_module.cpp

@@ -386,6 +386,100 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
     return ihipLogStatus(result);
 }
 
+namespace {
+// kernel for initializing GWS
+// nwm1 is the total number of work groups minus 1
+__global__ void init_gws(uint nwm1) {
+    __ockl_gws_init(nwm1, 0);
+}
+}
+
+hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim,
+        dim3 blockDimX, void** kernelParams, unsigned int sharedMemBytes,
+        hipStream_t stream) {
+
+    HIP_INIT_API(hipLaunchCooperativeKernel, f, gridDim, blockDimX, kernelParams, sharedMemBytes, stream);
+    hipError_t result;
+
+
+    if ((f == nullptr) || (stream == nullptr) || (kernelParams == nullptr)) {
+        return ihipLogStatus(hipErrorNotInitialized);
+    }
+
+    if (!stream->getDevice()->_props.cooperativeLaunch) {
+        return ihipLogStatus(hipErrorInvalidConfiguration);
+    }
+
+    // Prepare the kernel descriptor for initializing the GWS
+    hipFunction_t gwsKD = hip_impl::get_program_state().kernel_descriptor(
+            reinterpret_cast<std::uintptr_t>(&init_gws),
+            hip_impl::target_agent(stream));
+
+    if (gwsKD == nullptr) {
+        return ihipLogStatus(hipErrorInvalidValue);
+    }
+    hip_impl::kernargs_size_align gwsKargs =
+            hip_impl::get_program_state().get_kernargs_size_align(
+                    reinterpret_cast<std::uintptr_t>(&init_gws));
+
+    gwsKD->_kernarg_layout = *reinterpret_cast<const std::vector<
+            std::pair<std::size_t, std::size_t>>*>(gwsKargs.getHandle());
+
+    // Prepare the kernel descriptor for the main kernel
+    hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor(
+            reinterpret_cast<std::uintptr_t>(f),
+            hip_impl::target_agent(stream));
+    if (kd == nullptr) {
+        return ihipLogStatus(hipErrorInvalidValue);
+    }
+    hip_impl::kernargs_size_align kargs =
+            hip_impl::get_program_state().get_kernargs_size_align(
+                    reinterpret_cast<std::uintptr_t>(f));
+
+    kd->_kernarg_layout = *reinterpret_cast<const std::vector<
+            std::pair<std::size_t, std::size_t>>*>(kargs.getHandle());
+
+
+    void *gwsKernelParam[1];
+    // calculate total number of work groups minus 1 for the main kernel
+    uint nwm1 = (gridDim.x * gridDim.y * gridDim.z) - 1;
+    gwsKernelParam[0] = &nwm1;
+
+    LockedAccessor_StreamCrit_t streamCrit(stream->criticalData(), false);
+#if (__hcc_workweek__ >= 19213)
+    streamCrit->_av.acquire_locked_hsa_queue();
+#endif
+
+    // launch the init_gws kernel to initialize the GWS
+    result = ihipModuleLaunchKernel(tls, gwsKD, 1, 1, 1, 1, 1, 1,
+             0, stream, gwsKernelParam, nullptr, nullptr, nullptr, 0, true);
+
+    if (result != hipSuccess) {
+        stream->criticalData().unlock();
+#if (__hcc_workweek__ >= 19213)
+        stream->criticalData()._av.release_locked_hsa_queue();
+#endif
+
+        return ihipLogStatus(hipErrorLaunchFailure);
+    }
+
+    // launch the main kernel
+    result = ihipModuleLaunchKernel(tls, kd,
+            gridDim.x * blockDimX.x,
+            gridDim.y * blockDimX.y,
+            gridDim.z * blockDimX.z,
+            blockDimX.x, blockDimX.y, blockDimX.z,
+            sharedMemBytes, stream, kernelParams, nullptr, nullptr,
+            nullptr, 0, true);
+
+    stream->criticalData().unlock();
+#if (__hcc_workweek__ >= 19213)
+    stream->criticalData()._av.release_locked_hsa_queue();
+#endif
+
+    return ihipLogStatus(result);
+}
+
 namespace hip_impl {
     hsa_executable_t executable_for(hipModule_t hmod) {
         return hmod->executable;