IPC prototyps and part of the implementation included

Change-Id: Id88c7f155d23ec63f57a6ef05098fba43f8af336 [ROCm/hip commit: 17b98d59b8]
2016-12-06 14:09:53 -06:00
parent 2b557e949e
commit 6d2bf52f9f
@@ -771,7 +771,8 @@ extern "C" __device__ void __threadfence(void);
 *
 * @warning __threadfence_system is a stub and map to no-op.
 */
-__device__ void  __threadfence_system(void) __attribute__((deprecated("Provided with workaround configuration, see hip_kernel_language.md for details")));
+//__device__ void  __threadfence_system(void) __attribute__((deprecated("Provided with workaround configuration, see hip_kernel_language.md for details")));
+__device__ void  __threadfence_system(void) ;

 __device__ unsigned __hip_ds_bpermute(int index, unsigned src);
 __device__ float __hip_ds_bpermutef(int index, float src);
@@ -60,24 +60,13 @@ typedef struct ihipStream_t *hipStream_t;
 //TODO: IPC implementation

 #define hipIpcMemLazyEnablePeerAccess 0
-struct ihipIpcMemHandle_t;
-typedef struct ihipIpcMemHandle_t *hipIpcMemHandle_t;
+
+typedef struct ihipIpcMemHandle *hipIpcMemHandle_t;
+
+//TODO: IPC event handle currently unsupported
 struct ihipIpcEventHandle_t;
 typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t;

-typedef std::nullptr_t nullptr_t ;
-
-__device__ double  
-__longlong_as_double(long long int x)
-{
-  return (double)x;
-}
-__device__ long long int    
-__double_as_longlong(double x)
-{
-  return (long long int)x; 
-}
-

 //END TODO

@@ -1828,10 +1817,97 @@ hipError_t hipProfilerStop();
 */

 //TODO: implement IPC apis
-hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
-hipError_t hipIpcCloseMemHandle(void *devPtr);
-hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle);
-hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags);
+
+/**
+ * @brief Gets an interprocess memory handle for an existing device memory
+ *          allocation
+ *
+ * Takes a pointer to the base of an existing device memory allocation created 
+ * with hipMalloc and exports it for use in another process. This is a 
+ * lightweight operation and may be called multiple times on an allocation
+ * without adverse effects. 
+ *
+ * If a region of memory is freed with hipFree and a subsequent call
+ * to hipMalloc returns memory with the same device address,
+ * hipIpcGetMemHandle will return a unique handle for the
+ * new memory. 
+ *
+ * @param handle - Pointer to user allocated hipIpcMemHandle to return
+ *                    the handle in.
+ * @param devPtr - Base pointer to previously allocated device memory 
+ *
+ * @returns
+ * hipSuccess,
+ * hipErrorInvalidResourceHandle,
+ * hipErrorMemoryAllocation,
+ * hipErrorMapBufferObjectFailed,
+ *
+ */
+extern __host__ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t *handle, void *devPtr);
+
+/**
+ * @brief Opens an interprocess memory handle exported from another process
+ *          and returns a device pointer usable in the local process.
+ *
+ * Maps memory exported from another process with hipIpcGetMemHandle into
+ * the current device address space. For contexts on different devices 
+ * hipIpcOpenMemHandle can attempt to enable peer access between the
+ * devices as if the user called hipDeviceEnablePeerAccess. This behavior is 
+ * controlled by the hipIpcMemLazyEnablePeerAccess flag. 
+ * hipDeviceCanAccessPeer can determine if a mapping is possible.
+ *
+ * Contexts that may open hipIpcMemHandles are restricted in the following way.
+ * hipIpcMemHandles from each device in a given process may only be opened 
+ * by one context per device per other process.
+ *
+ * Memory returned from hipIpcOpenMemHandle must be freed with
+ * hipIpcCloseMemHandle.
+ *
+ * Calling hipFree on an exported memory region before calling
+ * hipIpcCloseMemHandle in the importing context will result in undefined
+ * behavior.
+ * 
+ * @param devPtr - Returned device pointer
+ * @param handle - hipIpcMemHandle to open
+ * @param flags  - Flags for this operation. Must be specified as hipIpcMemLazyEnablePeerAccess
+ *
+ * @returns
+ * hipSuccess,
+ * hipErrorMapBufferObjectFailed,
+ * hipErrorInvalidResourceHandle,
+ * hipErrorTooManyPeers
+ *
+ * @note No guarantees are made about the address returned in @p *devPtr.  
+ * In particular, multiple processes may not receive the same address for the same @p handle.
+ *
+ */
+extern __host__ hipError_t hipIpcOpenMemHandle(void **devPtr, hipIpcMemHandle_t handle, unsigned int flags);
+
+/**
+ * @brief Close memory mapped with hipIpcOpenMemHandle
+ * 
+ * Unmaps memory returnd by hipIpcOpenMemHandle. The original allocation
+ * in the exporting process as well as imported mappings in other processes
+ * will be unaffected.
+ *
+ * Any resources used to enable peer access will be freed if this is the
+ * last mapping using them.
+ *
+ * @param devPtr - Device pointer returned by hipIpcOpenMemHandle
+ * 
+ * @returns
+ * hipSuccess,
+ * hipErrorMapBufferObjectFailed,
+ * hipErrorInvalidResourceHandle,
+ *
+ */
+extern __host__ hipError_t hipIpcCloseMemHandle(void *devPtr);
+
+
+// hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
+// hipError_t hipIpcCloseMemHandle(void *devPtr);
+// // hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle);
+// hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags);


 #ifdef __cplusplus
