/**
* @return #hipSuccess, #hipErrorInvalidResourceHandle
*/
hipError_t hipStreamDestroy(hipStream_t stream)
{
HIP_INIT_API(stream);
hipError_t e = hipSuccess;
//--- Drain the stream:
if (stream == NULL) {
ihipDevice_t *device = ihipGetTlsDefaultDevice();
device->locked_syncDefaultStream(true/*waitOnSelf*/);
} else {
stream->locked_wait();
e = hipSuccess;
}
ihipDevice_t *device = stream->getDevice();
if (device) {
device->locked_removeStream(stream);
delete stream;
} else {
e = hipErrorInvalidResourceHandle;
}
return ihipLogStatus(e);
}
/**
* @return #hipSuccess, #hipErrorInvalidResourceHandle
*/
hipError_t hipStreamDestroy(hipStream_t stream) {
HIP_INIT_API(hipStreamDestroy, stream);
hipError_t e = hipSuccess;
//--- Drain the stream:
if (stream == NULL) {
if (!HIP_FORCE_NULL_STREAM) {
e = hipErrorInvalidResourceHandle;
}
} else {
stream->locked_wait();
ihipCtx_t* ctx = stream->getCtx();
if (ctx) {
ctx->locked_removeStream(stream);
delete stream;
} else {
e = hipErrorInvalidResourceHandle;
}
}
return ihipLogStatus(e);
}
/** @return #hipErrorInvalidValue
*/
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream )
{
HIP_INIT_API(dst, value, sizeBytes, stream);
hipError_t e = hipSuccess;
stream = ihipSyncAndResolveStream(stream);
if (stream) {
stream->lockopen_preKernelCommand();
hc::completion_future cf ;
if ((sizeBytes & 0x3) == 0) {
// use a faster dword-per-workitem copy:
try {
value = value & 0xff;
unsigned value32 = (value << 24) | (value << 16) | (value << 8) | (value) ;
cf = ihipMemsetKernel<unsigned> (stream, static_cast<unsigned*> (dst), value32, sizeBytes/sizeof(unsigned));
}
catch (std::exception &ex) {
e = hipErrorInvalidValue;
}
} else {
// use a slow byte-per-workitem copy:
try {
cf = ihipMemsetKernel<char> (stream, static_cast<char*> (dst), value, sizeBytes);
}
catch (std::exception &ex) {
e = hipErrorInvalidValue;
}
}
stream->lockclose_postKernelCommand(cf);
if (HIP_LAUNCH_BLOCKING) {
tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING wait for memset [stream:%p].\n", __func__, (void*)stream);
cf.wait();
tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed [stream:%p].\n", __func__, (void*)stream);
}
} else {
e = hipErrorInvalidValue;
}
return ihipLogStatus(e);
};
//---
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
{
std::call_once(hip_initialized, ihipInit);
ihipEvent_t *eh = event._handle;
if (eh && eh->_state != hipEventStatusUnitialized) {
eh->_stream = stream;
if (stream == NULL) {
// If stream == NULL, wait on all queues.
// Behavior matches "use standard default semantics".
// TODO-HCC fix this - conservative or use device timestamps?
// TODO-HCC can we use barrier or event marker to implement better solution?
ihipDevice_t *device = ihipGetTlsDefaultDevice();
device->locked_syncDefaultStream(true);
eh->_timestamp = hc::get_system_ticks();
eh->_state = hipEventStatusRecorded;
return ihipLogStatus(hipSuccess);
} else {
eh->_state = hipEventStatusRecording;
// Clear timestamps
eh->_timestamp = 0;
eh->_marker = stream->_av.create_marker();
eh->_copy_seq_id = stream->locked_lastCopySeqId();
return ihipLogStatus(hipSuccess);
}
} else {
return ihipLogStatus(hipErrorInvalidResourceHandle);
}
}
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) {
HIP_INIT_SPECIAL_API(hipStreamWaitEvent, TRACE_SYNC, stream, event, flags);
hipError_t e = hipSuccess;
auto ecd = event->locked_copyCrit();
if (event == nullptr) {
e = hipErrorInvalidResourceHandle;
} else if ((ecd._state != hipEventStatusUnitialized) && (ecd._state != hipEventStatusCreated)) {
if (HIP_SYNC_STREAM_WAIT || (HIP_SYNC_NULL_STREAM && (stream == 0))) {
// conservative wait on host for the specified event to complete:
// return _stream->locked_eventWaitComplete(this, waitMode);
//
ecd._stream->locked_eventWaitComplete(
ecd.marker(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
: hc::hcWaitModeActive);
} else {
stream = ihipSyncAndResolveStream(stream);
// This will use create_blocking_marker to wait on the specified queue.
stream->locked_streamWaitEvent(ecd);
}
} // else event not recorded, return immediately and don't create marker.
return ihipLogStatus(e);
};
/**
* @bug This function conservatively waits for all work in the specified stream to complete.
*/
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags)
{
HIP_INIT_API(stream, event, flags);
hipError_t e = hipSuccess;
{
// TODO-hcc Convert to use create_blocking_marker(...) functionality.
// Currently we have a super-conservative version of this - block on host, and drain the queue.
// This should create a barrier packet in the target queue.
stream->locked_wait();
e = hipSuccess;
}
return ihipLogStatus(e);
};
//---
hipError_t hipStreamSynchronize(hipStream_t stream)
{
HIP_INIT_API(stream);
hipError_t e = hipSuccess;
if (stream == NULL) {
ihipDevice_t *device = ihipGetTlsDefaultDevice();
device->locked_syncDefaultStream(true/*waitOnSelf*/);
} else {
stream->locked_wait();
e = hipSuccess;
}
return ihipLogStatus(e);
};
//---
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream)
{
HIP_INIT_API(dst, src, sizeBytes, kind, stream);
hipError_t e = hipSuccess;
stream = ihipSyncAndResolveStream(stream);
if ((dst == NULL) || (src == NULL)) {
e= hipErrorInvalidValue;
} else if (stream) {
try {
stream->copyAsync(dst, src, sizeBytes, kind);
}
catch (ihipException ex) {
e = ex._code;
}
} else {
e = hipErrorInvalidValue;
}
return ihipLogStatus(e);
}
