SWDEV-323441 - Phase-II : per thread default stream

Change-Id: I3c796ddaebcf0223d7faf50c425c1674de215f9d

hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h

@@ -25,6 +25,8 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H
 #define HIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H
 
+#if defined(__clang__) && defined(__HIP__)
+
 /// hipStreamPerThread implementation
 #if defined(HIP_API_PER_THREAD_DEFAULT_STREAM)
     #define __HIP_STREAM_PER_THREAD
@@ -46,6 +48,17 @@ THE SOFTWARE.
     #define hipMemset2D                   __HIP_API_SPT(hipMemset2D)
     #define hipMemset3D                   __HIP_API_SPT(hipMemset3D)
     #define hipMemcpyAsync                __HIP_API_SPT(hipMemcpyAsync)
+    #define hipMemset3DAsync              __HIP_API_SPT(hipMemset3DAsync)
+    #define hipMemset2DAsync              __HIP_API_SPT(hipMemset2DAsync)
+    #define hipMemsetAsync                __HIP_API_SPT(hipMemsetAsync)
+    #define hipMemcpy3DAsync              __HIP_API_SPT(hipMemcpy3DAsync)
+    #define hipMemcpy2DAsync              __HIP_API_SPT(hipMemcpy2DAsync)
+    #define hipMemcpyFromSymbolAsync      __HIP_API_SPT(hipMemcpyFromSymbolAsync)
+    #define hipMemcpyToSymbolAsync        __HIP_API_SPT(hipMemcpyToSymbolAsync)
+    #define hipMemcpyFromArray            __HIP_API_SPT(hipMemcpyFromArray)
+    #define hipMemcpy2DToArray            __HIP_API_SPT(hipMemcpy2DToArray)
+    #define hipMemcpy2DFromArrayAsync     __HIP_API_SPT(hipMemcpy2DFromArrayAsync)
+    #define hipMemcpy2DToArrayAsync       __HIP_API_SPT(hipMemcpy2DToArrayAsync)
 
     // Stream APIs
     #define hipStreamSynchronize          __HIP_API_SPT(hipStreamSynchronize)
@@ -53,6 +66,7 @@ THE SOFTWARE.
     #define hipStreamGetFlags             __HIP_API_SPT(hipStreamGetFlags)
     #define hipStreamGetPriority          __HIP_API_SPT(hipStreamGetPriority)
     #define hipStreamWaitEvent            __HIP_API_SPT(hipStreamWaitEvent)
+    #define hipStreamAddCallback          __HIP_API_SPT(hipStreamAddCallback)
 
     // Event APIs
     #define hipEventRecord               __HIP_API_SPT(hipEventRecord)
@@ -60,6 +74,14 @@ THE SOFTWARE.
     // Launch APIs
     #define hipLaunchKernel               __HIP_API_SPT(hipLaunchKernel)
     #define hipLaunchCooperativeKernel    __HIP_API_SPT(hipLaunchCooperativeKernel)
+
+    // Graph APIs
+    #define hipGraphLaunch                __HIP_API_SPT(hipGraphLaunch)
+    #define hipStreamBeginCapture         __HIP_API_SPT(hipStreamBeginCapture)
+    #define hipStreamEndCapture           __HIP_API_SPT(hipStreamEndCapture)
+    #define hipStreamIsCapturing          __HIP_API_SPT(hipStreamIsCapturing)
+    #define hipStreamGetCaptureInfo       __HIP_API_SPT(hipStreamGetCaptureInfo)
+    #define hipStreamGetCaptureInfo_v2    __HIP_API_SPT(hipStreamGetCaptureInfo_v2)
 #endif
 
 hipError_t hipMemcpy_spt(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind);
@@ -83,13 +105,45 @@ hipError_t hipMemcpy3D_spt(const struct hipMemcpy3DParms* p);
 
 hipError_t hipMemset_spt(void* dst, int value, size_t sizeBytes);
 
+hipError_t hipMemsetAsync_spt(void* dst, int value, size_t sizeBytes, hipStream_t stream);
+
 hipError_t hipMemset2D_spt(void* dst, size_t pitch, int value, size_t width, size_t height);
 
+hipError_t hipMemset2DAsync_spt(void* dst, size_t pitch, int value,
+                            size_t width, size_t height, hipStream_t stream);
+
+hipError_t hipMemset3DAsync_spt(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream);
+
 hipError_t hipMemset3D_spt(hipPitchedPtr pitchedDevPtr, int  value, hipExtent extent );
 
 hipError_t hipMemcpyAsync_spt(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
                           hipStream_t stream);
 
