/**
* @returns #hipSuccess #hipErrorMemoryAllocation
*/
hipError_t hipMalloc(void** ptr, size_t sizeBytes)
{
HIP_INIT_API(ptr, sizeBytes);
hipError_t hip_status = hipSuccess;
auto device = ihipGetTlsDefaultDevice();
if (device) {
const unsigned am_flags = 0;
*ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags);
if (sizeBytes && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
} else {
hc::am_memtracker_update(*ptr, device->_device_index, 0);
{
LockedAccessor_DeviceCrit_t crit(device->criticalData());
if (crit->peerCnt()) {
hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
}
}
}
} else {
hip_status = hipErrorMemoryAllocation;
}
return ihipLogStatus(hip_status);
}
/**
* @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);
}
/*
* @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue (if free != NULL due to bug)S
* @warning On HCC, the free memory only accounts for memory allocated by this process and may be optimistic.
*/
hipError_t hipMemGetInfo (size_t *free, size_t *total)
{
HIP_INIT_API(free, total);
hipError_t e = hipSuccess;
ihipDevice_t * hipDevice = ihipGetTlsDefaultDevice();
if (hipDevice) {
if (total) {
*total = hipDevice->_props.totalGlobalMem;
}
if (free) {
// TODO - replace with kernel-level for reporting free memory:
size_t deviceMemSize, hostMemSize, userMemSize;
hc::am_memtracker_sizeinfo(hipDevice->_acc, &deviceMemSize, &hostMemSize, &userMemSize);
printf ("deviceMemSize=%zu\n", deviceMemSize);
*free = hipDevice->_props.totalGlobalMem - deviceMemSize;
}
} else {
e = hipErrorInvalidDevice;
}
return ihipLogStatus(e);
}
hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
{
HIP_INIT_API(ptr, sizeBytes, flags);
hipError_t hip_status = hipSuccess;
auto device = ihipGetTlsDefaultDevice();
if(device){
if(flags == hipHostMallocDefault){
*ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned);
if(sizeBytes < 1 && (*ptr == NULL)){
hip_status = hipErrorMemoryAllocation;
}else{
hc::am_memtracker_update(*ptr, device->_device_index, amHostPinned);
}
tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
} else if(flags & hipHostMallocMapped){
*ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned);
if(sizeBytes && (*ptr == NULL)){
hip_status = hipErrorMemoryAllocation;
}else{
hc::am_memtracker_update(*ptr, device->_device_index, flags);
{
LockedAccessor_DeviceCrit_t crit(device->criticalData());
if (crit->peerCnt()) {
hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
}
}
}
tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
}
}
return ihipLogStatus(hip_status);
}
/**
* @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);
}
//---
hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority) {
HIP_INIT_API(hipDeviceGetStreamPriorityRange, leastPriority, greatestPriority);
if (leastPriority != NULL) *leastPriority = priority_low;
if (greatestPriority != NULL) *greatestPriority = priority_high;
return ihipLogStatus(hipSuccess);
}
//---
hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority) {
HIP_INIT_API(hipStreamCreateWithPriority, stream, flags, priority);
// clamp priority to range [priority_high:priority_low]
priority = (priority < priority_high ? priority_high : (priority > priority_low ? priority_low : priority));
return ihipLogStatus(ihipStreamCreate(stream, flags, priority));
}
//---
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
{
HIP_INIT_API(stream, flags);
return ihipLogStatus(ihipStreamCreate(stream, flags));
}
hipError_t hipPeekAtLastError()
{
HIP_INIT_API();
// peek at last error, but don't reset it.
return ihipLogStatus(tls_lastHipError);
}
hipError_t hipGetLastError() {
HIP_INIT_API(hipGetLastError);
// Return last error, but then reset the state:
hipError_t e = ihipLogStatus(tls_lastHipError);
tls_lastHipError = hipSuccess;
return e;
}
const char *hipGetErrorString(hipError_t hip_error)
{
HIP_INIT_API(hip_error);
// TODO - return a message explaining the error.
// TODO - This should be set up to return the same string reported in the the doxygen comments, somehow.
return hipGetErrorName(hip_error);
}
//---
hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
unsigned int flags) {
HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags);
hipError_t e = hipSuccess;
// Create a thread in detached mode to handle callback
ihipStreamCallback_t* cb = new ihipStreamCallback_t(stream, callback, userData);
std::thread(ihipStreamCallbackHandler, cb).detach();
return ihipLogStatus(e);
}
//---
hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags) {
HIP_INIT_API(hipStreamGetFlags, stream, flags);
if (flags == NULL) {
return ihipLogStatus(hipErrorInvalidValue);
} else if (stream == hipStreamNull) {
return ihipLogStatus(hipErrorInvalidResourceHandle);
} else {
*flags = stream->_flags;
return ihipLogStatus(hipSuccess);
}
}
hipError_t hipMemset(void* dst, int value, size_t sizeBytes )
{
hipStream_t stream = hipStreamNull;
// TODO - call an ihip memset so HIP_TRACE is correct.
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;
}
}
cf.wait();
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);
}
/**
* @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 hipStreamGetPriority(hipStream_t stream, int* priority) {
HIP_INIT_API(hipStreamGetPriority, stream, priority);
if (priority == NULL) {
return ihipLogStatus(hipErrorInvalidValue);
} else if (stream == hipStreamNull) {
return ihipLogStatus(hipErrorInvalidResourceHandle);
} else {
#if defined(__HCC__) && (__hcc_minor__ < 3)
*priority = 0;
#else
LockedAccessor_StreamCrit_t crit(stream->_criticalData);
*priority = crit->_av.get_queue_priority();
#endif
return ihipLogStatus(hipSuccess);
}
}
//---
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 hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags)
{
HIP_INIT_API(hostPtr, sizeBytes, flags);
hipError_t hip_status = hipSuccess;
auto device = ihipGetTlsDefaultDevice();
if(hostPtr == NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hc::accelerator acc;
hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
am_status_t am_status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);
if(am_status == AM_SUCCESS){
hip_status = hipErrorHostMemoryAlreadyRegistered;
}else{
auto device = ihipGetTlsDefaultDevice();
if(hostPtr == NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
if(device){
if(flags == hipHostRegisterDefault || flags == hipHostRegisterPortable || flags == hipHostRegisterMapped){
std::vector<hc::accelerator>vecAcc;
for(int i=0;i<g_deviceCnt;i++){
vecAcc.push_back(g_devices[i]._acc);
}
#if USE_HCC_LOCK_API
am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0], vecAcc.size());
#else
am_status = AM_ERROR_MISC;
#endif
if(am_status == AM_SUCCESS){
hip_status = hipSuccess;
}else{
hip_status = hipErrorMemoryAllocation;
}
}else{
hip_status = hipErrorInvalidValue;
}
}
}
return ihipLogStatus(hip_status);
}
//---
hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind)
{
HIP_INIT_API(symbolName, src, count, offset, kind);
#ifdef USE_MEMCPYTOSYMBOL
if(kind != hipMemcpyHostToDevice)
{
return ihipLogStatus(hipErrorInvalidValue);
}
auto device = ihipGetTlsDefaultDevice();
//hsa_signal_t depSignal;
//int depSignalCnt = device._default_stream->preCopyCommand(NULL, &depSignal, ihipCommandCopyH2D);
assert(0); // Need to properly synchronize the copy - do something with depSignal if != NULL.
device->_acc.memcpy_symbol(symbolName, (void*) src,count, offset);
#endif
return ihipLogStatus(hipSuccess);
}
//---
hipError_t hipHostUnregister(void *hostPtr)
{
HIP_INIT_API(hostPtr);
auto device = ihipGetTlsDefaultDevice();
hipError_t hip_status = hipSuccess;
if(hostPtr == NULL){
hip_status = hipErrorInvalidValue;
}else{
#if USE_HCC_LOCK_API
am_status_t am_status = hc::am_memory_host_unlock(device->_acc, hostPtr);
#else
am_status_t am_status = AM_ERROR_MISC;
#endif
if(am_status != AM_SUCCESS){
hip_status = hipErrorHostMemoryNotRegistered;
}
}
return ihipLogStatus(hip_status);
}
//---
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind)
{
HIP_INIT_API(dst, src, sizeBytes, kind);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
hipError_t e = hipSuccess;
try {
stream->locked_copySync(dst, src, sizeBytes, kind);
}
catch (ihipException ex) {
e = ex._code;
}
return ihipLogStatus(e);
}
hipError_t hipHostFree(void* ptr)
{
HIP_INIT_API(ptr);
// TODO - ensure this pointer was created by hipMallocHost and not hipMalloc
std::call_once(hip_initialized, ihipInit);
hipError_t hipStatus = hipErrorInvalidDevicePointer;
if (ptr) {
hc::accelerator acc;
hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr);
if(status == AM_SUCCESS){
if(amPointerInfo._hostPointer == ptr){
hc::am_free(ptr);
hipStatus = hipSuccess;
}
}
}
return ihipLogStatus(hipStatus);
};
//---
hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr)
{
HIP_INIT_API(flagsPtr, hostPtr);
hipError_t hip_status = hipSuccess;
hc::accelerator acc;
hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);
if(status == AM_SUCCESS){
*flagsPtr = amPointerInfo._appAllocationFlags;
if(*flagsPtr == 0){
hip_status = hipErrorInvalidValue;
}
else{
hip_status = hipSuccess;
}
tprintf(DB_MEM, " %s: host ptr=%p\n", __func__, hostPtr);
}else{
hip_status = hipErrorInvalidValue;
}
return ihipLogStatus(hip_status);
}
//---
hipError_t hipFree(void* ptr)
{
HIP_INIT_API(ptr);
hipError_t hipStatus = hipErrorInvalidDevicePointer;
// Synchronize to ensure all work has finished.
ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
if (ptr) {
hc::accelerator acc;
hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr);
if(status == AM_SUCCESS){
if(amPointerInfo._hostPointer == NULL){
hc::am_free(ptr);
hipStatus = hipSuccess;
}
}
}
return ihipLogStatus(hipStatus);
}
//---
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);
}
//---
hipError_t hipStreamCreate(hipStream_t *stream)
{
HIP_INIT_API(stream);
return ihipLogStatus(ihipStreamCreate(stream, hipStreamDefault));
}
/**
* @warning : flags must be 0.
*/
hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags)
{
HIP_INIT_API(event, flags);
return ihipLogStatus(ihipEventCreate(event, flags));
}
hipError_t hipEventCreate(hipEvent_t* event)
{
HIP_INIT_API(event);
return ihipLogStatus(ihipEventCreate(event, 0));
}
//---
hipError_t hipStreamCreate(hipStream_t* stream) {
HIP_INIT_API(hipStreamCreate, stream);
return ihipLogStatus(ihipStreamCreate(stream, hipStreamDefault, priority_normal));
}
const char *hipGetErrorName(hipError_t hip_error)
{
HIP_INIT_API(hip_error);
return ihipErrorString(hip_error);
}
