Merge 'master' into 'amd-master'

Change-Id: I759c530530011ee9cff5aef00944c190e4204df2


[ROCm/hip commit: 0573395340]

projects/hip/include/hip/hcc_detail/hip_runtime_api.h

@@ -177,6 +177,8 @@ enum hipLimit_t {
     0x80000000  ///< Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific
                 ///< allocation.
 
+#define hipDeviceMallocDefault 0x0
+#define hipDeviceMallocFinegrained 0x1  ///< Memory is allocated in fine grained region of device.
 
 //! Flags that can be used with hipHostRegister
 #define hipHostRegisterDefault 0x0   ///< Memory is Mapped and Portable
@@ -1055,6 +1057,22 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void
  */
 hipError_t hipMalloc(void** ptr, size_t size);
 
+/**
+ *  @brief Allocate memory on the default accelerator
+ *
+ *  @param[out] ptr Pointer to the allocated memory
+ *  @param[in]  size Requested memory size
+ *  @param[in]  flags Type of memory allocation
+ *
+ *  If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
+ *
+ *  @return #hipSuccess, #hipErrorMemoryAllocation, #hipErrorInvalidValue (bad context, null *ptr)
+ *
+ *  @see hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D, hipMalloc3DArray,
+ * hipHostFree, hipHostMalloc
+ */
+hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags);
+
 /**
  *  @brief Allocate pinned host memory [Deprecated]
  *

projects/hip/src/hip_hcc.cpp

@@ -2117,8 +2117,9 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes,
     }
 }
 
-void ihipStream_t::locked_copy2DSync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind,
+bool ihipStream_t::locked_copy2DSync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind,
                                    bool resolveOn) {
+    bool retStatus = true;
     ihipCtx_t* ctx = this->getCtx();
     const ihipDevice_t* device = ctx->getDevice();
 
@@ -2163,11 +2164,20 @@ void ihipStream_t::locked_copy2DSync(void* dst, const void* src, size_t width, s
                 sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy);
         printPointerInfo(DB_COPY, "  dst", dst, dstPtrInfo);
         printPointerInfo(DB_COPY, "  src", src, srcPtrInfo);
-
+#if (__hcc_workweek__ >= 19101)
+        if(!crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo,
+                           copyDevice ? &copyDevice->getDevice()->_acc : nullptr,
+                           forceUnpinnedCopy)) {
+            tprintf(DB_COPY,"locked_copy2DSync failed to use SDMA\n");
+            retStatus = false;
+        }
+#else
         crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo,
                            copyDevice ? &copyDevice->getDevice()->_acc : nullptr,
                            forceUnpinnedCopy);
+#endif
     }
+    return retStatus;
 }
 
 void ihipStream_t::addSymbolPtrToTracker(hc::accelerator& acc, void* ptr, size_t sizeBytes) {
@@ -2340,8 +2350,9 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
     }
 }
 
-void ihipStream_t::locked_copy2DAsync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind)
+bool ihipStream_t::locked_copy2DAsync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind)
 {
+    bool retStatus = true;
     const ihipCtx_t* ctx = this->getCtx();
 
     if ((ctx == nullptr) || (ctx->getDevice() == nullptr)) {
@@ -2376,13 +2387,26 @@ void ihipStream_t::locked_copy2DAsync(void* dst, const void* src, size_t width,
 
         try {
              if (HIP_FORCE_SYNC_COPY) {
+#if (__hcc_workweek__ >= 19101)
+                 if(!crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo,
+                           &copyDevice->getDevice()->_acc,
+                           forceUnpinnedCopy)){
+                     tprintf(DB_COPY,"locked_copy2DASync with HIP_FORCE_SYNC_COPY failed to use SDMA\n");
+                     retStatus = false;
+                 }
+#else
                  crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo,
                            &copyDevice->getDevice()->_acc,
                            forceUnpinnedCopy);
+#endif
 
              } else {
-                 crit->_av.copy2d_async_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo,
+                 const auto& future = crit->_av.copy2d_async_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo,
                                           &copyDevice->getDevice()->_acc);
+                 if(!future.valid()) {
+                     tprintf(DB_COPY, "locked_copy2DAsync failed to use SDMA\n");
+                     retStatus = false;
+                 }
              }
          } catch (Kalmar::runtime_exception) {
                 throw ihipException(hipErrorRuntimeOther);
@@ -2397,10 +2421,20 @@ void ihipStream_t::locked_copy2DAsync(void* dst, const void* src, size_t width,
     } else {
          //Do sync 2D copy
          LockedAccessor_StreamCrit_t crit(_criticalData);
+#if (__hcc_workweek__ >= 19101)
+         if(!crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo,
+                           copyDevice ? &copyDevice->getDevice()->_acc : nullptr,
+                           forceUnpinnedCopy)){
+             tprintf(DB_COPY, "locked_copy2DAsync Sync copy failed to use SDMA\n");
+             retStatus = false;
+         }
+#else
          crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo,
                            copyDevice ? &copyDevice->getDevice()->_acc : nullptr,
                            forceUnpinnedCopy);
+#endif
     }
+    return retStatus;
 }
 
 //-------------------------------------------------------------------------------------------------

projects/hip/src/hip_hcc_internal.h

@@ -529,12 +529,12 @@ class ihipStream_t {
     void locked_copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind,
                          bool resolveOn = true);
 
-    void locked_copy2DSync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind,
+    bool locked_copy2DSync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind,
                          bool resolveOn = true);
 
     void locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind);
 
-    void locked_copy2DAsync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind);
+    bool locked_copy2DAsync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind);
 
     void lockedSymbolCopySync(hc::accelerator& acc, void* dst, void* src, size_t sizeBytes,
                               size_t offset, unsigned kind);

projects/hip/src/hip_memory.cpp

@@ -264,6 +264,42 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
     return ihipLogStatus(hip_status);
 }
 
+hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags) {
+    HIP_INIT_SPECIAL_API(hipExtMallocWithFlags, (TRACE_MEM), ptr, sizeBytes, flags);
+    HIP_SET_DEVICE();
+
+#if (__hcc_workweek__ >= 19115)
+    hipError_t hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    // return NULL pointer when malloc size is 0
+    if (sizeBytes == 0) {
+        *ptr = NULL;
+        hip_status = hipSuccess;
+    } else if ((ctx == nullptr) || (ptr == nullptr)) {
+        hip_status = hipErrorInvalidValue;
+    } else {
+        unsigned amFlags = 0;
+        if (flags & hipDeviceMallocFinegrained) {
+            amFlags = amDeviceFinegrained;
+        } else if (flags != hipDeviceMallocDefault) {
+            hip_status = hipErrorInvalidValue;
+            return ihipLogStatus(hip_status);
+        }
+        auto device = ctx->getWriteableDevice();
+        *ptr = hip_internal::allocAndSharePtr("device_mem", sizeBytes, ctx, false /*shareWithAll*/,
+                                              amFlags /*amFlags*/, 0 /*hipFlags*/, 0);
+
+        if (sizeBytes && (*ptr == NULL)) {
+            hip_status = hipErrorMemoryAllocation;
+        }
+    }
+#else
+    hipError_t hip_status = hipErrorMemoryAllocation;
+#endif
+
+    return ihipLogStatus(hip_status);
+}
 
 hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
     HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags);