+hipError_t hipMemcpy3DAsync_spt(const hipMemcpy3DParms* p, hipStream_t stream);
+
+hipError_t hipMemcpy2DAsync_spt(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
+                            size_t height, hipMemcpyKind kind, hipStream_t stream);
+
+hipError_t hipMemcpyFromSymbolAsync_spt(void* dst, const void* symbol, size_t sizeBytes,
+                                    size_t offset, hipMemcpyKind kind, hipStream_t stream);
+
+hipError_t hipMemcpyToSymbolAsync_spt(const void* symbol, const void* src, size_t sizeBytes,
+                                  size_t offset, hipMemcpyKind kind, hipStream_t stream);
+
+hipError_t hipMemcpyFromArray_spt(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset,
+                                  size_t count, hipMemcpyKind kind);
+
+hipError_t hipMemcpy2DToArray_spt(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
+                                  size_t spitch, size_t width, size_t height, hipMemcpyKind kind);
+
+hipError_t hipMemcpy2DFromArrayAsync_spt(void* dst, size_t dpitch, hipArray_const_t src,
+                                  size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height,
+                                  hipMemcpyKind kind, hipStream_t stream);
+
+hipError_t hipMemcpy2DToArrayAsync_spt(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
+                                  size_t spitch, size_t width, size_t height, hipMemcpyKind kind,
+                                  hipStream_t stream);
+
 hipError_t hipStreamQuery_spt(hipStream_t stream);
 
 hipError_t hipStreamSynchronize_spt(hipStream_t stream);
@@ -100,6 +154,11 @@ hipError_t hipStreamWaitEvent_spt(hipStream_t stream, hipEvent_t event, unsigned
 
 hipError_t hipStreamGetFlags_spt(hipStream_t stream, unsigned int* flags);
 
+hipError_t hipStreamAddCallback_spt(hipStream_t stream, hipStreamCallback_t callback, void* userData,
+                                unsigned int flags);
+
+hipError_t hipEventRecord_spt(hipEvent_t event, hipStream_t stream);
+
 hipError_t hipLaunchCooperativeKernel_spt(const void* f,
                                       dim3 gridDim, dim3 blockDim,
                                       void **kernelParams, uint32_t sharedMemBytes, hipStream_t hStream);
@@ -114,4 +173,17 @@ hipError_t hipLaunchKernel_spt(const void* function_address,
 #ifdef __cplusplus
 }
 #endif // extern "C"
+
+hipError_t hipGraphLaunch_spt(hipGraphExec_t graphExec, hipStream_t stream);
+hipError_t hipStreamBeginCapture_spt(hipStream_t stream, hipStreamCaptureMode mode);
+hipError_t hipStreamEndCapture_spt(hipStream_t stream, hipGraph_t* pGraph);
+hipError_t hipStreamIsCapturing_spt(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus);
+hipError_t hipStreamGetCaptureInfo_spt(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus,
+                                   unsigned long long* pId);
+hipError_t hipStreamGetCaptureInfo_v2_spt(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out,
+                                      unsigned long long* id_out, hipGraph_t* graph_out,
+                                      const hipGraphNode_t** dependencies_out,
+                                      size_t* numDependencies_out);
+
+#endif //defined(__clang__) && defined(__HIP__)
 #endif //HIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H

hipamd/src/amdhip.def

@@ -406,3 +406,21 @@ hipStreamWriteValue64
 hipStreamWaitValue32
 hipStreamWaitValue64
 hipGetStreamDeviceId
+hipGraphLaunch_spt
+hipStreamBeginCapture_spt
+hipStreamEndCapture_spt
+hipStreamIsCapturing_spt
+hipStreamGetCaptureInfo_spt
+hipStreamGetCaptureInfo_v2_spt
+hipStreamAddCallback_spt
+hipMemsetAsync_spt
+hipMemset2DAsync_spt
+hipMemset3DAsync_spt
+hipMemcpy3DAsync_spt
+hipMemcpy2DAsync_spt
+hipMemcpyFromSymbolAsync_spt
+hipMemcpyToSymbolAsync_spt
+hipMemcpyFromArray_spt
+hipMemcpy2DToArray_spt
+hipMemcpy2DFromArrayAsync_spt
+hipMemcpy2DToArrayAsync_spt

hipamd/src/hip_graph.cpp

@@ -726,17 +726,27 @@ hipError_t capturehipLaunchHostFunc(hipStream_t& stream, hipHostFn_t& fn, void*&
   return hipSuccess;
 }
 
-hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) {
-  HIP_INIT_API(hipStreamIsCapturing, stream, pCaptureStatus);
+hipError_t hipStreamIsCapturing_common(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) {
   if (pCaptureStatus == nullptr || !hip::isValid(stream)) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
   if (stream == nullptr) {
     *pCaptureStatus = hipStreamCaptureStatusNone;
   } else {
     *pCaptureStatus = reinterpret_cast<hip::Stream*>(stream)->GetCaptureStatus();
   }
-  HIP_RETURN(hipSuccess);
+  return hipSuccess;
+}
+
+hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) {
+  HIP_INIT_API(hipStreamIsCapturing, stream, pCaptureStatus);
+  HIP_RETURN(hipStreamIsCapturing_common(stream, pCaptureStatus));
+}
+
+hipError_t hipStreamIsCapturing_spt(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) {
+  HIP_INIT_API(hipStreamIsCapturing, stream, pCaptureStatus);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN(hipStreamIsCapturing_common(stream, pCaptureStatus));
 }
 
 hipError_t hipThreadExchangeStreamCaptureMode(hipStreamCaptureMode* mode) {
@@ -754,22 +764,22 @@ hipError_t hipThreadExchangeStreamCaptureMode(hipStreamCaptureMode* mode) {
   HIP_RETURN_DURATION(hipSuccess);
 }
 
-hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode) {
-  HIP_INIT_API(hipStreamBeginCapture, stream, mode);
+hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode mode) {
   if (!hip::isValid(stream)) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
   // capture cannot be initiated on legacy stream
   if (stream == nullptr) {
-    HIP_RETURN(hipErrorStreamCaptureUnsupported);
+    return hipErrorStreamCaptureUnsupported;
   }
-  if (mode < hipStreamCaptureModeGlobal || mode > hipStreamCaptureModeRelaxed) {
-    HIP_RETURN(hipErrorInvalidValue);
+  if (mode < hipStreamCaptureModeGlobal ||
+      mode > hipStreamCaptureModeRelaxed) {
+    return hipErrorInvalidValue;
   }
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   // It can be initiated if the stream is not already in capture mode
   if (s->GetCaptureStatus() == hipStreamCaptureStatusActive) {
-    HIP_RETURN(hipErrorIllegalState);
+    return hipErrorIllegalState;
   }
 
   s->SetCaptureGraph(new ihipGraph());
@@ -782,35 +792,45 @@ hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode)
     amd::ScopedLock lock(g_captureStreamsLock);
     g_captureStreams.push_back(s);
   }
