/** * @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); }