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

Change-Id: Ibbb31c0a6d1d92e96a6c9abed0d4d4eb0b357008

include/hip/hcc_detail/hip_runtime_api.h

@@ -232,6 +232,9 @@ enum hipLimit_t {
 #define hipCpuDeviceId ((int)-1)
 #define hipInvalidDeviceId ((int)-2)
 
+// Flags that can be used with hipExtLaunch Set of APIs
+#define hipExtAnyOrderLaunch 0x01  ///< AnyOrderLaunch of kernels
+
 /*
  * @brief HIP Memory Advise values
  * @enum

rocclr/hip_device.cpp

@@ -49,15 +49,6 @@ hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) {
   HIP_RETURN(hipSuccess);
 };
 
-hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig) {
-
-  HIP_INIT_API(hipFuncSetCacheConfig, cacheConfig);
-
-  // No way to set cache config yet.
-
-  HIP_RETURN(hipSuccess);
-}
-
 hipError_t hipDeviceTotalMem (size_t *bytes, hipDevice_t device) {
 
   HIP_INIT_API(hipDeviceTotalMem, bytes, device);

rocclr/hip_module.cpp

@@ -201,6 +201,15 @@ hipError_t hipFuncSetAttribute ( const void* func, hipFuncAttribute attr, int va
   HIP_RETURN(hipSuccess);
 }
 
+hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig) {
+
+  HIP_INIT_API(hipFuncSetCacheConfig, cacheConfig);
+
+  // No way to set cache config yet.
+
+  HIP_RETURN(hipSuccess);
+}
+
 hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                  uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
                                  uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
@@ -288,6 +297,11 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
 
   profileNDRange = (startEvent != nullptr && stopEvent != nullptr);
 
+  // Flag set to 1 signifies that kernel can be launched in anyorder
+  if (flags & hipExtAnyOrderLaunch) {
+      params |= amd::NDRangeKernelCommand::AnyOrderLaunch;
+  }
+
   amd::NDRangeKernelCommand* command = new amd::NDRangeKernelCommand(
     *queue, waitList, *kernel, ndrange, sharedMemBytes,
     params, gridId, numGrids, prevGridSum, allGridSum, firstDevice, profileNDRange);
@@ -456,14 +470,17 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
 
   hipError_t result = hipErrorUnknown;
   uint64_t allGridSize = 0;
+  uint32_t blockDims = 0;
   std::vector<const amd::Device*> mgpu_list(numDevices);
 
   for (int i = 0; i < numDevices; ++i) {
     const hipLaunchParams& launch = launchParamsList[i];
-    allGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z;
+    blockDims = launch.blockDim.x * launch.blockDim.y * launch.blockDim.z;
+    allGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z *
+                   blockDims;
 
     // Make sure block dimensions are valid
-    if (0 == launch.blockDim.x * launch.blockDim.y * launch.blockDim.z) {
+    if (0 == blockDims) {
       return hipErrorInvalidConfiguration;
     }
     if (launch.stream != nullptr) {