-  HIP_RETURN_DURATION(hipSuccess);
+  return hipSuccess;
 }
 
-hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) {
-  HIP_INIT_API(hipStreamEndCapture, stream, pGraph);
+hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode) {
+  HIP_INIT_API(hipStreamBeginCapture, stream, mode);
+  HIP_RETURN_DURATION(hipStreamBeginCapture_common(stream, mode));
+}
+
+hipError_t hipStreamBeginCapture_spt(hipStream_t stream, hipStreamCaptureMode mode) {
+  HIP_INIT_API(hipStreamBeginCapture, stream, mode);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN_DURATION(hipStreamBeginCapture_common(stream, mode));
+}
+
+hipError_t hipStreamEndCapture_common(hipStream_t stream, hipGraph_t* pGraph) {
   if (pGraph == nullptr) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
   if (stream == nullptr) {
-    HIP_RETURN(hipErrorIllegalState);
+    return hipErrorIllegalState;
   }
   if (!hip::isValid(stream)) {
-    HIP_RETURN(hipErrorContextIsDestroyed);
+    return hipErrorContextIsDestroyed;
   }
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   // Capture status must be active before endCapture can be initiated
   if (s->GetCaptureStatus() == hipStreamCaptureStatusNone) {
-    HIP_RETURN(hipErrorIllegalState);
+    return hipErrorIllegalState;
   }
   // Capture must be ended on the same stream in which it was initiated
   if (!s->IsOriginStream()) {
-    HIP_RETURN(hipErrorStreamCaptureUnmatched);
+    return hipErrorStreamCaptureUnmatched;
   }
   // If mode is not hipStreamCaptureModeRelaxed, hipStreamEndCapture must be called on the stream
   // from the same thread
   const auto& it = std::find(l_captureStreams.begin(), l_captureStreams.end(), s);
   if (s->GetCaptureMode() != hipStreamCaptureModeRelaxed) {
     if (it == l_captureStreams.end()) {
-      HIP_RETURN(hipErrorStreamCaptureWrongThread);
+      return hipErrorStreamCaptureWrongThread;
     }
     l_captureStreams.erase(it);
   }
@@ -821,7 +841,7 @@ hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) {
   // If capture was invalidated, due to a violation of the rules of stream capture
   if (s->GetCaptureStatus() == hipStreamCaptureStatusInvalidated) {
     *pGraph = nullptr;
-    HIP_RETURN(hipErrorStreamCaptureInvalidated);
+    return hipErrorStreamCaptureInvalidated;
   }
   // check if all parallel streams have joined
   // Nodes that are removed from the dependency set via API hipStreamUpdateCaptureDependencies do
@@ -838,12 +858,23 @@ hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) {
       }
     }
     if (foundInRemovedDep == false) {
-      HIP_RETURN(hipErrorStreamCaptureUnjoined);
+      return hipErrorStreamCaptureUnjoined;
     }
   }
   *pGraph = s->GetCaptureGraph();
   // end capture on all streams/events part of graph capture
-  HIP_RETURN_DURATION(s->EndCapture());
+  return s->EndCapture();
+}
+
+hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) {
+  HIP_INIT_API(hipStreamEndCapture, stream, pGraph);
+  HIP_RETURN_DURATION(hipStreamEndCapture_common(stream, pGraph));
+}
+
+hipError_t hipStreamEndCapture_spt(hipStream_t stream, hipGraph_t* pGraph) {
+  HIP_INIT_API(hipStreamEndCapture, stream, pGraph);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN_DURATION(hipStreamEndCapture_common(stream, pGraph));
 }
 
 hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags) {
@@ -1038,12 +1069,22 @@ hipError_t ihipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) {
   return graphExec->Run(stream);
 }
 