@@ -343,6 +343,10 @@ __HIP_DEVICE__ double2 make_double2(double, double );
 __HIP_DEVICE__ double3 make_double3(double, double, double );
 __HIP_DEVICE__ double4 make_double4(double, double, double, double );

+extern __HIP_DEVICE__ double  __longlong_as_double(long long int x);
+extern __HIP_DEVICE__ long long int __double_as_longlong(double x);
+
+
 /*
 ///---
 // Inline functions for creating vector types from basic types
@@ -14,8 +14,7 @@ all copies or substantial portions of the Software.
 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
-AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
@@ -2587,7 +2586,18 @@ __HIP_DEVICE__ double4 make_double4(double x, double y, double z, double w)
    return d4;
 }

-__device__ void  __threadfence_system(void){
+
+__HIP_DEVICE__ double  __longlong_as_double(long long int x)
+{
+  return static_cast<double>(x);
+}
+
+__HIP_DEVICE__ long long __double_as_longlong(double x)
+{
+  return static_cast<long long>(x);
+}
+
+__HIP_DEVICE__ void  __threadfence_system(void){
    // no-op
 }

@@ -3380,3 +3390,5 @@ __host__ double norm4d(double a, double b, double c, double d)
 {
    return std::sqrt(a*a + b*b + c*c + d*d);
 }
+
+
@@ -25,6 +25,7 @@ THE SOFTWARE.

 #include <hc.hpp>
 #include <hsa/hsa.h>
+#include "hsa/hsa_ext_amd.h"
 #include "hip_util.h"


@@ -367,6 +368,26 @@ struct LockedBase {
    MUTEX_TYPE  _mutex;
 };

+/**
+ * HIP IPC Handle Size
+ */
+#define HIP_IPC_HANDLE_SIZE 64
+struct __HIP_DEVICE__ ihipIpcMemHandle
+{
+    volatile hsa_amd_ipc_memory_t handle; ///< ipc memory handle on ROCr
+    char reserved[HIP_IPC_HANDLE_SIZE];
+};
+
+
+class ihipModule_t{
+public:
+  hsa_executable_t executable;
+  hsa_code_object_t object;
+  std::string fileName;
+  void *ptr;
+  size_t size;
+};
+
 class ihipFunction_t{
 public:
    ihipFunction_t(const char *name) {
@@ -507,9 +528,9 @@ private:

    // The unsigned return is hipMemcpyKind
    unsigned resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem);
-    void resolveHcMemcpyDirection(unsigned hipMemKind, 
-                                  const hc::AmPointerInfo *dstPtrInfo, const hc::AmPointerInfo *srcPtrInfo, 
-                                  hc::hcCommandKind *hcCopyDir, 
+    void resolveHcMemcpyDirection(unsigned hipMemKind,
+                                  const hc::AmPointerInfo *dstPtrInfo, const hc::AmPointerInfo *srcPtrInfo,
+                                  hc::hcCommandKind *hcCopyDir,
                                  ihipCtx_t **copyDevice,
                                  bool *forceUnpinnedCopy);

@@ -1037,15 +1037,24 @@ hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDevi


 //TODO: IPC implementaiton:
-hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
-    return hipSuccess;
-}
-hipError_t hipIpcCloseMemHandle(void *devPtr){
-    return hipSuccess;
-}
-hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle){
-    return hipSuccess;
-}
 hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags){
-    return hipSuccess;
+    // HIP_INIT_API ( devPtr, handle.handle , flags);
+    hipError_t hipStatus = hipSuccess;
+    return hipStatus;
 }
+
+hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
+    HIP_INIT_API ( handle, devPtr);
+    hipError_t hipStatus = hipSuccess;
+    return hipStatus;
+}
+
+hipError_t hipIpcCloseMemHandle(void *devPtr){
+    HIP_INIT_API ( devPtr );
+    hipError_t hipStatus = hipSuccess;
+    return hipStatus;
+}
+
+// hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle){
+//     return hipSuccess;
+// }