From abc3e6df5c4000bc51dd02b6e32b8f1b2e92e963 Mon Sep 17 00:00:00 2001
From: German Andryeyev <German.Andryeyev@amd.com>
Date: Mon, 6 Apr 2020 13:54:03 -0400
Subject: [PATCH] SWDEV-184710

Support hipLaunchCooperativeKernelMultiDevice()

- Add validation logic for MGPU launches to pass a cuda test

Change-Id: Iccca7fde43493fc3bc6685512d39202271ae3e92


[ROCm/clr commit: 8be723e199495e41945bd0e490525b5685863603]
---
 .../clr/hipamd/include/hip/hip_runtime_api.h  | 18 ++++++++++++++++-
 projects/clr/hipamd/vdi/hip_device.cpp        |  5 +++++
 .../clr/hipamd/vdi/hip_device_runtime.cpp     | 12 +++++++++++
 projects/clr/hipamd/vdi/hip_module.cpp        | 20 +++++++++++++++++++
 4 files changed, 54 insertions(+), 1 deletion(-)

diff --git a/projects/clr/hipamd/include/hip/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hip_runtime_api.h
index 025688e98c..cf6a64ad65 100644
--- a/projects/clr/hipamd/include/hip/hip_runtime_api.h
+++ b/projects/clr/hipamd/include/hip/hip_runtime_api.h
@@ -128,6 +128,14 @@ typedef struct hipDeviceProp_t {
     int kernelExecTimeoutEnabled;    ///<Run time limit for kernels executed on the device
     int ECCEnabled;                  ///<Device has ECC support enabled
     int tccDriver;                   ///< 1:If device is Tesla device using TCC driver, else 0
+    int cooperativeMultiDeviceUnmatchedFunc;        ///< HIP device supports cooperative launch on multiple
+                                                    ///devices with unmatched functions
+    int cooperativeMultiDeviceUnmatchedGridDim;     ///< HIP device supports cooperative launch on multiple
+                                                    ///devices with unmatched grid dimensions
+    int cooperativeMultiDeviceUnmatchedBlockDim;    ///< HIP device supports cooperative launch on multiple
+                                                    ///devices with unmatched block dimensions
+    int cooperativeMultiDeviceUnmatchedSharedMem;   ///< HIP device supports cooperative launch on multiple
+                                                    ///devices with unmatched shared memories
 
 } hipDeviceProp_t;
 
@@ -329,8 +337,16 @@ typedef enum hipDeviceAttribute_t {
     hipDeviceAttributeTexturePitchAlignment, ///<Pitch alignment requirement for 2D texture references bound to pitched memory;
     hipDeviceAttributeKernelExecTimeout,    ///<Run time limit for kernels executed on the device
     hipDeviceAttributeCanMapHostMemory,     ///<Device can map host memory into device address space
-    hipDeviceAttributeEccEnabled            ///<Device has ECC support enabled
+    hipDeviceAttributeEccEnabled,           ///<Device has ECC support enabled
 
+    hipDeviceAttributeCooperativeMultiDeviceUnmatchedFunc,        ///< Supports cooperative launch on multiple
+                                                                  ///devices with unmatched functions
+    hipDeviceAttributeCooperativeMultiDeviceUnmatchedGridDim,     ///< Supports cooperative launch on multiple
+                                                                  ///devices with unmatched grid dimensions
+    hipDeviceAttributeCooperativeMultiDeviceUnmatchedBlockDim,    ///< Supports cooperative launch on multiple
+                                                                  ///devices with unmatched block dimensions
+    hipDeviceAttributeCooperativeMultiDeviceUnmatchedSharedMem    ///< Supports cooperative launch on multiple
+                                                                  ///devices with unmatched shared memories
 } hipDeviceAttribute_t;
 
 enum hipComputeMode {
diff --git a/projects/clr/hipamd/vdi/hip_device.cpp b/projects/clr/hipamd/vdi/hip_device.cpp
index d1df525877..30b2292271 100644
--- a/projects/clr/hipamd/vdi/hip_device.cpp
+++ b/projects/clr/hipamd/vdi/hip_device.cpp
@@ -197,6 +197,11 @@ hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device )
   deviceProps.cooperativeLaunch = info.cooperativeGroups_;
   deviceProps.cooperativeMultiDeviceLaunch = info.cooperativeMultiDeviceGroups_;
 