+hipError_t hipGraphLaunch_common( hipGraphExec_t graphExec, hipStream_t stream) {
+  if (graphExec == nullptr || !hip::isValid(stream) || !hipGraphExec::isGraphExecValid(graphExec)) {
+    return hipErrorInvalidValue;
+  }
+  return ihipGraphLaunch(graphExec, stream);
+}
+
 hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) {
   HIP_INIT_API(hipGraphLaunch, graphExec, stream);
-  if (graphExec == nullptr || !hip::isValid(stream) || !hipGraphExec::isGraphExecValid(graphExec)) {
-    HIP_RETURN(hipErrorInvalidValue);
-  }
-  HIP_RETURN_DURATION(ihipGraphLaunch(graphExec, stream));
+  HIP_RETURN_DURATION(hipGraphLaunch_common(graphExec, stream));
+}
+
+hipError_t hipGraphLaunch_spt(hipGraphExec_t graphExec, hipStream_t stream) {
+  HIP_INIT_API(hipGraphLaunch, graphExec, stream);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN_DURATION(hipGraphLaunch_common(graphExec, stream));
 }
 
 hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* numNodes) {
@@ -1310,37 +1351,47 @@ hipError_t hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, hipGra
   HIP_RETURN(reinterpret_cast<hipChildGraphNode*>(clonedNode)->SetParams(childGraph));
 }
 
-hipError_t hipStreamGetCaptureInfo(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus,
+hipError_t hipStreamGetCaptureInfo_common(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus,
                                    unsigned long long* pId) {
-  HIP_INIT_API(hipStreamGetCaptureInfo, stream, pCaptureStatus, pId);
   if (pCaptureStatus == nullptr || !hip::isValid(stream)) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
   if (stream == nullptr) {
-    HIP_RETURN(hipErrorUnknown);
+    return hipErrorUnknown;
   }
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   *pCaptureStatus = s->GetCaptureStatus();
   if (*pCaptureStatus == hipStreamCaptureStatusActive && pId != nullptr) {
     *pId = s->GetCaptureID();
   }
-  HIP_RETURN(hipSuccess);
+  return hipSuccess;
 }
 
-hipError_t hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out,
+hipError_t hipStreamGetCaptureInfo(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus,
+                                   unsigned long long* pId) {
+  HIP_INIT_API(hipStreamGetCaptureInfo, stream, pCaptureStatus, pId);
+  HIP_RETURN(hipStreamGetCaptureInfo_common(stream, pCaptureStatus, pId));
+}
+
+hipError_t hipStreamGetCaptureInfo_spt(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus,
+                                   unsigned long long* pId) {
+  HIP_INIT_API(hipStreamGetCaptureInfo, stream, pCaptureStatus, pId);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN(hipStreamGetCaptureInfo_common(stream, pCaptureStatus, pId));
+}
+
+hipError_t hipStreamGetCaptureInfo_v2_common(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out,
                                       unsigned long long* id_out, hipGraph_t* graph_out,
                                       const hipGraphNode_t** dependencies_out,
                                       size_t* numDependencies_out) {
-  HIP_INIT_API(hipStreamGetCaptureInfo_v2, stream, captureStatus_out, id_out, graph_out,
-               dependencies_out, numDependencies_out);
   if (captureStatus_out == nullptr) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
   if (stream == nullptr) {
-    HIP_RETURN(hipErrorUnknown);
+    return hipErrorUnknown;
   }
   if (!hip::isValid(stream)) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   *captureStatus_out = s->GetCaptureStatus();
@@ -1358,7 +1409,28 @@ hipError_t hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus
       *numDependencies_out = s->GetLastCapturedNodes().size();
     }
   }
-  HIP_RETURN(hipSuccess);
+  return hipSuccess;
+}
+
+hipError_t hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out,
+                                      unsigned long long* id_out, hipGraph_t* graph_out,
+                                      const hipGraphNode_t** dependencies_out,
+                                      size_t* numDependencies_out) {
+  HIP_INIT_API(hipStreamGetCaptureInfo_v2, stream, captureStatus_out, id_out, graph_out,
+               dependencies_out, numDependencies_out);
+  HIP_RETURN(hipStreamGetCaptureInfo_v2_common(stream, captureStatus_out, id_out, graph_out,
+             dependencies_out, numDependencies_out));
+}
+
+hipError_t hipStreamGetCaptureInfo_v2_spt(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out,
+                                      unsigned long long* id_out, hipGraph_t* graph_out,
+                                      const hipGraphNode_t** dependencies_out,
+                                      size_t* numDependencies_out) {
+  HIP_INIT_API(hipStreamGetCaptureInfo_v2, stream, captureStatus_out, id_out, graph_out,
+               dependencies_out, numDependencies_out);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN(hipStreamGetCaptureInfo_v2_common(stream, captureStatus_out, id_out, graph_out,
+             dependencies_out, numDependencies_out));
 }
 
 hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies,

hipamd/src/hip_hcc.def.in

@@ -421,4 +421,21 @@ hipStreamWriteValue64
 hipStreamWaitValue32
 hipStreamWaitValue64
 hipDeviceSetLimit
-hipGetStreamDeviceId
+hipGetStreamDeviceId
+hipGraphLaunch_spt
+hipStreamBeginCapture_spt
+hipStreamEndCapture_spt
+hipStreamIsCapturing_spt
+hipStreamGetCaptureInfo_spt
+hipStreamGetCaptureInfo_v2_spt
+hipStreamAddCallback_spt
+hipMemsetAsync_spt
+hipMemset2DAsync_spt
+hipMemset3DAsync_spt
+hipMemcpy3DAsync_spt
+hipMemcpy2DAsync_spt
+hipMemcpyFromSymbolAsync_spt
+hipMemcpyToSymbolAsync_spt
+hipMemcpyFromArray_spt
+hipMemcpy2DToArray_spt
+hipMemcpy2DFromArrayAsync_spt

