Merge 'master' into 'amd-master'

Change-Id: Icb2a5ad71466d9c3478f4ffd9ac96173284dcf23

bin/hipcc

@@ -116,7 +116,7 @@ if ($HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) {
 }
 
 if (defined $HIP_VDI_HOME) {
-    if (!defined $HIP_CLANG_PATH and -e "$HIP_VDI_HOME/bin/clang") {
+    if (!defined $HIP_CLANG_PATH and (-e "$HIP_VDI_HOME/bin/clang" or -e "$HIP_VDI_HOME/bin/clang.exe")) {
         $HIP_CLANG_PATH = "$HIP_VDI_HOME/bin";
         $HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang";
     }

include/hip/hcc_detail/hip_fp16.h

@@ -34,7 +34,7 @@ THE SOFTWARE.
     #include <utility>
 #endif
 
-#if defined(__clang__) && (__clang_major__ > 5)
+#if __HCC_OR_HIP_CLANG__
     typedef _Float16 _Float16_2 __attribute__((ext_vector_type(2)));
 
     struct __half_raw {

include/hip/hcc_detail/hip_runtime_api.h

@@ -275,6 +275,15 @@ typedef struct dim3 {
 #endif
 } dim3;
 
+typedef struct hipLaunchParams_t {
+    void* func;             ///< Device function symbol
+    dim3 gridDim;           ///< Grid dimentions
+    dim3 blockDim;          ///< Block dimentions
+    void **args;            ///< Arguments
+    size_t sharedMem;       ///< Shared memory
+    hipStream_t stream;     ///< Stream identifier
+} hipLaunchParams;
+
 
 // Doxygen end group GlobalDefs
 /**  @} */
@@ -2842,6 +2851,62 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigne
                                  unsigned int sharedMemBytes, hipStream_t stream,
                                  void** kernelParams, void** extra);
 
+/**
+ * @brief launches kernel f with launch parameters and shared memory on stream with arguments passed
+ * to kernelparams or extra, where thread blocks can cooperate and synchronize as they execute
+ *
+ * @param [in] f         Kernel to launch.
+ * @param [in] gridDim   Grid dimensions specified as multiple of blockDim.
+ * @param [in] blockDim  Block dimensions specified in work-items
+ * @param [in] kernelParams A list of kernel arguments
+ * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.  The
+ * kernel can access this with HIP_DYNAMIC_SHARED.
+ * @param [in] stream    Stream where the kernel should be dispatched.  May be 0, in which case th
+ * default stream is used with associated synchronization rules.
+ *
+ * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
+ */
+hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDimX,
+                                      void** kernelParams, unsigned int sharedMemBytes,
+                                      hipStream_t stream);
+
+/**
+ * @brief Launches kernels on multiple devices where thread blocks can cooperate and
+ * synchronize as they execute.
+ *
+ * @param [in] hipLaunchParams          List of launch parameters, one per device.
+ * @param [in] numDevices               Size of the launchParamsList array.
+ * @param [in] flags                    Flags to control launch behavior.
+ *
+ * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
+ */
+hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
+                                                 int  numDevices, unsigned int  flags);
+
+/**
+ * @brief Returns occupancy for a device function.
+ *
+ * @param [out] numBlocks        Returned occupancy
+ * @param [in]  func             Kernel function for which occupancy is calulated
+ * @param [in]  blockSize        Block size the kernel is intended to be launched with
+ * @param [in]  dynamicSMemSize  Per - block dynamic shared memory usage intended, in bytes
+ */
+hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
+   int* numBlocks, const void* f, int  blockSize, size_t dynamicSMemSize);
+
+/**
+ * @brief Returns occupancy for a device function.
+ *
+ * @param [out] numBlocks        Returned occupancy
+ * @param [in]  func             Kernel function for which occupancy is calulated
+ * @param [in]  blockSize        Block size the kernel is intended to be launched with
+ * @param [in]  dynamicSMemSize  Per - block dynamic shared memory usage intended, in bytes
+ * @param [in]  flags            Extra flags for occupancy calculation (currently ignored)
+ */
+hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+   int* numBlocks, const void* f, int  blockSize, size_t dynamicSMemSize, unsigned int flags);
+
+
 // doxygen end Version Management
 /**
  * @}
@@ -3170,6 +3235,34 @@ hipError_t hipBindTextureToMipmappedArray(const texture<T, dim, readMode>& tex,
     return hipSuccess;
 }
 
+template <class T>
+inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
+    int* numBlocks, T f, int  blockSize, size_t dynamicSMemSize) {
+    return hipOccupancyMaxActiveBlocksPerMultiprocessor(
+        numBlocks, reinterpret_cast<const void*>(f), blockSize, dynamicSMemSize);
+}
+
+template <class T>
+inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+    int* numBlocks, T f, int  blockSize, size_t dynamicSMemSize, unsigned int flags) {
+    return hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+        numBlocks, reinterpret_cast<const void*>(f), blockSize, dynamicSMemSize, flags);
+}
+
+template <class T>
+inline hipError_t hipLaunchCooperativeKernel(T f, dim3 gridDim, dim3 blockDim,
+                                             void** kernelParams, unsigned int sharedMemBytes, hipStream_t stream) {
+    return hipLaunchCooperativeKernel(reinterpret_cast<const void*>(f), gridDim,
+                                      blockDim, kernelParams, sharedMemBytes, stream);
+}
+
+template <class T>
+inline hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
+                                                        unsigned int  numDevices, unsigned int  flags = 0) {
+    return hipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags);
+}
+
+
 /*
  * @brief Unbinds the textuer bound to @p tex
  *

include/hip/hip_runtime_api.h

@@ -115,6 +115,8 @@ typedef struct hipDeviceProp_t {
     int canMapHostMemory;                     ///< Check whether HIP can map host memory
     int gcnArch;                              ///< AMD GCN Arch Value. Eg: 803, 701
     int integrated;            ///< APU vs dGPU
+    int cooperativeLaunch;            ///< HIP device supports cooperative launch
+    int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices 
 } hipDeviceProp_t;
 
 
@@ -291,6 +293,8 @@ typedef enum hipDeviceAttribute_t {
                                                          ///< Multiprocessor.
     hipDeviceAttributeIsMultiGpuBoard,                   ///< Multiple GPU devices.
     hipDeviceAttributeIntegrated,                        ///< iGPU
+    hipDeviceAttributeCooperativeLaunch,                 ///< Support cooperative launch
+    hipDeviceAttributeCooperativeMultiDeviceLaunch,      ///< Support cooperative launch on multiple devices
 } hipDeviceAttribute_t;
 
 enum hipComputeMode {