@@ -1573,7 +1609,10 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch
                     stream->locked_copySync((unsigned char*)dst + i * dpitch,
                                     (unsigned char*)src + i * spitch, width, kind);
             } else {
-                stream->locked_copy2DSync(dst, src, width, height, spitch, dpitch, kind);
+                if(!stream->locked_copy2DSync(dst, src, width, height, spitch, dpitch, kind)){
+                    ihipMemcpy2dKernel<uint8_t> (stream, static_cast<uint8_t*> (dst), static_cast<const uint8_t*> (src), width, height, dpitch, spitch);
+                    stream->locked_wait();
+                }
             }
         } catch (ihipException& ex) {
             e = ex._code;
@@ -1621,7 +1660,9 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
                     e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch,
                                           (unsigned char*)src + i * spitch, width, kind, stream);
             } else{
-                stream->locked_copy2DAsync(dst, src, width, height, spitch, dpitch, kind);
+                if(!stream->locked_copy2DAsync(dst, src, width, height, spitch, dpitch, kind)){
+                    ihipMemcpy2dKernel<uint8_t> (stream, static_cast<uint8_t*> (dst), static_cast<const uint8_t*> (src), width, height, dpitch, spitch);
+                }
             }
         } catch (ihipException& ex) {
             e = ex._code;