hipamd/src/hip_hcc.map.in

@@ -475,6 +475,24 @@ global:
     hipDeviceSetLimit;
     hiprtcGetBitcode;
     hiprtcGetBitcodeSize;
+    hipGraphLaunch_spt;
+    hipStreamBeginCapture_spt;
+    hipStreamEndCapture_spt;
+    hipStreamIsCapturing_spt;
+    hipStreamGetCaptureInfo_spt;
+    hipStreamGetCaptureInfo_v2_spt;
+    hipStreamAddCallback_spt;
+    hipMemsetAsync_spt;
+    hipMemset2DAsync_spt;
+    hipMemset3DAsync_spt;
+    hipMemcpy3DAsync_spt;
+    hipMemcpy2DAsync_spt;
+    hipMemcpyFromSymbolAsync_spt;
+    hipMemcpyToSymbolAsync_spt;
+    hipMemcpyFromArray_spt;
+    hipMemcpy2DToArray_spt;
+    hipMemcpy2DFromArrayAsync_spt;
+    hipMemcpy2DToArrayAsync_spt;
 local:
     *;
 } hip_5.2;

hipamd/src/hip_memory.cpp

@@ -1261,14 +1261,12 @@ hipError_t hipMemcpyFromSymbol_spt(void* dst, const void* symbol, size_t sizeByt
                       getPerThreadDefaultStream()));
 }
 
-hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t sizeBytes,
+hipError_t hipMemcpyToSymbolAsync_common(const void* symbol, const void* src, size_t sizeBytes,
                                   size_t offset, hipMemcpyKind kind, hipStream_t stream) {
-  HIP_INIT_API(hipMemcpyToSymbolAsync, symbol, src, sizeBytes, offset, kind, stream);
-
   STREAM_CAPTURE(hipMemcpyToSymbolAsync, stream, symbol, src, sizeBytes, offset, kind);
 
-  if (kind != hipMemcpyHostToDevice && kind != hipMemcpyDeviceToDevice) {
-    HIP_RETURN(hipErrorInvalidMemcpyDirection);
+  if (kind != hipMemcpyHostToDevice) {
+    return hipErrorInvalidMemcpyDirection;
   }
 
   size_t sym_size = 0;
@@ -1279,29 +1277,53 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t si
     return status;
   }
   /* Copy memory from source to destination address */
-  HIP_RETURN_DURATION(hipMemcpyAsync(device_ptr, src, sizeBytes, kind, stream));
+  return hipMemcpyAsync(device_ptr, src, sizeBytes, kind, stream);
+}
+
+hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t sizeBytes,
+                                  size_t offset, hipMemcpyKind kind, hipStream_t stream) {
+  HIP_INIT_API(hipMemcpyToSymbolAsync, symbol, src, sizeBytes, offset, kind, stream);
+  HIP_RETURN_DURATION(hipMemcpyToSymbolAsync_common(symbol, src, sizeBytes, offset, kind, stream));
+}
+
+hipError_t hipMemcpyToSymbolAsync_spt(const void* symbol, const void* src, size_t sizeBytes,
+                                  size_t offset, hipMemcpyKind kind, hipStream_t stream) {
+  HIP_INIT_API(hipMemcpyToSymbolAsync, symbol, src, sizeBytes, offset, kind, stream);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN_DURATION(hipMemcpyToSymbolAsync_common(symbol, src, sizeBytes, offset, kind, stream));
+}
+
+hipError_t hipMemcpyFromSymbolAsync_common(void* dst, const void* symbol, size_t sizeBytes,
+                                    size_t offset, hipMemcpyKind kind, hipStream_t stream) {
+  STREAM_CAPTURE(hipMemcpyFromSymbolAsync, stream, dst, symbol, sizeBytes, offset, kind);
+
+  if (kind != hipMemcpyDeviceToHost) {
+    return hipErrorInvalidMemcpyDirection;
+  }
+
+  size_t sym_size = 0;
+  hipDeviceptr_t device_ptr = nullptr;
+
+  hipError_t status = ihipMemcpySymbol_validate(symbol, sizeBytes, offset, sym_size, device_ptr);
+  if (status != hipSuccess) {
+    return status;
+  }
+
+  /* Copy memory from source to destination address */
+  return hipMemcpyAsync(dst, device_ptr, sizeBytes, kind, stream);
 }
 
 hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol, size_t sizeBytes,
                                     size_t offset, hipMemcpyKind kind, hipStream_t stream) {
-  HIP_INIT_API(hipMemcpyFromSymbolAsync, symbol, dst, sizeBytes, offset, kind, stream);
+  HIP_INIT_API(hipMemcpyFromSymbolAsync, dst, symbol, sizeBytes, offset, kind, stream);
+  HIP_RETURN_DURATION(hipMemcpyFromSymbolAsync_common(dst, symbol, sizeBytes, offset, kind, stream));
+}
 