+  deviceProps.cooperativeMultiDeviceUnmatchedFunc = info.cooperativeMultiDeviceGroups_;
+  deviceProps.cooperativeMultiDeviceUnmatchedGridDim = info.cooperativeMultiDeviceGroups_;
+  deviceProps.cooperativeMultiDeviceUnmatchedBlockDim = info.cooperativeMultiDeviceGroups_;
+  deviceProps.cooperativeMultiDeviceUnmatchedSharedMem = info.cooperativeMultiDeviceGroups_;
+
   deviceProps.maxTexture1D = info.imageMaxBufferSize_;
   deviceProps.maxTexture2D[0] = info.image2DMaxWidth_;
   deviceProps.maxTexture2D[1] = info.image2DMaxHeight_;
diff --git a/projects/clr/hipamd/vdi/hip_device_runtime.cpp b/projects/clr/hipamd/vdi/hip_device_runtime.cpp
index acf73fa656..febf64d116 100644
--- a/projects/clr/hipamd/vdi/hip_device_runtime.cpp
+++ b/projects/clr/hipamd/vdi/hip_device_runtime.cpp
@@ -281,6 +281,18 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
   case hipDeviceAttributeEccEnabled:
     *pi = prop.ECCEnabled;
     break;
+  case hipDeviceAttributeCooperativeMultiDeviceUnmatchedFunc:
+    *pi = prop.cooperativeMultiDeviceUnmatchedFunc;
+    break;
+  case hipDeviceAttributeCooperativeMultiDeviceUnmatchedGridDim:
+    *pi = prop.cooperativeMultiDeviceUnmatchedGridDim;
+    break;
+  case hipDeviceAttributeCooperativeMultiDeviceUnmatchedBlockDim:
+    *pi = prop.cooperativeMultiDeviceUnmatchedBlockDim;
+    break;
+  case hipDeviceAttributeCooperativeMultiDeviceUnmatchedSharedMem:
+    *pi = prop.cooperativeMultiDeviceUnmatchedSharedMem;
+    break;
   default:
     HIP_RETURN(hipErrorInvalidValue);
   }
diff --git a/projects/clr/hipamd/vdi/hip_module.cpp b/projects/clr/hipamd/vdi/hip_module.cpp
index 4bbb0b5bd4..1e5f7d8b31 100755
--- a/projects/clr/hipamd/vdi/hip_module.cpp
+++ b/projects/clr/hipamd/vdi/hip_module.cpp
@@ -481,9 +481,29 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
 
   hipError_t result = hipErrorUnknown;
   uint64_t allGridSize = 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;
+
+    // Make sure block dimensions are valid
+    if (0 == launch.blockDim.x * launch.blockDim.y * launch.blockDim.z) {
+      return hipErrorInvalidConfiguration;
+    }
+    if (launch.stream != nullptr) {
+      // Validate devices to make sure it dosn't have duplicates
+      amd::HostQueue* queue = reinterpret_cast<hip::Stream*>(launch.stream)->asHostQueue();
+      auto device = &queue->vdev()->device();
+      for (int j = 0; j < numDevices; ++j) {
+        if (mgpu_list[j] == device) {
+          return hipErrorInvalidDevice;
+        }
+      }
+      mgpu_list[i] = device;
+    } else {
+      return hipErrorInvalidResourceHandle;
+    }
   }
   uint64_t prevGridSize = 0;
   uint32_t firstDevice = 0;