Partition hip_hcc into sections
Separate files for different categories of HIP API.
Currently just #include into hip_hcc.cpp
[ROCm/hip commit: 655534b1ba]
이 커밋은 다음에 포함됨:
Ben Sander
@@ -0,0 +1,265 @@
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//Devices
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//---
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
hipError_t hipGetDevice(int *device)
|
||||
{
|
||||
HIP_INIT_API(device);
|
||||
|
||||
*device = tls_defaultDevice;
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorNoDevice
|
||||
*/
|
||||
hipError_t hipGetDeviceCount(int *count)
|
||||
{
|
||||
HIP_INIT_API(count);
|
||||
|
||||
*count = g_deviceCnt;
|
||||
|
||||
if (*count > 0) {
|
||||
return ihipLogStatus(hipSuccess);
|
||||
} else {
|
||||
return ihipLogStatus(hipErrorNoDevice);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
/**
|
||||
* @returns #hipSuccess
|
||||
*/
|
||||
hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig )
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
// Nop, AMD does not support variable cache configs.
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
/**
|
||||
* @returns #hipSuccess
|
||||
*/
|
||||
hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig )
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
*cacheConfig = hipFuncCachePreferNone;
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
/**
|
||||
* @returns #hipSuccess
|
||||
*/
|
||||
hipError_t hipFuncSetCacheConfig ( hipFuncCache cacheConfig )
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
// Nop, AMD does not support variable cache configs.
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
|
||||
//---
|
||||
/**
|
||||
* @returns #hipSuccess
|
||||
*/
|
||||
hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config )
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
// Nop, AMD does not support variable shared mem configs.
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
|
||||
//---
|
||||
/**
|
||||
* @returns #hipSuccess
|
||||
*/
|
||||
hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig )
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
*pConfig = hipSharedMemBankSizeFourByte;
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
//---
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidDevice
|
||||
*/
|
||||
hipError_t hipSetDevice(int device)
|
||||
{
|
||||
HIP_INIT_API(device);
|
||||
|
||||
if ((device < 0) || (device >= g_deviceCnt)) {
|
||||
return ihipLogStatus(hipErrorInvalidDevice);
|
||||
} else {
|
||||
tls_defaultDevice = device;
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
hipError_t hipDeviceSynchronize(void)
|
||||
{
|
||||
HIP_INIT_API();
|
||||
|
||||
ihipGetTlsDefaultDevice()->waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
|
||||
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
/**
|
||||
* @return @ref hipSuccess
|
||||
*/
|
||||
hipError_t hipDeviceReset(void)
|
||||
{
|
||||
HIP_INIT_API();
|
||||
|
||||
ihipDevice_t *device = ihipGetTlsDefaultDevice();
|
||||
|
||||
// TODO-HCC
|
||||
// This function currently does a user-level cleanup of known resources.
|
||||
// It could benefit from KFD support to perform a more "nuclear" clean that would include any associated kernel resources and page table entries.
|
||||
|
||||
|
||||
if (device) {
|
||||
//---
|
||||
//Wait for pending activity to complete?
|
||||
//TODO - check if this is required behavior:
|
||||
for (auto streamI=device->_streams.begin(); streamI!=device->_streams.end(); streamI++) {
|
||||
ihipStream_t *stream = *streamI;
|
||||
stream->wait();
|
||||
}
|
||||
|
||||
// Release device resources (streams and memory):
|
||||
device->reset();
|
||||
}
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
/**
|
||||
*
|
||||
*/
|
||||
hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
ihipDevice_t * hipDevice = ihipGetDevice(device);
|
||||
hipDeviceProp_t *prop = &hipDevice->_props;
|
||||
if (hipDevice) {
|
||||
switch (attr) {
|
||||
case hipDeviceAttributeMaxThreadsPerBlock:
|
||||
*pi = prop->maxThreadsPerBlock; break;
|
||||
case hipDeviceAttributeMaxBlockDimX:
|
||||
*pi = prop->maxThreadsDim[0]; break;
|
||||
case hipDeviceAttributeMaxBlockDimY:
|
||||
*pi = prop->maxThreadsDim[1]; break;
|
||||
case hipDeviceAttributeMaxBlockDimZ:
|
||||
*pi = prop->maxThreadsDim[2]; break;
|
||||
case hipDeviceAttributeMaxGridDimX:
|
||||
*pi = prop->maxGridSize[0]; break;
|
||||
case hipDeviceAttributeMaxGridDimY:
|
||||
*pi = prop->maxGridSize[1]; break;
|
||||
case hipDeviceAttributeMaxGridDimZ:
|
||||
*pi = prop->maxGridSize[2]; break;
|
||||
case hipDeviceAttributeMaxSharedMemoryPerBlock:
|
||||
*pi = prop->sharedMemPerBlock; break;
|
||||
case hipDeviceAttributeTotalConstantMemory:
|
||||
*pi = prop->totalConstMem; break;
|
||||
case hipDeviceAttributeWarpSize:
|
||||
*pi = prop->warpSize; break;
|
||||
case hipDeviceAttributeMaxRegistersPerBlock:
|
||||
*pi = prop->regsPerBlock; break;
|
||||
case hipDeviceAttributeClockRate:
|
||||
*pi = prop->clockRate; break;
|
||||
case hipDeviceAttributeMemoryClockRate:
|
||||
*pi = prop->memoryClockRate; break;
|
||||
case hipDeviceAttributeMemoryBusWidth:
|
||||
*pi = prop->memoryBusWidth; break;
|
||||
case hipDeviceAttributeMultiprocessorCount:
|
||||
*pi = prop->multiProcessorCount; break;
|
||||
case hipDeviceAttributeComputeMode:
|
||||
*pi = prop->computeMode; break;
|
||||
case hipDeviceAttributeL2CacheSize:
|
||||
*pi = prop->l2CacheSize; break;
|
||||
case hipDeviceAttributeMaxThreadsPerMultiProcessor:
|
||||
*pi = prop->maxThreadsPerMultiProcessor; break;
|
||||
case hipDeviceAttributeComputeCapabilityMajor:
|
||||
*pi = prop->major; break;
|
||||
case hipDeviceAttributeComputeCapabilityMinor:
|
||||
*pi = prop->minor; break;
|
||||
case hipDeviceAttributePciBusId:
|
||||
*pi = prop->pciBusID; break;
|
||||
case hipDeviceAttributeConcurrentKernels:
|
||||
*pi = prop->concurrentKernels; break;
|
||||
case hipDeviceAttributePciDeviceId:
|
||||
*pi = prop->pciDeviceID; break;
|
||||
case hipDeviceAttributeMaxSharedMemoryPerMultiprocessor:
|
||||
*pi = prop->maxSharedMemoryPerMultiProcessor; break;
|
||||
case hipDeviceAttributeIsMultiGpuBoard:
|
||||
*pi = prop->isMultiGpuBoard; break;
|
||||
default:
|
||||
e = hipErrorInvalidValue; break;
|
||||
}
|
||||
} else {
|
||||
e = hipErrorInvalidDevice;
|
||||
}
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidDevice
|
||||
* @bug HCC always returns 0 for maxThreadsPerMultiProcessor
|
||||
* @bug HCC always returns 0 for regsPerBlock
|
||||
* @bug HCC always returns 0 for l2CacheSize
|
||||
*/
|
||||
hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device)
|
||||
{
|
||||
HIP_INIT_API(props, device);
|
||||
|
||||
hipError_t e;
|
||||
|
||||
ihipDevice_t * hipDevice = ihipGetDevice(device);
|
||||
if (hipDevice) {
|
||||
// copy saved props
|
||||
*props = hipDevice->_props;
|
||||
e = hipSuccess;
|
||||
} else {
|
||||
e = hipErrorInvalidDevice;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
@@ -0,0 +1,75 @@
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
// Error Handling
|
||||
//---
|
||||
/**
|
||||
* @returns return code from last HIP called from the active host thread.
|
||||
*/
|
||||
hipError_t hipGetLastError()
|
||||
{
|
||||
HIP_INIT_API();
|
||||
|
||||
// Return last error, but then reset the state:
|
||||
hipError_t e = ihipLogStatus(tls_lastHipError);
|
||||
tls_lastHipError = hipSuccess;
|
||||
return e;
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
hipError_t hipPeakAtLastError()
|
||||
{
|
||||
HIP_INIT_API();
|
||||
|
||||
|
||||
// peak at last error, but don't reset it.
|
||||
return ihipLogStatus(tls_lastHipError);
|
||||
}
|
||||
|
||||
|
||||
const char *ihipErrorString(hipError_t hip_error)
|
||||
{
|
||||
switch (hip_error) {
|
||||
case hipSuccess : return "hipSuccess";
|
||||
case hipErrorMemoryAllocation : return "hipErrorMemoryAllocation";
|
||||
case hipErrorMemoryFree : return "hipErrorMemoryFree";
|
||||
case hipErrorUnknownSymbol : return "hipErrorUnknownSymbol";
|
||||
case hipErrorOutOfResources : return "hipErrorOutOfResources";
|
||||
case hipErrorInvalidValue : return "hipErrorInvalidValue";
|
||||
case hipErrorInvalidResourceHandle : return "hipErrorInvalidResourceHandle";
|
||||
case hipErrorInvalidDevice : return "hipErrorInvalidDevice";
|
||||
case hipErrorInvalidMemcpyDirection : return "hipErrorInvalidMemcpyDirection";
|
||||
case hipErrorNoDevice : return "hipErrorNoDevice";
|
||||
case hipErrorNotReady : return "hipErrorNotReady";
|
||||
case hipErrorRuntimeMemory : return "hipErrorRuntimeMemory";
|
||||
case hipErrorRuntimeOther : return "hipErrorRuntimeOther";
|
||||
case hipErrorUnknown : return "hipErrorUnknown";
|
||||
case hipErrorTbd : return "hipErrorTbd";
|
||||
default : return "hipErrorUnknown";
|
||||
};
|
||||
};
|
||||
|
||||
|
||||
|
||||
//---
|
||||
const char *hipGetErrorName(hipError_t hip_error)
|
||||
{
|
||||
HIP_INIT_API(hip_error);
|
||||
|
||||
return ihipErrorString(hip_error);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* @warning : hipGetErrorString returns string from hipGetErrorName
|
||||
*/
|
||||
|
||||
//---
|
||||
const char *hipGetErrorString(hipError_t hip_error)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
// 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);
|
||||
}
|
||||
@@ -0,0 +1,197 @@
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
// Events
|
||||
//---
|
||||
/**
|
||||
* @warning : flags must be 0.
|
||||
*/
|
||||
hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags)
|
||||
{
|
||||
// TODO - support hipEventDefault, hipEventBlockingSync, hipEventDisableTiming
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
if (flags == 0) {
|
||||
ihipEvent_t *eh = event->_handle = new ihipEvent_t();
|
||||
|
||||
eh->_state = hipEventStatusCreated;
|
||||
eh->_stream = NULL;
|
||||
eh->_flags = flags;
|
||||
eh->_timestamp = 0;
|
||||
eh->_copy_seq_id = 0;
|
||||
} 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.
|
||||
// This matches behavior described in CUDA 7 RT APIs, which say that "This function uses standard default stream semantics".
|
||||
// TODO-HCC fix this - is CUDA this conservative or still uses device timestamps?
|
||||
// TODO-HCC can we use barrier or event marker to implement better solution?
|
||||
ihipDevice_t *device = ihipGetTlsDefaultDevice();
|
||||
device->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->lastCopySeqId();
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
} else {
|
||||
return ihipLogStatus(hipErrorInvalidResourceHandle);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
hipError_t hipEventDestroy(hipEvent_t event)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
event._handle->_state = hipEventStatusUnitialized;
|
||||
|
||||
delete event._handle;
|
||||
event._handle = NULL;
|
||||
|
||||
// TODO - examine return additional error codes
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
hipError_t hipEventSynchronize(hipEvent_t event)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
ihipEvent_t *eh = event._handle;
|
||||
|
||||
if (eh) {
|
||||
if (eh->_state == hipEventStatusUnitialized) {
|
||||
return ihipLogStatus(hipErrorInvalidResourceHandle);
|
||||
} else if (eh->_state == hipEventStatusCreated ) {
|
||||
// Created but not actually recorded on any device:
|
||||
return ihipLogStatus(hipSuccess);
|
||||
} else if (eh->_stream == NULL) {
|
||||
ihipDevice_t *device = ihipGetTlsDefaultDevice();
|
||||
device->syncDefaultStream(true);
|
||||
return ihipLogStatus(hipSuccess);
|
||||
} else {
|
||||
#if __hcc_workweek__ >= 16033
|
||||
eh->_marker.wait((eh->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
|
||||
#else
|
||||
eh->_marker.wait();
|
||||
#endif
|
||||
eh->_stream->reclaimSignals_ts(eh->_copy_seq_id);
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
} else {
|
||||
return ihipLogStatus(hipErrorInvalidResourceHandle);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void ihipSetTs(hipEvent_t e)
|
||||
{
|
||||
ihipEvent_t *eh = e._handle;
|
||||
if (eh->_state == hipEventStatusRecorded) {
|
||||
// already recorded, done:
|
||||
return;
|
||||
} else {
|
||||
// TODO - use completion-future functions to obtain ticks and timestamps:
|
||||
hsa_signal_t *sig = static_cast<hsa_signal_t*> (eh->_marker.get_native_handle());
|
||||
if (sig) {
|
||||
if (hsa_signal_load_acquire(*sig) == 0) {
|
||||
eh->_timestamp = eh->_marker.get_end_tick();
|
||||
eh->_state = hipEventStatusRecorded;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
ihipEvent_t *start_eh = start._handle;
|
||||
ihipEvent_t *stop_eh = stop._handle;
|
||||
|
||||
ihipSetTs(start);
|
||||
ihipSetTs(stop);
|
||||
|
||||
hipError_t status = hipSuccess;
|
||||
*ms = 0.0f;
|
||||
|
||||
if (start_eh && stop_eh) {
|
||||
if ((start_eh->_state == hipEventStatusRecorded) && (stop_eh->_state == hipEventStatusRecorded)) {
|
||||
// Common case, we have good information for both events.
|
||||
|
||||
int64_t tickDiff = (stop_eh->_timestamp - start_eh->_timestamp);
|
||||
|
||||
// TODO-move this to a variable saved with each agent.
|
||||
uint64_t freqHz;
|
||||
hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz);
|
||||
if (freqHz) {
|
||||
*ms = ((double)(tickDiff) / (double)(freqHz)) * 1000.0f;
|
||||
status = hipSuccess;
|
||||
} else {
|
||||
* ms = 0.0f;
|
||||
status = hipErrorInvalidValue;
|
||||
}
|
||||
|
||||
|
||||
} else if ((start_eh->_state == hipEventStatusRecording) ||
|
||||
(stop_eh->_state == hipEventStatusRecording)) {
|
||||
status = hipErrorNotReady;
|
||||
} else if ((start_eh->_state == hipEventStatusUnitialized) ||
|
||||
(stop_eh->_state == hipEventStatusUnitialized)) {
|
||||
status = hipErrorInvalidResourceHandle;
|
||||
}
|
||||
}
|
||||
|
||||
return ihipLogStatus(status);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
hipError_t hipEventQuery(hipEvent_t event)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
ihipEvent_t *eh = event._handle;
|
||||
|
||||
// TODO-stream - need to read state of signal here: The event may have become ready after recording..
|
||||
// TODO-HCC - use get_hsa_signal here.
|
||||
|
||||
if (eh->_state == hipEventStatusRecording) {
|
||||
return ihipLogStatus(hipErrorNotReady);
|
||||
} else {
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
+7
-1559
파일 크기가 너무 크기때문에 변경 상태를 표시하지 않습니다.
Diff 로드
파일 크기가 너무 크기때문에 변경 상태를 표시하지 않습니다.
Diff 로드
@@ -0,0 +1,781 @@
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
// Memory
|
||||
//
|
||||
//
|
||||
//
|
||||
|
||||
//---
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice
|
||||
*/
|
||||
hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
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) {
|
||||
|
||||
attributes->memoryType = amPointerInfo._isInDeviceMem ? hipMemoryTypeDevice: hipMemoryTypeHost;
|
||||
attributes->hostPointer = amPointerInfo._hostPointer;
|
||||
attributes->devicePointer = amPointerInfo._devicePointer;
|
||||
attributes->isManaged = 0;
|
||||
if(attributes->memoryType == hipMemoryTypeHost){
|
||||
attributes->hostPointer = ptr;
|
||||
}
|
||||
if(attributes->memoryType == hipMemoryTypeDevice){
|
||||
attributes->devicePointer = ptr;
|
||||
}
|
||||
attributes->allocationFlags = amPointerInfo._appAllocationFlags;
|
||||
attributes->device = amPointerInfo._appId;
|
||||
|
||||
if (attributes->device < 0) {
|
||||
e = hipErrorInvalidDevice;
|
||||
}
|
||||
|
||||
|
||||
} else {
|
||||
attributes->memoryType = hipMemoryTypeDevice;
|
||||
attributes->hostPointer = 0;
|
||||
attributes->devicePointer = 0;
|
||||
attributes->device = -1;
|
||||
attributes->isManaged = 0;
|
||||
attributes->allocationFlags = 0;
|
||||
|
||||
e = hipErrorUnknown; // TODO - should be hipErrorInvalidValue ?
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* @returns #hipSuccess,
|
||||
* @returns #hipErrorInvalidValue if flags are not 0
|
||||
* @returns #hipErrorMemoryAllocation if hostPointer is not a tracked allocation.
|
||||
*/
|
||||
hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsigned flags)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
// Flags must be 0:
|
||||
if (flags != 0) {
|
||||
e = hipErrorInvalidValue;
|
||||
} else {
|
||||
hc::accelerator acc;
|
||||
hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPointer);
|
||||
if (status == AM_SUCCESS) {
|
||||
*devicePointer = amPointerInfo._devicePointer;
|
||||
} else {
|
||||
e = hipErrorMemoryAllocation;
|
||||
*devicePointer = NULL;
|
||||
}
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
|
||||
|
||||
// kernel for launching memcpy operations:
|
||||
template <typename T>
|
||||
hc::completion_future
|
||||
ihipMemcpyKernel(hipStream_t stream, T * c, const T * a, size_t sizeBytes)
|
||||
{
|
||||
int wg = std::min((unsigned)8, stream->getDevice()->_compute_units);
|
||||
const int threads_per_wg = 256;
|
||||
|
||||
int threads = wg * threads_per_wg;
|
||||
if (threads > sizeBytes) {
|
||||
threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg;
|
||||
}
|
||||
|
||||
|
||||
hc::extent<1> ext(threads);
|
||||
auto ext_tile = ext.tile(threads_per_wg);
|
||||
|
||||
hc::completion_future cf =
|
||||
hc::parallel_for_each(
|
||||
stream->_av,
|
||||
ext_tile,
|
||||
[=] (hc::tiled_index<1> idx)
|
||||
__attribute__((hc))
|
||||
{
|
||||
int offset = amp_get_global_id(0);
|
||||
// TODO-HCC - change to hc_get_local_size()
|
||||
int stride = amp_get_local_size(0) * hc_get_num_groups(0) ;
|
||||
|
||||
for (int i=offset; i<sizeBytes; i+=stride) {
|
||||
c[i] = a[i];
|
||||
}
|
||||
});
|
||||
|
||||
return cf;
|
||||
}
|
||||
|
||||
|
||||
// kernel for launching memset operations:
|
||||
template <typename T>
|
||||
hc::completion_future
|
||||
ihipMemsetKernel(hipStream_t stream, T * ptr, T val, size_t sizeBytes)
|
||||
{
|
||||
int wg = std::min((unsigned)8, stream->getDevice()->_compute_units);
|
||||
const int threads_per_wg = 256;
|
||||
|
||||
int threads = wg * threads_per_wg;
|
||||
if (threads > sizeBytes) {
|
||||
threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg;
|
||||
}
|
||||
|
||||
|
||||
hc::extent<1> ext(threads);
|
||||
auto ext_tile = ext.tile(threads_per_wg);
|
||||
|
||||
hc::completion_future cf =
|
||||
hc::parallel_for_each(
|
||||
stream->_av,
|
||||
ext_tile,
|
||||
[=] (hc::tiled_index<1> idx)
|
||||
__attribute__((hc))
|
||||
{
|
||||
int offset = amp_get_global_id(0);
|
||||
// TODO-HCC - change to hc_get_local_size()
|
||||
int stride = amp_get_local_size(0) * hc_get_num_groups(0) ;
|
||||
|
||||
for (int i=offset; i<sizeBytes; i+=stride) {
|
||||
ptr[i] = val;
|
||||
}
|
||||
});
|
||||
|
||||
return cf;
|
||||
|
||||
}
|
||||
|
||||
//---
|
||||
/**
|
||||
* @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);
|
||||
}
|
||||
} else {
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
}
|
||||
|
||||
return ihipLogStatus(hip_status);
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipMallocHost(void** ptr, size_t sizeBytes)
|
||||
{
|
||||
HIP_INIT_API(ptr, sizeBytes);
|
||||
|
||||
hipError_t hip_status = hipSuccess;
|
||||
|
||||
const unsigned am_flags = amHostPinned;
|
||||
auto device = ihipGetTlsDefaultDevice();
|
||||
|
||||
if (device) {
|
||||
*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);
|
||||
}
|
||||
|
||||
tprintf (DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
|
||||
}
|
||||
|
||||
return ihipLogStatus(hip_status);
|
||||
}
|
||||
|
||||
|
||||
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 && (*ptr == NULL)){
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
}else{
|
||||
hc::am_memtracker_update(*ptr, device->_device_index, 0);
|
||||
}
|
||||
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);
|
||||
}
|
||||
tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
|
||||
}
|
||||
}
|
||||
return ihipLogStatus(hip_status);
|
||||
}
|
||||
|
||||
|
||||
// TODO - remove me, this is deprecated.
|
||||
hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags)
|
||||
{
|
||||
return hipHostMalloc(ptr, sizeBytes, flags);
|
||||
};
|
||||
|
||||
|
||||
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 hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags)
|
||||
{
|
||||
HIP_INIT_API(hostPtr, sizeBytes, flags);
|
||||
|
||||
hipError_t hip_status = hipSuccess;
|
||||
|
||||
auto device = ihipGetTlsDefaultDevice();
|
||||
void* srcPtr;
|
||||
if(hostPtr == NULL){
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
if(device){
|
||||
if(flags == hipHostRegisterDefault){
|
||||
hsa_status_t hsa_status = hsa_amd_memory_lock(hostPtr, sizeBytes, &device->_hsa_agent, 1, &srcPtr);
|
||||
if(hsa_status == HSA_STATUS_SUCCESS){
|
||||
hip_status = hipSuccess;
|
||||
}else{
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
}
|
||||
}
|
||||
else if (flags | hipHostRegisterMapped){
|
||||
hsa_status_t hsa_status = hsa_amd_memory_lock(hostPtr, sizeBytes, &device->_hsa_agent, 1, &srcPtr);
|
||||
//TODO: Added feature for actual host pointer being tracked
|
||||
if(hsa_status != HSA_STATUS_SUCCESS){
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
}
|
||||
}
|
||||
}
|
||||
return ihipLogStatus(hip_status);
|
||||
}
|
||||
|
||||
hipError_t hipHostUnregister(void *hostPtr)
|
||||
{
|
||||
HIP_INIT_API(hostPtr);
|
||||
|
||||
hipError_t hip_status = hipSuccess;
|
||||
if(hostPtr == NULL){
|
||||
hip_status = hipErrorInvalidValue;
|
||||
}else{
|
||||
hsa_status_t hsa_status = hsa_amd_memory_unlock(hostPtr);
|
||||
if(hsa_status != HSA_STATUS_SUCCESS){
|
||||
hip_status = hipErrorInvalidValue;
|
||||
// TODO: Add a different return error. This is not true
|
||||
}
|
||||
}
|
||||
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);
|
||||
}
|
||||
|
||||
|
||||
// Resolve hipMemcpyDefault to a known type.
|
||||
unsigned ihipStream_t::resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem)
|
||||
{
|
||||
hipMemcpyKind kind = hipMemcpyDefault;
|
||||
|
||||
if (!srcInDeviceMem && !dstInDeviceMem) {
|
||||
kind = hipMemcpyHostToHost;
|
||||
} else if (!srcInDeviceMem && dstInDeviceMem) {
|
||||
kind = hipMemcpyHostToDevice;
|
||||
} else if (srcInDeviceMem && !dstInDeviceMem) {
|
||||
kind = hipMemcpyDeviceToHost;
|
||||
} else if (srcInDeviceMem && dstInDeviceMem) {
|
||||
kind = hipMemcpyDeviceToDevice;
|
||||
}
|
||||
|
||||
assert (kind != hipMemcpyDefault);
|
||||
|
||||
return kind;
|
||||
}
|
||||
|
||||
|
||||
// Setup the copyCommandType and the copy agents (for hsa_amd_memory_async_copy)
|
||||
void ihipStream_t::setCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent)
|
||||
{
|
||||
ihipDevice_t *device = this->getDevice();
|
||||
hsa_agent_t deviceAgent = device->_hsa_agent;
|
||||
|
||||
switch (kind) {
|
||||
case hipMemcpyHostToHost : *commandType = ihipCommandCopyH2H; *srcAgent=g_cpu_agent; *dstAgent=g_cpu_agent; break;
|
||||
case hipMemcpyHostToDevice : *commandType = ihipCommandCopyH2D; *srcAgent=g_cpu_agent; *dstAgent=deviceAgent; break;
|
||||
case hipMemcpyDeviceToHost : *commandType = ihipCommandCopyD2H; *srcAgent=deviceAgent; *dstAgent=g_cpu_agent; break;
|
||||
case hipMemcpyDeviceToDevice : *commandType = ihipCommandCopyD2D; *srcAgent=deviceAgent; *dstAgent=deviceAgent; break;
|
||||
default: throw ihipException(hipErrorInvalidMemcpyDirection);
|
||||
};
|
||||
}
|
||||
|
||||
|
||||
void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
|
||||
{
|
||||
ihipDevice_t *device = this->getDevice();
|
||||
|
||||
if (device == NULL) {
|
||||
throw ihipException(hipErrorInvalidDevice);
|
||||
}
|
||||
|
||||
hc::accelerator acc;
|
||||
hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
|
||||
bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
|
||||
bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
|
||||
|
||||
|
||||
// Resolve default to a specific Kind so we know which algorithm to use:
|
||||
if (kind == hipMemcpyDefault) {
|
||||
bool srcInDeviceMem = (srcTracked && srcPtrInfo._isInDeviceMem);
|
||||
bool dstInDeviceMem = (dstTracked && dstPtrInfo._isInDeviceMem);
|
||||
kind = resolveMemcpyDirection(srcInDeviceMem, dstInDeviceMem);
|
||||
};
|
||||
|
||||
hsa_signal_t depSignal;
|
||||
|
||||
if ((kind == hipMemcpyHostToDevice) && (!srcTracked)) {
|
||||
int depSignalCnt = preCopyCommand(NULL, &depSignal, ihipCommandCopyH2D);
|
||||
if (HIP_STAGING_BUFFERS) {
|
||||
tprintf(DB_COPY1, "D2H && !dstTracked: staged copy H2D dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
|
||||
if (HIP_PININPLACE) {
|
||||
device->_staging_buffer[0]->CopyHostToDevicePinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
|
||||
} else {
|
||||
device->_staging_buffer[0]->CopyHostToDevice(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
|
||||
}
|
||||
|
||||
// The copy waits for inputs and then completes before returning so can reset queue to empty:
|
||||
this->wait(true);
|
||||
} else {
|
||||
// TODO - remove, slow path.
|
||||
tprintf(DB_COPY1, "H2D && ! srcTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
#if USE_AV_COPY
|
||||
_av.copy(src,dst,sizeBytes);
|
||||
#else
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
#endif
|
||||
}
|
||||
} else if ((kind == hipMemcpyDeviceToHost) && (!dstTracked)) {
|
||||
int depSignalCnt = preCopyCommand(NULL, &depSignal, ihipCommandCopyD2H);
|
||||
if (HIP_STAGING_BUFFERS) {
|
||||
tprintf(DB_COPY1, "D2H && !dstTracked: staged copy D2H dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
//printf ("staged-copy- read dep signals\n");
|
||||
device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
|
||||
|
||||
// The copy waits for inputs and then completes before returning so can reset queue to empty:
|
||||
this->wait(true);
|
||||
|
||||
} else {
|
||||
// TODO - remove, slow path.
|
||||
tprintf(DB_COPY1, "D2H && !dstTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
#if USE_AV_COPY
|
||||
_av.copy(src, dst, sizeBytes);
|
||||
#else
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
#endif
|
||||
}
|
||||
} else if (kind == hipMemcpyHostToHost) {
|
||||
int depSignalCnt = preCopyCommand(NULL, &depSignal, ihipCommandCopyH2H);
|
||||
|
||||
if (depSignalCnt) {
|
||||
// host waits before doing host memory copy.
|
||||
hsa_signal_wait_acquire(depSignal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
|
||||
}
|
||||
tprintf(DB_COPY1, "H2H memcpy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
memcpy(dst, src, sizeBytes);
|
||||
|
||||
} else {
|
||||
// If not special case - these can all be handled by the hsa async copy:
|
||||
ihipCommand_t commandType;
|
||||
hsa_agent_t srcAgent, dstAgent;
|
||||
setCopyAgents(kind, &commandType, &srcAgent, &dstAgent);
|
||||
|
||||
int depSignalCnt = preCopyCommand(NULL, &depSignal, commandType);
|
||||
|
||||
// Get a completion signal:
|
||||
ihipSignal_t *ihipSignal = allocSignal();
|
||||
hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal;
|
||||
|
||||
hsa_signal_store_relaxed(copyCompleteSignal, 1);
|
||||
|
||||
tprintf(DB_COPY1, "HSA Async_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
|
||||
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, copyCompleteSignal);
|
||||
|
||||
// This is sync copy, so let's wait for copy right here:
|
||||
if (hsa_status == HSA_STATUS_SUCCESS) {
|
||||
waitCopy(ihipSignal); // wait for copy, and return to pool.
|
||||
} else {
|
||||
throw ihipException(hipErrorInvalidValue);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
|
||||
{
|
||||
ihipDevice_t *device = this->getDevice();
|
||||
|
||||
if (device == NULL) {
|
||||
throw ihipException(hipErrorInvalidDevice);
|
||||
}
|
||||
|
||||
if (kind == hipMemcpyHostToHost) {
|
||||
tprintf (DB_COPY2, "Asyc: H2H with memcpy");
|
||||
|
||||
// TODO - consider if we want to perhaps use the GPU SDMA engines anyway, to avoid the host-side sync here and keep everything flowing on the GPU.
|
||||
/* As this is a CPU op, we need to wait until all
|
||||
the commands in current stream are finished.
|
||||
*/
|
||||
this->wait();
|
||||
|
||||
memcpy(dst, src, sizeBytes);
|
||||
|
||||
} else {
|
||||
bool trueAsync = true;
|
||||
|
||||
hc::accelerator acc;
|
||||
hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
|
||||
bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
|
||||
|
||||
|
||||
// "tracked" really indicates if the pointer's virtual address is available in the GPU address space.
|
||||
// If both pointers are not tracked, we need to fall back to a sync copy.
|
||||
if (!dstTracked || !srcTracked) {
|
||||
trueAsync = false;
|
||||
}
|
||||
|
||||
if (kind == hipMemcpyDefault) {
|
||||
bool srcInDeviceMem = (srcTracked && srcPtrInfo._isInDeviceMem);
|
||||
bool dstInDeviceMem = (dstTracked && dstPtrInfo._isInDeviceMem);
|
||||
kind = resolveMemcpyDirection(srcInDeviceMem, dstInDeviceMem);
|
||||
}
|
||||
|
||||
|
||||
|
||||
ihipSignal_t *ihip_signal = allocSignal();
|
||||
hsa_signal_store_relaxed(ihip_signal->_hsa_signal, 1);
|
||||
|
||||
|
||||
if(trueAsync == true){
|
||||
|
||||
ihipCommand_t commandType;
|
||||
hsa_agent_t srcAgent, dstAgent;
|
||||
setCopyAgents(kind, &commandType, &srcAgent, &dstAgent);
|
||||
|
||||
hsa_signal_t depSignal;
|
||||
int depSignalCnt = preCopyCommand(ihip_signal, &depSignal, commandType);
|
||||
|
||||
tprintf (DB_SYNC, " copy-async, waitFor=%lu completion=#%lu(%lu)\n", depSignalCnt? depSignal.handle:0x0, ihip_signal->_sig_id, ihip_signal->_hsa_signal.handle);
|
||||
|
||||
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, ihip_signal->_hsa_signal);
|
||||
|
||||
|
||||
if (hsa_status == HSA_STATUS_SUCCESS) {
|
||||
if (HIP_LAUNCH_BLOCKING) {
|
||||
tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes);
|
||||
this->wait();
|
||||
}
|
||||
} else {
|
||||
// This path can be hit if src or dst point to unpinned host memory.
|
||||
// TODO-stream - does async-copy fall back to sync if input pointers are not pinned?
|
||||
throw ihipException(hipErrorInvalidValue);
|
||||
}
|
||||
} else {
|
||||
copySync(dst, src, sizeBytes, kind);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
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->copySync(dst, src, sizeBytes, kind);
|
||||
}
|
||||
catch (ihipException ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
|
||||
if (HIP_LAUNCH_BLOCKING) {
|
||||
tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpy\n");
|
||||
stream->wait();
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* @result #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidMemcpyDirection,
|
||||
* @result #hipErrorInvalidValue : If dst==NULL or src==NULL, or other bad argument.
|
||||
* @warning on HCC hipMemcpyAsync does not support overlapped H2D and D2H copies.
|
||||
* @warning on HCC hipMemcpyAsync requires that any host pointers are pinned (ie via the hipMallocHost call).
|
||||
*/
|
||||
//---
|
||||
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);
|
||||
}
|
||||
|
||||
|
||||
// TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
|
||||
/** @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);
|
||||
stream->preKernelCommand();
|
||||
|
||||
if (stream) {
|
||||
|
||||
hc::completion_future cf ;
|
||||
|
||||
if ((sizeBytes & 0x3) == 0) {
|
||||
// use a faster word-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->postKernelCommand(cf);
|
||||
|
||||
|
||||
if (HIP_LAUNCH_BLOCKING) {
|
||||
tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING wait for completion [stream:%p].\n", __func__, (void*)stream);
|
||||
cf.wait();
|
||||
tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING completed [stream:%p].\n", __func__, (void*)stream);
|
||||
}
|
||||
} else {
|
||||
e = hipErrorInvalidValue;
|
||||
}
|
||||
|
||||
|
||||
return ihipLogStatus(e);
|
||||
};
|
||||
|
||||
|
||||
hipError_t hipMemset(void* dst, int value, size_t sizeBytes )
|
||||
{
|
||||
HIP_INIT_API(dst, value, sizeBytes);
|
||||
|
||||
// TODO - call an ihip memset so HIP_TRACE is correct.
|
||||
return hipMemsetAsync(dst, value, sizeBytes, hipStreamNull);
|
||||
}
|
||||
|
||||
|
||||
/*
|
||||
* @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 hipFree(void* ptr)
|
||||
{
|
||||
HIP_INIT_API(ptr);
|
||||
|
||||
hipError_t hipStatus = hipErrorInvalidDevicePointer;
|
||||
|
||||
// Synchronize to ensure all work has finished.
|
||||
ihipGetTlsDefaultDevice()->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 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);
|
||||
};
|
||||
|
||||
|
||||
// TODO - deprecated function.
|
||||
hipError_t hipFreeHost(void* ptr)
|
||||
{
|
||||
return hipHostFree(ptr);
|
||||
}
|
||||
|
||||
|
||||
@@ -0,0 +1,60 @@
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipDriverGetVersion(int *driverVersion)
|
||||
{
|
||||
HIP_INIT_API(driverVersion);
|
||||
|
||||
if (driverVersion) {
|
||||
*driverVersion = 4;
|
||||
}
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
// HCC-specific accessor functions:
|
||||
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidDevice
|
||||
*/
|
||||
//---
|
||||
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
ihipDevice_t *d = ihipGetDevice(deviceId);
|
||||
hipError_t err;
|
||||
if (d == NULL) {
|
||||
err = hipErrorInvalidDevice;
|
||||
} else {
|
||||
*acc = d->_acc;
|
||||
err = hipSuccess;
|
||||
}
|
||||
return ihipLogStatus(err);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
if (stream == hipStreamNull ) {
|
||||
ihipDevice_t *device = ihipGetTlsDefaultDevice();
|
||||
stream = device->_default_stream;
|
||||
}
|
||||
|
||||
*av = &(stream->_av);
|
||||
|
||||
hipError_t err = hipSuccess;
|
||||
return ihipLogStatus(err);
|
||||
}
|
||||
@@ -0,0 +1,58 @@
|
||||
|
||||
/**
|
||||
* @warning HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P
|
||||
*/
|
||||
//---
|
||||
hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int device, int peerDevice )
|
||||
{
|
||||
HIP_INIT_API(canAccessPeer, device, peerDevice);
|
||||
|
||||
*canAccessPeer = false;
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* @warning Need to update this function when RT supports P2P
|
||||
*/
|
||||
//---
|
||||
hipError_t hipDeviceDisablePeerAccess ( int peerDevice )
|
||||
{
|
||||
HIP_INIT_API(peerDevice);
|
||||
|
||||
// TODO-p2p
|
||||
return ihipLogStatus(hipSuccess);
|
||||
};
|
||||
|
||||
|
||||
/**
|
||||
* @warning Need to update this function when RT supports P2P
|
||||
*/
|
||||
//---
|
||||
hipError_t hipDeviceEnablePeerAccess ( int peerDevice, unsigned int flags )
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
// TODO-p2p
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes )
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
// HCC has a unified memory architecture so device specifiers are not required.
|
||||
return hipMemcpy(dst, src, sizeBytes, hipMemcpyDefault);
|
||||
};
|
||||
|
||||
|
||||
/**
|
||||
* @bug This function uses a synchronous copy
|
||||
*/
|
||||
//---
|
||||
hipError_t hipMemcpyPeerAsync ( void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream )
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
// HCC has a unified memory architecture so device specifiers are not required.
|
||||
return hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream);
|
||||
};
|
||||
@@ -0,0 +1,118 @@
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
// Stream
|
||||
//
|
||||
|
||||
//---
|
||||
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
ihipDevice_t *device = ihipGetTlsDefaultDevice();
|
||||
hc::accelerator acc = device->_acc;
|
||||
|
||||
// TODO - se try-catch loop to detect memory exception?
|
||||
//
|
||||
//
|
||||
//Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete:
|
||||
//This matches CUDA stream behavior:
|
||||
|
||||
auto istream = new ihipStream_t(device->_device_index, acc.create_view(), device->_stream_id, flags);
|
||||
device->_streams.push_back(istream);
|
||||
*stream = istream;
|
||||
tprintf(DB_SYNC, "hipStreamCreate, stream=%p\n", *stream);
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
/**
|
||||
* @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)
|
||||
{
|
||||
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
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->wait();
|
||||
e = hipSuccess;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
};
|
||||
|
||||
|
||||
//---
|
||||
hipError_t hipStreamSynchronize(hipStream_t stream)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
if (stream == NULL) {
|
||||
ihipDevice_t *device = ihipGetTlsDefaultDevice();
|
||||
device->syncDefaultStream(true/*waitOnSelf*/);
|
||||
} else {
|
||||
stream->wait();
|
||||
e = hipSuccess;
|
||||
}
|
||||
|
||||
|
||||
return ihipLogStatus(e);
|
||||
};
|
||||
|
||||
|
||||
//---
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidResourceHandle
|
||||
*/
|
||||
hipError_t hipStreamDestroy(hipStream_t stream)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
//--- Drain the stream:
|
||||
if (stream == NULL) {
|
||||
ihipDevice_t *device = ihipGetTlsDefaultDevice();
|
||||
device->syncDefaultStream(true/*waitOnSelf*/);
|
||||
} else {
|
||||
stream->wait();
|
||||
e = hipSuccess;
|
||||
}
|
||||
|
||||
ihipDevice_t *device = stream->getDevice();
|
||||
|
||||
if (device) {
|
||||
device->_streams.remove(stream);
|
||||
delete stream;
|
||||
} else {
|
||||
e = hipErrorInvalidResourceHandle;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
if (flags == NULL) {
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
} else if (stream == NULL) {
|
||||
return ihipLogStatus(hipErrorInvalidResourceHandle);
|
||||
} else {
|
||||
*flags = stream->_flags;
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
}
|
||||
새 이슈에서 참조
사용자 차단