-  STREAM_CAPTURE(hipMemcpyFromSymbolAsync, stream, dst, symbol, sizeBytes, offset, kind);
-
-  if (kind != hipMemcpyDeviceToHost && kind != hipMemcpyDeviceToDevice) {
-    HIP_RETURN(hipErrorInvalidMemcpyDirection);
-  }
-
-  size_t sym_size = 0;
-  hipDeviceptr_t device_ptr = nullptr;
-
-  hipError_t status = ihipMemcpySymbol_validate(symbol, sizeBytes, offset, sym_size, device_ptr);
-  if (status != hipSuccess) {
-    return status;
-  }
-
-  /* Copy memory from source to destination address */
-  HIP_RETURN_DURATION(hipMemcpyAsync(dst, device_ptr, sizeBytes, kind, stream));
+hipError_t hipMemcpyFromSymbolAsync_spt(void* dst, const void* symbol, size_t sizeBytes,
+                                    size_t offset, hipMemcpyKind kind, hipStream_t stream) {
+  HIP_INIT_API(hipMemcpyFromSymbolAsync, dst, symbol, sizeBytes, offset, kind, stream);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN_DURATION(hipMemcpyFromSymbolAsync_common(dst, symbol, sizeBytes, offset, kind, stream));
 }
 
 hipError_t hipMemcpyHtoD(hipDeviceptr_t dstDevice,
@@ -2158,11 +2180,23 @@ hipError_t hipMemcpy2D_spt(void* dst, size_t dpitch, const void* src, size_t spi
                             getPerThreadDefaultStream()));
 }
 
+hipError_t hipMemcpy2DAsync_common(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
+                            size_t height, hipMemcpyKind kind, hipStream_t stream) {
+  STREAM_CAPTURE(hipMemcpy2DAsync, stream, dst, dpitch, src, spitch, width, height, kind);
+  return ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, stream, true);
+}
+
 hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
                             size_t height, hipMemcpyKind kind, hipStream_t stream) {
   HIP_INIT_API(hipMemcpy2DAsync, dst, dpitch, src, spitch, width, height, kind, stream);
-  STREAM_CAPTURE(hipMemcpy2DAsync, stream, dst, dpitch, src, spitch, width, height, kind);
-  HIP_RETURN_DURATION(ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, stream, true));
+  HIP_RETURN_DURATION(hipMemcpy2DAsync_common(dst, dpitch, src, spitch, width, height, kind, stream));
+}
+
+hipError_t hipMemcpy2DAsync_spt(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
+                            size_t height, hipMemcpyKind kind, hipStream_t stream) {
+  HIP_INIT_API(hipMemcpy2DAsync, dst, dpitch, src, spitch, width, height, kind, stream);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN_DURATION(hipMemcpy2DAsync_common(dst, dpitch, src, spitch, width, height, kind, stream));
 }
 
 hipError_t ihipMemcpy2DToArray(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream, bool isAsync = false) {
@@ -2259,11 +2293,11 @@ hipError_t ihipMemcpy2DFromArray(void* dst, size_t dpitch, hipArray_const_t src,
   return ihipMemcpyParam2D(&desc, stream, isAsync);
 }
 
-hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset, size_t count, hipMemcpyKind kind) {
-  HIP_INIT_API(hipMemcpyFromArray, dst, src, wOffsetSrc, hOffset, count, kind);
+hipError_t hipMemcpyFromArray_common(void* dst, hipArray_const_t src, size_t wOffsetSrc,
+                                    size_t hOffset, size_t count, hipMemcpyKind kind, hipStream_t stream) {
   CHECK_STREAM_CAPTURING();
   if (src == nullptr) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
 
   const size_t arrayHeight = (src->height != 0) ? src->height : 1;
@@ -2271,7 +2305,18 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t src, size_t wOffsetSrc
 
   const size_t height = (count / src->width) / hip::getElementSize(src);
 
-  HIP_RETURN_DURATION(ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffset, witdthInBytes, height, kind, nullptr));
+  return ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffset, witdthInBytes, height, kind, stream);
+}
+
+hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset, size_t count, hipMemcpyKind kind) {
+  HIP_INIT_API(hipMemcpyFromArray, dst, src, wOffsetSrc, hOffset, count, kind);
+  HIP_RETURN_DURATION(hipMemcpyFromArray_common(dst, src, wOffsetSrc, hOffset, count, kind, nullptr));
+}
+
+hipError_t hipMemcpyFromArray_spt(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset, size_t count, hipMemcpyKind kind) {
+  HIP_INIT_API(hipMemcpyFromArray, dst, src, wOffsetSrc, hOffset, count, kind);
+  HIP_RETURN_DURATION(hipMemcpyFromArray_common(dst, src, wOffsetSrc, hOffset, count, kind,
+                      getPerThreadDefaultStream()));
 }
 
 hipError_t ihipMemcpyAtoD(hipArray* srcArray, void* dstDevice, amd::Coord3D srcOrigin,
@@ -2429,12 +2474,20 @@ hipError_t hipMemcpy3D_spt(const hipMemcpy3DParms* p) {
   HIP_RETURN_DURATION(hipMemcpy3D_common(p, getPerThreadDefaultStream()));
 }
 
+hipError_t hipMemcpy3DAsync_common(const hipMemcpy3DParms* p, hipStream_t stream) {
+  STREAM_CAPTURE(hipMemcpy3DAsync, stream, p);
+  return ihipMemcpy3D(p, stream, true);
+}
+
 hipError_t hipMemcpy3DAsync(const hipMemcpy3DParms* p, hipStream_t stream) {
   HIP_INIT_API(hipMemcpy3DAsync, p, stream);
+  HIP_RETURN_DURATION(hipMemcpy3DAsync_common(p, stream));
+}
 
-  STREAM_CAPTURE(hipMemcpy3DAsync, stream, p);
-
-  HIP_RETURN_DURATION(ihipMemcpy3D(p, stream, true));
+hipError_t hipMemcpy3DAsync_spt(const hipMemcpy3DParms* p, hipStream_t stream) {
+  HIP_INIT_API(hipMemcpy3DAsync, p, stream);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN_DURATION(hipMemcpy3DAsync_common(p, stream));
 }
 
 hipError_t hipDrvMemcpy3D(const HIP_MEMCPY3D* pCopy) {
@@ -2578,11 +2631,21 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
   HIP_RETURN(hipMemset_common(dst, value, sizeBytes));
 }
 
-hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
-  HIP_INIT_API(hipMemsetAsync, dst, value, sizeBytes, stream);
+hipError_t hipMemsetAsync_common(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
   size_t valueSize = sizeof(int8_t);
   STREAM_CAPTURE(hipMemsetAsync, stream, dst, value, valueSize, sizeBytes);
-  HIP_RETURN(ihipMemset(dst, value, sizeof(int8_t), sizeBytes, stream, true));
+  return ihipMemset(dst, value, sizeof(int8_t), sizeBytes, stream, true);
+}
+
+hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
+  HIP_INIT_API(hipMemsetAsync, dst, value, sizeBytes, stream);
+  HIP_RETURN(hipMemsetAsync_common(dst, value, sizeBytes, stream));
+}
+
+hipError_t hipMemsetAsync_spt(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
+  HIP_INIT_API(hipMemsetAsync, dst, value, sizeBytes, stream);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN(hipMemsetAsync_common(dst, value, sizeBytes, stream));
 }
 
 hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t count) {
@@ -2727,14 +2790,24 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t
   HIP_RETURN(hipMemset2D_common(dst, pitch, value, width, height));
 }
 
+hipError_t hipMemset2DAsync_common(void* dst, size_t pitch, int value,
+                            size_t width, size_t height, hipStream_t stream) {
+  STREAM_CAPTURE(hipMemset2DAsync, stream, dst, pitch, value, width, height);
+
+  return ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream, true);
+}
 
 hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value,
                             size_t width, size_t height, hipStream_t stream) {
   HIP_INIT_API(hipMemset2DAsync, dst, pitch, value, width, height, stream);
+  HIP_RETURN(hipMemset2DAsync_common(dst, pitch, value, width, height, stream));
+}
 
-  STREAM_CAPTURE(hipMemset2DAsync, stream, dst, pitch, value, width, height);
-
-  HIP_RETURN(ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream, true));
+hipError_t hipMemset2DAsync_spt(void* dst, size_t pitch, int value,
+                            size_t width, size_t height, hipStream_t stream) {
+  HIP_INIT_API(hipMemset2DAsync, dst, pitch, value, width, height, stream);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN(hipMemset2DAsync_common(dst, pitch, value, width, height, stream));
 }
 
 hipError_t hipMemset3D_common(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream=nullptr) {
@@ -2753,12 +2826,20 @@ hipError_t hipMemset3D_spt(hipPitchedPtr pitchedDevPtr, int value, hipExtent ext
   HIP_RETURN(hipMemset3D_common(pitchedDevPtr, value, extent,stream));
 }
 
+hipError_t hipMemset3DAsync_common(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream) {
+  STREAM_CAPTURE(hipMemset3DAsync, stream, pitchedDevPtr, value, extent);
+  return ihipMemset3D(pitchedDevPtr, value, extent, stream, true);
+}
+
 hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream) {
   HIP_INIT_API(hipMemset3DAsync, pitchedDevPtr, value, extent, stream);
+  HIP_RETURN(hipMemset3DAsync_common(pitchedDevPtr, value, extent, stream));
+}
 
-  STREAM_CAPTURE(hipMemset3DAsync, stream, pitchedDevPtr, value, extent);
-
-  HIP_RETURN(ihipMemset3D(pitchedDevPtr, value, extent, stream, true));
+hipError_t hipMemset3DAsync_spt(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream) {
+  HIP_INIT_API(hipMemset3DAsync, pitchedDevPtr, value, extent, stream);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN(hipMemset3DAsync_common(pitchedDevPtr, value, extent, stream));
 }
 
 hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes,
@@ -3236,16 +3317,24 @@ hipError_t hipMemcpy2DFromArray_spt(void* dst, size_t dpitch, hipArray_const_t s
   HIP_RETURN_DURATION(hipMemcpy2DFromArray_common(dst, dpitch, src, wOffsetSrc, hOffset, width, height, kind, stream));
 }
 
-hipError_t hipMemcpy2DFromArrayAsync(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) {
-  HIP_INIT_API(hipMemcpy2DFromArrayAsync, dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream);
+hipError_t hipMemcpy2DFromArrayAsync_common(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) {
   STREAM_CAPTURE(hipMemcpy2DFromArrayAsync, stream, dst, dpitch, src, wOffsetSrc, hOffsetSrc, width,
                  height, kind);
-
   if (dpitch == 0) {
-    HIP_RETURN(hipErrorInvalidPitchValue);
+    return hipErrorInvalidPitchValue;
   }
+  return ihipMemcpy2DFromArray(dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream, true);
+}
 
-  HIP_RETURN_DURATION(ihipMemcpy2DFromArray(dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream, true));
+hipError_t hipMemcpy2DFromArrayAsync(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) {
+  HIP_INIT_API(hipMemcpy2DFromArrayAsync, dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream);
+  HIP_RETURN_DURATION(hipMemcpy2DFromArrayAsync_common(dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream));
+}
+
+hipError_t hipMemcpy2DFromArrayAsync_spt(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) {
+  HIP_INIT_API(hipMemcpy2DFromArrayAsync, dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN_DURATION(hipMemcpy2DFromArrayAsync_common(dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream));
 }
 
 hipError_t hipMemcpyFromArrayAsync(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, hipMemcpyKind kind, hipStream_t stream) {
@@ -3264,15 +3353,24 @@ hipError_t hipMemcpyFromArrayAsync(void* dst, hipArray_const_t src, size_t wOffs
   HIP_RETURN_DURATION(ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffsetSrc, widthInBytes, height, kind, stream, true));
 }
 
-hipError_t hipMemcpy2DToArrayAsync(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) {
-  HIP_INIT_API(hipMemcpy2DToArrayAsync, dst, wOffset, hOffset, src, spitch, width, height, kind);
+hipError_t hipMemcpy2DToArrayAsync_common(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) {
   STREAM_CAPTURE(hipMemcpy2DToArrayAsync, stream, dst, wOffset, hOffset, src, spitch, width, height,
                  kind);
   if (spitch == 0) {
-    HIP_RETURN(hipErrorInvalidPitchValue);
+    return hipErrorInvalidPitchValue;
   }
+  return ihipMemcpy2DToArray(dst, wOffset, hOffset, src, spitch, width, height, kind, stream, true);
+}
 
-  HIP_RETURN_DURATION(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, spitch, width, height, kind, stream, true));
+hipError_t hipMemcpy2DToArrayAsync(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) {
+  HIP_INIT_API(hipMemcpy2DToArrayAsync, dst, wOffset, hOffset, src, spitch, width, height, kind, stream);
+  HIP_RETURN_DURATION(hipMemcpy2DToArrayAsync_common(dst, wOffset, hOffset, src, spitch, width, height, kind, stream));
+}
+
+hipError_t hipMemcpy2DToArrayAsync_spt(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) {
+  HIP_INIT_API(hipMemcpy2DToArrayAsync, dst, wOffset, hOffset, src, spitch, width, height, kind, stream);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN_DURATION(hipMemcpy2DToArrayAsync_common(dst, wOffset, hOffset, src, spitch, width, height, kind, stream));
 }
 
 hipError_t hipMemcpyToArrayAsync(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind, hipStream_t stream) {

hipamd/src/hip_stream.cpp

@@ -574,16 +574,15 @@ hipError_t hipStreamQuery_spt(hipStream_t stream) {
 }
 
 // ================================================================================================
-hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
+hipError_t hipStreamAddCallback_common(hipStream_t stream, hipStreamCallback_t callback, void* userData,
                                 unsigned int flags) {
-  HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags);
   //flags - Reserved for future use, must be 0
   if (callback == nullptr || flags != 0) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
 
   if (!hip::isValid(stream)) {
-    HIP_RETURN(hipErrorContextIsDestroyed);
+    return hipErrorContextIsDestroyed;
   }
 
   amd::HostQueue* hostQueue = hip::getQueue(stream);
@@ -594,7 +593,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
   }
   amd::Command* command = new amd::Marker(*hostQueue, !kMarkerDisableFlush, eventWaitList);
   if (command == nullptr) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
   StreamCallback* cbo = new StreamCallback(stream, callback, userData, command);
 
@@ -617,14 +616,29 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
   eventWaitList.push_back(command);
   amd::Command* block_command = new amd::Marker(*hostQueue, !kMarkerDisableFlush, eventWaitList);
   if (block_command == nullptr) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
   block_command->enqueue();
   block_command->release();
   // Release the callback marker
   command->release();
 
-  HIP_RETURN(hipSuccess);
+  return hipSuccess;
+}
+
+// ================================================================================================
+hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
+                                unsigned int flags) {
+  HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags);
+  HIP_RETURN(hipStreamAddCallback_common(stream, callback, userData, flags));
+}
+
+// ================================================================================================
+hipError_t hipStreamAddCallback_spt(hipStream_t stream, hipStreamCallback_t callback, void* userData,
+                                unsigned int flags) {
+  HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags);
+  PER_THREAD_DEFAULT_STREAM(stream);
+  HIP_RETURN(hipStreamAddCallback_common(stream, callback, userData, flags));
 }
 
 // ================================================================================================