From c1bb44fa73d41efa4b6899d969319f947934602d Mon Sep 17 00:00:00 2001 From: foreman Date: Tue, 14 Aug 2018 18:54:13 -0400 Subject: [PATCH] P4 to Git Change 1593706 by skudchad@skudchad_test2_win_opencl on 2018/08/14 18:44:29 SWDEV-145570 - [HIP] Implement hipError* ReviewBoardURL = http://ocltc.amd.com/reviews/r/15619/diff/ Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#13 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_device.cpp#14 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_device_runtime.cpp#11 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_error.cpp#2 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.cpp#5 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#13 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#41 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#17 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_peer.cpp#2 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#17 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_profile.cpp#2 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#12 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_surface.cpp#2 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#9 edit [ROCm/clr commit: 2ec2cfd3c1b10873e5be9423028b4a27e41c1b05] --- projects/clr/hipamd/api/hip/hip_context.cpp | 41 +++-- projects/clr/hipamd/api/hip/hip_device.cpp | 36 ++-- .../clr/hipamd/api/hip/hip_device_runtime.cpp | 76 ++++---- projects/clr/hipamd/api/hip/hip_error.cpp | 132 +++++++++++++- projects/clr/hipamd/api/hip/hip_event.cpp | 34 ++-- projects/clr/hipamd/api/hip/hip_internal.hpp | 7 +- projects/clr/hipamd/api/hip/hip_memory.cpp | 172 +++++++++--------- projects/clr/hipamd/api/hip/hip_module.cpp | 32 ++-- projects/clr/hipamd/api/hip/hip_peer.cpp | 20 +- projects/clr/hipamd/api/hip/hip_platform.cpp | 17 +- projects/clr/hipamd/api/hip/hip_profile.cpp | 4 +- projects/clr/hipamd/api/hip/hip_stream.cpp | 28 +-- projects/clr/hipamd/api/hip/hip_surface.cpp | 4 +- projects/clr/hipamd/api/hip/hip_texture.cpp | 58 +++--- 14 files changed, 396 insertions(+), 265 deletions(-) diff --git a/projects/clr/hipamd/api/hip/hip_context.cpp b/projects/clr/hipamd/api/hip/hip_context.cpp index e75c8e5b5e..0ac257df83 100644 --- a/projects/clr/hipamd/api/hip/hip_context.cpp +++ b/projects/clr/hipamd/api/hip/hip_context.cpp @@ -31,6 +31,7 @@ namespace hip { thread_local amd::Context* g_context = nullptr; thread_local std::stack g_ctxtStack; +thread_local hipError_t g_lastError = hipSuccess; std::once_flag g_ihipInitialized; std::map g_nullStreams; @@ -85,14 +86,14 @@ using namespace hip; hipError_t hipInit(unsigned int flags) { HIP_INIT_API(flags); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) { HIP_INIT_API(ctx, flags, device); if (static_cast(device) >= g_devices.size()) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *ctx = reinterpret_cast(g_devices[device]); @@ -101,7 +102,7 @@ hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) g_devices[device]->retain(); g_ctxtStack.push(g_devices[device]); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipCtxSetCurrent(hipCtx_t ctx) { @@ -119,7 +120,7 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) { g_ctxtStack.push(hip::getCurrentContext()); } - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipCtxGetCurrent(hipCtx_t* ctx) { @@ -127,19 +128,19 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx) { *ctx = reinterpret_cast(hip::getCurrentContext()); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipRuntimeGetVersion(int *runtimeVersion) { HIP_INIT_API(runtimeVersion); if (!runtimeVersion) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *runtimeVersion = AMD_PLATFORM_BUILD_NUMBER; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipCtxDestroy(hipCtx_t ctx) { @@ -147,7 +148,7 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) { amd::Context* amdContext = reinterpret_cast(as_amd(ctx)); if (amdContext == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } // Need to remove the ctx of calling thread if its the top one @@ -163,7 +164,7 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) { } } - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { @@ -171,17 +172,17 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { amd::Context* amdContext = reinterpret_cast(as_amd(ctx)); if (amdContext == nullptr) { - return hipErrorInvalidContext; + HIP_RETURN(hipErrorInvalidContext); } if (!g_ctxtStack.empty()) { amdContext = g_ctxtStack.top(); g_ctxtStack.pop(); } else { - return hipErrorInvalidContext; + HIP_RETURN(hipErrorInvalidContext); } - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipCtxPushCurrent(hipCtx_t ctx) { @@ -189,13 +190,13 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) { amd::Context* amdContext = reinterpret_cast(as_amd(ctx)); if (amdContext == nullptr) { - return hipErrorInvalidContext; + HIP_RETURN(hipErrorInvalidContext); } hip::g_context = amdContext; g_ctxtStack.push(hip::getCurrentContext()); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDriverGetVersion(int* driverVersion) { @@ -205,13 +206,13 @@ hipError_t hipDriverGetVersion(int* driverVersion) { const auto& info = deviceHandle->info(); if (driverVersion) { - *driverVersion = AMD_PLATFORM_BUILD_NUMBER * 100 + - AMD_PLATFORM_REVISION_NUMBER; + *driverVersion = AMD_PLATFORM_BUILD_NUMBER * 100 + + AMD_PLATFORM_REVISION_NUMBER; } else { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } - return hipSuccess;; + HIP_RETURN(hipSuccess); } hipError_t hipCtxGetDevice(hipDevice_t* device) { @@ -221,11 +222,11 @@ hipError_t hipCtxGetDevice(hipDevice_t* device) { for (unsigned int i = 0; i < g_devices.size(); i++) { if (g_devices[i] == hip::getCurrentContext()) { *device = static_cast(i); - return hipSuccess; + HIP_RETURN(hipSuccess); } } } else { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } return hipErrorUnknown; diff --git a/projects/clr/hipamd/api/hip/hip_device.cpp b/projects/clr/hipamd/api/hip/hip_device.cpp index 5c03f13cc5..30ad2f3af1 100644 --- a/projects/clr/hipamd/api/hip/hip_device.cpp +++ b/projects/clr/hipamd/api/hip/hip_device.cpp @@ -30,10 +30,10 @@ hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) { if (device != nullptr) { *device = deviceId; } else { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } - return hipSuccess; + HIP_RETURN(hipSuccess); }; hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig) { @@ -42,7 +42,7 @@ hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig) // No way to set cache config yet. - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceTotalMem (size_t *bytes, hipDevice_t device) { @@ -50,11 +50,11 @@ hipError_t hipDeviceTotalMem (size_t *bytes, hipDevice_t device) { HIP_INIT_API(bytes, device); if (device < 0 || static_cast(device) >= g_devices.size()) { - return hipErrorInvalidDevice; + HIP_RETURN(hipErrorInvalidDevice); } if (bytes == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } auto* deviceHandle = g_devices[device]->devices()[0]; @@ -62,7 +62,7 @@ hipError_t hipDeviceTotalMem (size_t *bytes, hipDevice_t device) { *bytes = info.globalMemSize_; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device) { @@ -70,11 +70,11 @@ hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device HIP_INIT_API(major, minor, device); if (device < 0 || static_cast(device) >= g_devices.size()) { - return hipErrorInvalidDevice; + HIP_RETURN(hipErrorInvalidDevice); } if (major == nullptr || minor == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } auto* deviceHandle = g_devices[device]->devices()[0]; @@ -82,13 +82,13 @@ hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device *major = info.gfxipVersion_ / 100; *minor = info.gfxipVersion_ % 100; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceGetCount(int* count) { HIP_INIT_API(count); - return ihipDeviceGetCount(count); + HIP_RETURN(ihipDeviceGetCount(count)); } hipError_t ihipDeviceGetCount(int* count) { @@ -107,11 +107,11 @@ hipError_t hipDeviceGetName(char *name, int len, hipDevice_t device) { HIP_INIT_API((void*)name, len, device); if (device < 0 || static_cast(device) >= g_devices.size()) { - return hipErrorInvalidDevice; + HIP_RETURN(hipErrorInvalidDevice); } if (name == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } auto* deviceHandle = g_devices[device]->devices()[0]; @@ -120,18 +120,18 @@ hipError_t hipDeviceGetName(char *name, int len, hipDevice_t device) { len = ((cl_uint)len < ::strlen(info.boardName_)) ? len : 128; ::strncpy(name, info.boardName_, len); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device ) { HIP_INIT_API(props, device); if (props == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } if (unsigned(device) >= g_devices.size()) { - return hipErrorInvalidDevice; + HIP_RETURN(hipErrorInvalidDevice); } auto* deviceHandle = g_devices[device]->devices()[0]; @@ -188,7 +188,7 @@ hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device ) deviceProps.gcnArch = info.gfxipVersion_; *props = deviceProps; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) { @@ -196,7 +196,7 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av) { @@ -204,5 +204,5 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } diff --git a/projects/clr/hipamd/api/hip/hip_device_runtime.cpp b/projects/clr/hipamd/api/hip/hip_device_runtime.cpp index 5b60623761..87a149a10a 100644 --- a/projects/clr/hipamd/api/hip/hip_device_runtime.cpp +++ b/projects/clr/hipamd/api/hip/hip_device_runtime.cpp @@ -29,7 +29,7 @@ hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* properties) { HIP_INIT_API(device, properties); if (device == nullptr || properties == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *device = 0; @@ -135,7 +135,7 @@ hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* properties) { } } - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) { @@ -143,19 +143,19 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) HIP_INIT_API(pi, attr, device); if (pi == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } int count = 0; ihipDeviceGetCount(&count); if (device < 0 || device >= count) { - return hipErrorInvalidDevice; + HIP_RETURN(hipErrorInvalidDevice); } //FIXME: should we cache the props, or just select from deviceHandle->info_? hipDeviceProp_t prop = {0}; hipError_t err = hipGetDeviceProperties(&prop, device); - if (err != hipSuccess) return err; + if (err != hipSuccess) HIP_RETURN(err); switch (attr) { case hipDeviceAttributeMaxThreadsPerBlock: @@ -234,10 +234,10 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) *pi = prop.isMultiGpuBoard; break; default: - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceGetByPCIBusId(int* device, const char*pciBusIdstr) { @@ -245,7 +245,7 @@ hipError_t hipDeviceGetByPCIBusId(int* device, const char*pciBusIdstr) { HIP_INIT_API(device, pciBusIdstr); if (device == nullptr || pciBusIdstr == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } int pciBusID = -1; @@ -268,19 +268,19 @@ hipError_t hipDeviceGetByPCIBusId(int* device, const char*pciBusIdstr) { } } - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceGetCacheConfig ( hipFuncCache_t * cacheConfig ) { HIP_INIT_API(cacheConfig); if(cacheConfig == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *cacheConfig = hipFuncCache_t(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceGetLimit ( size_t* pValue, hipLimit_t limit ) { @@ -288,23 +288,23 @@ hipError_t hipDeviceGetLimit ( size_t* pValue, hipLimit_t limit ) { HIP_INIT_API(pValue, limit); if(pValue == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } if(limit == hipLimitMallocHeapSize) { hipDeviceProp_t prop; hipGetDeviceProperties(&prop, 0); *pValue = prop.totalGlobalMem; - return hipSuccess; + HIP_RETURN(hipSuccess); } else { - return hipErrorUnsupportedLimit; + HIP_RETURN(hipErrorUnsupportedLimit); } } /** hipError_t hipDeviceGetP2PAttribute ( int* value, hipDeviceP2PAttr attr, int srcDevice, int dstDevice ) { assert(0); - return hipSuccess; + HIP_RETURN(hipSuccess); } **/ @@ -315,11 +315,11 @@ hipError_t hipDeviceGetPCIBusId ( char* pciBusId, int len, int device ) { int count; ihipDeviceGetCount(&count); if (device < 0 || device > count) { - return hipErrorInvalidDevice; + HIP_RETURN(hipErrorInvalidDevice); } if (pciBusId == nullptr || len < 0) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } hipDeviceProp_t prop; @@ -330,7 +330,7 @@ hipError_t hipDeviceGetPCIBusId ( char* pciBusId, int len, int device ) { prop.pciBusID, prop.pciDeviceID); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ) { @@ -338,12 +338,12 @@ hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ) { *pConfig = hipSharedMemBankSizeFourByte; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceGetStreamPriorityRange ( int* leastPriority, int* greatestPriority ) { assert(0); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceReset ( void ) { @@ -351,7 +351,7 @@ hipError_t hipDeviceReset ( void ) { /* FIXME */ - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceSetCacheConfig ( hipFuncCache_t cacheConfig ) { @@ -359,11 +359,11 @@ hipError_t hipDeviceSetCacheConfig ( hipFuncCache_t cacheConfig ) { // No way to set cache config yet. - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceSetLimit ( hipLimit_t limit, size_t value ) { - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) { @@ -371,7 +371,7 @@ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) { // No way to set cache config yet. - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipDeviceSynchronize ( void ) { @@ -382,11 +382,11 @@ hipError_t hipDeviceSynchronize ( void ) { amd::HostQueue* queue = hip::getNullStream(); if (!queue) { - return hipErrorOutOfMemory; + HIP_RETURN(hipErrorOutOfMemory); } queue->finish(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipGetDevice ( int* deviceId ) { @@ -396,24 +396,24 @@ hipError_t hipGetDevice ( int* deviceId ) { for (unsigned int i = 0; i < g_devices.size(); i++) { if (g_devices[i] == hip::getCurrentContext()) { *deviceId = i; - return hipSuccess; + HIP_RETURN(hipSuccess); } } } else { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipGetDeviceCount ( int* count ) { HIP_INIT_API(count); - return ihipDeviceGetCount(count); + HIP_RETURN(ihipDeviceGetCount(count)); } hipError_t hipGetDeviceFlags ( unsigned int* flags ) { - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipIpcGetEventHandle ( hipIpcEventHandle_t* handle, hipEvent_t event ) { @@ -421,7 +421,7 @@ hipError_t hipIpcGetEventHandle ( hipIpcEventHandle_t* handle, hipEvent_t event assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipIpcOpenEventHandle ( hipEvent_t* event, hipIpcEventHandle_t handle ) { @@ -429,7 +429,7 @@ hipError_t hipIpcOpenEventHandle ( hipEvent_t* event, hipIpcEventHandle_t handle assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipSetDevice ( int device ) { @@ -438,9 +438,9 @@ hipError_t hipSetDevice ( int device ) { if (static_cast(device) < g_devices.size()) { hip::setCurrentContext(device); - return hipSuccess; + HIP_RETURN(hipSuccess); } - return hipErrorInvalidDevice; + HIP_RETURN(hipErrorInvalidDevice); } hipError_t hipSetDeviceFlags ( unsigned int flags ) { @@ -453,10 +453,10 @@ hipError_t hipSetDeviceFlags ( unsigned int flags ) { hipDeviceScheduleMask | hipDeviceMapHost | hipDeviceLmemResizeToMax; if (flags & (~supportedFlags)) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipSetValidDevices ( int* device_arr, int len ) { @@ -464,6 +464,6 @@ hipError_t hipSetValidDevices ( int* device_arr, int len ) { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } diff --git a/projects/clr/hipamd/api/hip/hip_error.cpp b/projects/clr/hipamd/api/hip/hip_error.cpp index 2a8785c375..5f76e560c0 100644 --- a/projects/clr/hipamd/api/hip/hip_error.cpp +++ b/projects/clr/hipamd/api/hip/hip_error.cpp @@ -27,22 +27,146 @@ THE SOFTWARE. hipError_t hipGetLastError() { HIP_INIT_API(); - return hipErrorUnknown; + hipError_t err = hip::g_lastError; + hip::g_lastError = hipSuccess; + return err; } hipError_t hipPeekAtLastError() { HIP_INIT_API(); - return hipErrorUnknown; + hipError_t err = hip::g_lastError; + HIP_RETURN(err); } const char *hipGetErrorName(hipError_t hip_error) { - return ""; + switch (hip_error) { + case hipSuccess: + return "hipSuccess"; + case hipErrorOutOfMemory: + return "hipErrorOutOfMemory"; + case hipErrorNotInitialized: + return "hipErrorNotInitialized"; + case hipErrorDeinitialized: + return "hipErrorDeinitialized"; + case hipErrorProfilerDisabled: + return "hipErrorProfilerDisabled"; + case hipErrorProfilerNotInitialized: + return "hipErrorProfilerNotInitialized"; + case hipErrorProfilerAlreadyStarted: + return "hipErrorProfilerAlreadyStarted"; + case hipErrorProfilerAlreadyStopped: + return "hipErrorProfilerAlreadyStopped"; + case hipErrorInvalidImage: + return "hipErrorInvalidImage"; + case hipErrorInvalidContext: + return "hipErrorInvalidContext"; + case hipErrorContextAlreadyCurrent: + return "hipErrorContextAlreadyCurrent"; + case hipErrorMapFailed: + return "hipErrorMapFailed"; + case hipErrorUnmapFailed: + return "hipErrorUnmapFailed"; + case hipErrorArrayIsMapped: + return "hipErrorArrayIsMapped"; + case hipErrorAlreadyMapped: + return "hipErrorAlreadyMapped"; + case hipErrorNoBinaryForGpu: + return "hipErrorNoBinaryForGpu"; + case hipErrorAlreadyAcquired: + return "hipErrorAlreadyAcquired"; + case hipErrorNotMapped: + return "hipErrorNotMapped"; + case hipErrorNotMappedAsArray: + return "hipErrorNotMappedAsArray"; + case hipErrorNotMappedAsPointer: + return "hipErrorNotMappedAsPointer"; + case hipErrorECCNotCorrectable: + return "hipErrorECCNotCorrectable"; + case hipErrorUnsupportedLimit: + return "hipErrorUnsupportedLimit"; + case hipErrorContextAlreadyInUse: + return "hipErrorContextAlreadyInUse"; + case hipErrorPeerAccessUnsupported: + return "hipErrorPeerAccessUnsupported"; + case hipErrorInvalidKernelFile: + return "hipErrorInvalidKernelFile"; + case hipErrorInvalidGraphicsContext: + return "hipErrorInvalidGraphicsContext"; + case hipErrorInvalidSource: + return "hipErrorInvalidSource"; + case hipErrorFileNotFound: + return "hipErrorFileNotFound"; + case hipErrorSharedObjectSymbolNotFound: + return "hipErrorSharedObjectSymbolNotFound"; + case hipErrorSharedObjectInitFailed: + return "hipErrorSharedObjectInitFailed"; + case hipErrorOperatingSystem: + return "hipErrorOperatingSystem"; + case hipErrorSetOnActiveProcess: + return "hipErrorSetOnActiveProcess"; + case hipErrorInvalidHandle: + return "hipErrorInvalidHandle"; + case hipErrorNotFound: + return "hipErrorNotFound"; + case hipErrorIllegalAddress: + return "hipErrorIllegalAddress"; + case hipErrorMissingConfiguration: + return "hipErrorMissingConfiguration"; + case hipErrorMemoryAllocation: + return "hipErrorMemoryAllocation"; + case hipErrorInitializationError: + return "hipErrorInitializationError"; + case hipErrorLaunchFailure: + return "hipErrorLaunchFailure"; + case hipErrorPriorLaunchFailure: + return "hipErrorPriorLaunchFailure"; + case hipErrorLaunchTimeOut: + return "hipErrorLaunchTimeOut"; + case hipErrorLaunchOutOfResources: + return "hipErrorLaunchOutOfResources"; + case hipErrorInvalidDeviceFunction: + return "hipErrorInvalidDeviceFunction"; + case hipErrorInvalidConfiguration: + return "hipErrorInvalidConfiguration"; + case hipErrorInvalidDevice: + return "hipErrorInvalidDevice"; + case hipErrorInvalidValue: + return "hipErrorInvalidValue"; + case hipErrorInvalidDevicePointer: + return "hipErrorInvalidDevicePointer"; + case hipErrorInvalidMemcpyDirection: + return "hipErrorInvalidMemcpyDirection"; + case hipErrorUnknown: + return "hipErrorUnknown"; + case hipErrorInvalidResourceHandle: + return "hipErrorInvalidResourceHandle"; + case hipErrorNotReady: + return "hipErrorNotReady"; + case hipErrorNoDevice: + return "hipErrorNoDevice"; + case hipErrorPeerAccessAlreadyEnabled: + return "hipErrorPeerAccessAlreadyEnabled"; + case hipErrorPeerAccessNotEnabled: + return "hipErrorPeerAccessNotEnabled"; + case hipErrorRuntimeMemory: + return "hipErrorRuntimeMemory"; + case hipErrorRuntimeOther: + return "hipErrorRuntimeOther"; + case hipErrorHostMemoryAlreadyRegistered: + return "hipErrorHostMemoryAlreadyRegistered"; + case hipErrorHostMemoryNotRegistered: + return "hipErrorHostMemoryNotRegistered"; + case hipErrorTbd: + return "hipErrorTbd"; + default: + return "hipErrorUnknown"; + }; } const char *hipGetErrorString(hipError_t hip_error) { - return ""; + return hipGetErrorName(hip_error); } diff --git a/projects/clr/hipamd/api/hip/hip_event.cpp b/projects/clr/hipamd/api/hip/hip_event.cpp index 127148e63c..197eac0579 100644 --- a/projects/clr/hipamd/api/hip/hip_event.cpp +++ b/projects/clr/hipamd/api/hip/hip_event.cpp @@ -64,38 +64,38 @@ hipError_t ihipEventQuery(hipEvent_t event) { e->event_->notifyCmdQueue(); - return (e->event_->status() == CL_COMPLETE)? hipSuccess : hipErrorNotReady; + return (e->event_->status() == CL_COMPLETE) ? hipSuccess : hipErrorNotReady; } hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) { HIP_INIT_API(event, flags); - return ihipEventCreateWithFlags(event, flags); + HIP_RETURN(ihipEventCreateWithFlags(event, flags)); } hipError_t hipEventCreate(hipEvent_t* event) { HIP_INIT_API(event); - return ihipEventCreateWithFlags(event, 0); + HIP_RETURN(ihipEventCreateWithFlags(event, 0)); } hipError_t hipEventDestroy(hipEvent_t event) { HIP_INIT_API(event); if (event == nullptr) { - return hipErrorInvalidResourceHandle; + HIP_RETURN(hipErrorInvalidResourceHandle); } delete reinterpret_cast(event); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) { HIP_INIT_API(ms, start, stop); if (start == nullptr || stop == nullptr) { - return hipErrorInvalidResourceHandle; + HIP_RETURN(hipErrorInvalidResourceHandle); } hip::Event* eStart = reinterpret_cast(start); @@ -103,32 +103,32 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) { if (eStart->event_ == nullptr || eStop->event_ == nullptr) { - return hipErrorInvalidResourceHandle; + HIP_RETURN(hipErrorInvalidResourceHandle); } if ((eStart->flags | eStop->flags) & hipEventDisableTiming) { - return hipErrorInvalidResourceHandle; + HIP_RETURN(hipErrorInvalidResourceHandle); } if (ihipEventQuery(start) == hipErrorNotReady || ihipEventQuery(stop) == hipErrorNotReady) { - return hipErrorNotReady; + HIP_RETURN(hipErrorNotReady); } if (ms == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *ms = static_cast(eStop->event_->profilingInfo().submitted_ - eStart->event_->profilingInfo().submitted_)/1000000.f; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { HIP_INIT_API(event, stream); if (event == nullptr) { - return hipErrorInvalidResourceHandle; + HIP_RETURN(hipErrorInvalidResourceHandle); } hip::Event* e = reinterpret_cast(event); @@ -148,29 +148,29 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { e->event_ = &command->event(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipEventSynchronize(hipEvent_t event) { HIP_INIT_API(event); if (event == nullptr) { - return hipErrorInvalidResourceHandle; + HIP_RETURN(hipErrorInvalidResourceHandle); } hip::Event* e = reinterpret_cast(event); if (e->event_ == nullptr) { - return hipErrorInvalidResourceHandle; + HIP_RETURN(hipErrorInvalidResourceHandle); } e->event_->awaitCompletion(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipEventQuery(hipEvent_t event) { HIP_INIT_API(event); - return ihipEventQuery(event); + HIP_RETURN(ihipEventQuery(event)); } diff --git a/projects/clr/hipamd/api/hip/hip_internal.hpp b/projects/clr/hipamd/api/hip/hip_internal.hpp index e82cb13a2a..a8a9c3eeb3 100644 --- a/projects/clr/hipamd/api/hip/hip_internal.hpp +++ b/projects/clr/hipamd/api/hip/hip_internal.hpp @@ -39,7 +39,7 @@ THE SOFTWARE. #define HIP_INIT_API(...) \ amd::Thread* thread = amd::Thread::current(); \ if (!CL_CHECK_THREAD(thread)) { \ - return hipErrorOutOfMemory; \ + HIP_RETURN(hipErrorOutOfMemory); \ } \ HIP_INIT(); @@ -51,6 +51,7 @@ class accelerator_view; namespace hip { extern std::once_flag g_ihipInitialized; extern thread_local amd::Context* g_context; + extern thread_local hipError_t g_lastError; extern void init(); @@ -64,5 +65,9 @@ extern std::vector g_devices; extern hipError_t ihipDeviceGetCount(int* count); extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset); +#define HIP_RETURN(ret) \ + hip::g_lastError = ret; \ + return ret; \ + #endif // HIP_SRC_HIP_INTERNAL_H diff --git a/projects/clr/hipamd/api/hip/hip_memory.cpp b/projects/clr/hipamd/api/hip/hip_memory.cpp index 24ed647acb..b800ffe713 100644 --- a/projects/clr/hipamd/api/hip/hip_memory.cpp +++ b/projects/clr/hipamd/api/hip/hip_memory.cpp @@ -151,14 +151,14 @@ hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, amd::HostQueue& qu hipError_t hipMalloc(void** ptr, size_t sizeBytes) { HIP_INIT_API(ptr, sizeBytes); - return ihipMalloc(ptr, sizeBytes, 0); + HIP_RETURN(ihipMalloc(ptr, sizeBytes, 0)); } hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_INIT_API(ptr, sizeBytes, flags); if (ptr == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *ptr = nullptr; @@ -166,10 +166,10 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { // can't have both Coherent and NonCoherent flags set at the same time if ((flags & coherentFlags) == coherentFlags) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } - return ihipMalloc(ptr, sizeBytes, CL_MEM_SVM_FINE_GRAIN_BUFFER | (flags << 16)); + HIP_RETURN(ihipMalloc(ptr, sizeBytes, CL_MEM_SVM_FINE_GRAIN_BUFFER | (flags << 16))); } hipError_t hipFree(void* ptr) { @@ -177,9 +177,9 @@ hipError_t hipFree(void* ptr) { hip::syncStreams(); hip::getNullStream()->finish(); amd::SvmBuffer::free(*hip::getCurrentContext(), ptr); - return hipSuccess; + HIP_RETURN(hipSuccess); } - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) { @@ -187,7 +187,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); - return ihipMemcpy(dst, src, sizeBytes, kind, *queue); + HIP_RETURN(ihipMemcpy(dst, src, sizeBytes, kind, *queue)); } hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { @@ -203,7 +203,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } - return ihipMemset(dst, value, sizeBytes, *queue, true); + HIP_RETURN(ihipMemset(dst, value, sizeBytes, *queue, true)); } hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { @@ -212,7 +212,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); - return ihipMemset(dst, value, sizeBytes, *queue); + HIP_RETURN(ihipMemset(dst, value, sizeBytes, *queue)); } hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) { @@ -222,12 +222,12 @@ hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) { amd::Memory* svmMem = getMemoryObject(ptr, offset); if (svmMem == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *size = svmMem->getSize(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipHostFree(void* ptr) { @@ -235,9 +235,9 @@ hipError_t hipHostFree(void* ptr) { if (amd::SvmBuffer::malloced(ptr)) { amd::SvmBuffer::free(*hip::getCurrentContext(), ptr); - return hipSuccess; + HIP_RETURN(hipSuccess); } - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } hipError_t hipFreeArray(hipArray* array) { @@ -245,9 +245,9 @@ hipError_t hipFreeArray(hipArray* array) { if (amd::SvmBuffer::malloced(array->data)) { amd::SvmBuffer::free(*hip::getCurrentContext(), array->data); - return hipSuccess; + HIP_RETURN(hipSuccess); } - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr) { @@ -259,13 +259,13 @@ hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDevice amd::Memory* svmMem = getMemoryObject(ptr, offset); if (svmMem == nullptr) { - return hipErrorInvalidDevicePointer; + HIP_RETURN(hipErrorInvalidDevicePointer); } *pbase = svmMem->getSvmPtr(); *psize = svmMem->getSize(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipMemGetInfo(size_t* free, size_t* total) { @@ -274,17 +274,17 @@ hipError_t hipMemGetInfo(size_t* free, size_t* total) { size_t freeMemory[2]; amd::Device* device = hip::getCurrentContext()->devices()[0]; if(device == nullptr) { - return hipErrorInvalidDevice; + HIP_RETURN(hipErrorInvalidDevice); } if(!device->globalFreeMemory(freeMemory)) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *free = freeMemory[0] * Ki; *total = device->info().globalMemSize_; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height, size_t depth, @@ -325,7 +325,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height HIP_INIT_API(ptr, pitch, width, height); const cl_image_format image_format = { CL_R, CL_UNSIGNED_INT8 }; - return ihipMallocPitch(ptr, pitch, width, height, 1, CL_MEM_OBJECT_IMAGE2D, &image_format); + HIP_RETURN(ihipMallocPitch(ptr, pitch, width, height, 1, CL_MEM_OBJECT_IMAGE2D, &image_format)); } hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { @@ -334,7 +334,7 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { size_t pitch = 0; if (pitchedDevPtr == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } const cl_image_format image_format = { CL_R, CL_UNSIGNED_INT8 }; @@ -348,7 +348,7 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { pitchedDevPtr->ysize = extent.height; } - return status; + HIP_RETURN(status); } hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent) { @@ -360,14 +360,14 @@ hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent) hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); - return ihipMemset(dst, value, sizeBytes, *queue); + HIP_RETURN(ihipMemset(dst, value, sizeBytes, *queue)); } hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { HIP_INIT_API(array, pAllocateArray); if (array[0]->width == 0) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *array = (hipArray*)malloc(sizeof(hipArray)); @@ -393,7 +393,7 @@ hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocat hipError_t status = ihipMallocPitch(ptr, &pitch, array[0]->width, array[0]->height, 1, CL_MEM_OBJECT_IMAGE2D, &image_format); - return status; + HIP_RETURN(status); } hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, @@ -401,7 +401,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, HIP_INIT_API(array, desc, width, height, flags); if (width == 0) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *array = (hipArray*)malloc(sizeof(hipArray)); @@ -436,7 +436,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, hipError_t status = ihipMallocPitch(ptr, &pitch, width, height, 1, CL_MEM_OBJECT_IMAGE2D, &image_format); - return status; + HIP_RETURN(status); } hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc* desc, @@ -475,7 +475,7 @@ hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc hipError_t status = ihipMallocPitch(ptr, &pitch, extent.width, extent.height, extent.depth, CL_MEM_OBJECT_IMAGE3D, &image_format); - return status; + HIP_RETURN(status); } hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) { @@ -483,19 +483,19 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) { if (flagsPtr == nullptr || hostPtr == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } size_t offset = 0; amd::Memory* svmMem = getMemoryObject(hostPtr, offset); if (svmMem == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *flagsPtr = svmMem->getMemFlags() >> 16; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) { @@ -506,12 +506,12 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) if (!mem->create(hostPtr)) { mem->release(); - return hipErrorMemoryAllocation; + HIP_RETURN(hipErrorMemoryAllocation); } amd::MemObjMap::AddMemObj(hostPtr, mem); - return hipSuccess; + HIP_RETURN(hipSuccess); } else { - return ihipMalloc(&hostPtr, sizeBytes, flags); + HIP_RETURN(ihipMalloc(&hostPtr, sizeBytes, flags)); } } @@ -522,7 +522,7 @@ hipError_t hipHostUnregister(void* hostPtr) { hip::syncStreams(); hip::getNullStream()->finish(); amd::SvmBuffer::free(*hip::getCurrentContext(), hostPtr); - return hipSuccess; + HIP_RETURN(hipSuccess); } else { size_t offset = 0; amd::Memory* mem = getMemoryObject(hostPtr, offset); @@ -532,16 +532,16 @@ hipError_t hipHostUnregister(void* hostPtr) { hip::getNullStream()->finish(); amd::MemObjMap::RemoveMemObj(hostPtr); mem->release(); - return hipSuccess; + HIP_RETURN(hipSuccess); } } - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } // Deprecated function: hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) { - return ihipMalloc(ptr, sizeBytes, flags); + HIP_RETURN(ihipMalloc(ptr, sizeBytes, flags)); }; @@ -551,7 +551,7 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t cou assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, @@ -560,7 +560,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_t count, @@ -569,7 +569,7 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_ assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t count, @@ -578,7 +578,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) { @@ -587,7 +587,7 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) { hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); - return ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToDevice, *queue); + HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToDevice, *queue)); } hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { @@ -596,7 +596,7 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); - return ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToHost, *queue); + HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToHost, *queue)); } hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) { @@ -605,7 +605,7 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); - return ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToDevice, *queue); + HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToDevice, *queue)); } hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { @@ -614,7 +614,7 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); - return ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToHost, *queue); + HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToHost, *queue)); } hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, @@ -631,7 +631,7 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } - return ihipMemcpy(dst, src, sizeBytes, kind, *queue, true); + HIP_RETURN(ihipMemcpy(dst, src, sizeBytes, kind, *queue, true)); } @@ -649,8 +649,8 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } - return ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToDevice, - *queue, true); + HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToDevice, + *queue, true)); } hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, @@ -667,8 +667,8 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } - return ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToDevice, - *queue, true); + HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToDevice, + *queue, true)); } hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, @@ -685,8 +685,8 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } - return ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToHost, - *queue, true); + HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToHost, + *queue, true)); } hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { @@ -694,7 +694,7 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, @@ -763,7 +763,7 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); - return ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, *queue); + HIP_RETURN(ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, *queue)); } @@ -781,7 +781,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } - return ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, *queue, true); + HIP_RETURN(ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, *queue, true)); } hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, @@ -789,7 +789,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con HIP_INIT_API(dst, wOffset, hOffset, src, spitch, width, height, kind); if (dst->data == nullptr) { - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hip::syncStreams(); @@ -816,7 +816,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con } if ((wOffset + width > (dpitch)) || width > spitch) { - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } // Create buffer rectangle info structure @@ -836,7 +836,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con if (!srcRect.create(sOrigin, region, spitch, src_slice_pitch) || !dstRect.create(dOrigin, region, dpitch, dst_slice_pitch)) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } amd::Command* command = nullptr; @@ -855,7 +855,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con void* pSrc = reinterpret_cast(reinterpret_cast(src) + y * spitch); memcpy(pDst, pSrc, width); } - return hipSuccess; + HIP_RETURN(hipSuccess); } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) { command = new amd::WriteMemoryCommand(*queue, CL_COMMAND_WRITE_BUFFER_RECT, waitList, *dstMemory->asBuffer(), dstStart, size, src, dstRect, srcRect); @@ -868,14 +868,14 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con } if (command == nullptr) { - return hipErrorOutOfMemory; + HIP_RETURN(hipErrorOutOfMemory); } command->enqueue(); command->awaitCompletion(); command->release(); - return hipSuccess; + HIP_RETURN(hipSuccess); } @@ -909,14 +909,14 @@ hipError_t hipMemcpyToArray(hipArray* dstArray, size_t wOffset, size_t hOffset, } if (command == nullptr) { - return hipErrorOutOfMemory; + HIP_RETURN(hipErrorOutOfMemory); } command->enqueue(); command->awaitCompletion(); command->release(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset, @@ -949,14 +949,14 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs } if (command == nullptr) { - return hipErrorOutOfMemory; + HIP_RETURN(hipErrorOutOfMemory); } command->enqueue(); command->awaitCompletion(); command->release(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) { @@ -986,14 +986,14 @@ hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHo } if (command == nullptr) { - return hipErrorOutOfMemory; + HIP_RETURN(hipErrorOutOfMemory); } command->enqueue(); command->awaitCompletion(); command->release(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count) { @@ -1023,14 +1023,14 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t } if (command == nullptr) { - return hipErrorOutOfMemory; + HIP_RETURN(hipErrorOutOfMemory); } command->enqueue(); command->awaitCompletion(); command->release(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) { @@ -1108,7 +1108,7 @@ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) { if (!srcRect.create(srcOrigin, region, srcPitchInBytes, src_slice_pitch) || !dstRect.create(dstOrigin, region, dstPitchInbytes, dst_slice_pitch)) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } hipMemcpyKind kind = p->kind; @@ -1123,7 +1123,7 @@ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) { if (((srcMemory == nullptr) && (dstMemory == nullptr)) || (kind == hipMemcpyHostToHost)) { memcpy(dstPtr, srcPtr, region[0] * region[1] * region[2]); - return hipSuccess; + HIP_RETURN(hipSuccess); } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) { command = new amd::WriteMemoryCommand(*queue, CL_COMMAND_WRITE_BUFFER_RECT, waitList, *dstMemory->asBuffer(), srcStart, size, srcPtr, srcRect, dstRect); @@ -1137,14 +1137,14 @@ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) { } if (command == nullptr) { - return hipErrorOutOfMemory; + HIP_RETURN(hipErrorOutOfMemory); } command->enqueue(); command->awaitCompletion(); command->release(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t ihipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height, @@ -1192,7 +1192,7 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); - return ihipMemset2D(dst, pitch, value, width, height, *queue); + HIP_RETURN(ihipMemset2D(dst, pitch, value, width, height, *queue)); } hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, @@ -1208,14 +1208,14 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } - return ihipMemset2D(dst, pitch, value, width, height, *queue, true); + HIP_RETURN(ihipMemset2D(dst, pitch, value, width, height, *queue, true)); } hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) { HIP_INIT_API(dst, value, sizeBytes); if (dst == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } hip::syncStreams(); @@ -1232,7 +1232,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes &value, sizeof(char), fillOffset, fillSize); if (command == nullptr) { - return hipErrorOutOfMemory; + HIP_RETURN(hipErrorOutOfMemory); } command->enqueue(); @@ -1243,7 +1243,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes memset(dst, value, sizeBytes); } - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) { @@ -1251,7 +1251,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags) { @@ -1259,7 +1259,7 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipIpcCloseMemHandle(void* devPtr) { @@ -1267,7 +1267,7 @@ hipError_t hipIpcCloseMemHandle(void* devPtr) { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { @@ -1287,11 +1287,11 @@ hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsi amd::Memory* memObj = getMemoryObject(hostPointer, offset); if (!memObj) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } *devicePointer = reinterpret_cast(memObj->getDeviceMemory(*hip::getCurrentContext()->devices()[0])->virtualAddress() + offset); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr) { @@ -1325,5 +1325,5 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void attributes->allocationFlags = 0; } - return hipSuccess; + HIP_RETURN(hipSuccess); } diff --git a/projects/clr/hipamd/api/hip/hip_module.cpp b/projects/clr/hipamd/api/hip/hip_module.cpp index 67c4f06c6c..9830a97dfb 100644 --- a/projects/clr/hipamd/api/hip/hip_module.cpp +++ b/projects/clr/hipamd/api/hip/hip_module.cpp @@ -56,18 +56,18 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname) HIP_INIT_API(module, fname); if (!fname) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } std::ifstream file{fname}; if (!file.is_open()) { - return hipErrorFileNotFound; + HIP_RETURN(hipErrorFileNotFound); } std::vector tmp{std::istreambuf_iterator{file}, std::istreambuf_iterator{}}; - return ihipModuleLoadData(module, tmp.data()); + HIP_RETURN(ihipModuleLoadData(module, tmp.data())); } @@ -76,21 +76,21 @@ hipError_t hipModuleUnload(hipModule_t hmod) HIP_INIT_API(hmod); if (hmod == nullptr) { - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } amd::Program* program = as_amd(reinterpret_cast(hmod)); program->release(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipModuleLoadData(hipModule_t *module, const void *image) { HIP_INIT_API(module, image); - return ihipModuleLoadData(module, image); + HIP_RETURN(ihipModuleLoadData(module, image)); } hipError_t ihipModuleLoadData(hipModule_t *module, const void *image) @@ -118,24 +118,24 @@ hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, const ch const amd::Symbol* symbol = program->findSymbol(name); if (!symbol) { - return hipErrorNotFound; + HIP_RETURN(hipErrorNotFound); } amd::Kernel* kernel = new amd::Kernel(*program, *symbol, name); if (!kernel) { - return hipErrorOutOfMemory; + HIP_RETURN(hipErrorOutOfMemory); } *hfunc = reinterpret_cast(as_cl(kernel)); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) { HIP_INIT_API(attr, func); - return hipErrorInvalidDeviceFunction; + HIP_RETURN(hipErrorInvalidDeviceFunction); } @@ -242,9 +242,9 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, uint32_t sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra) { - return ihipModuleLaunchKernel(f, gridDimX * blockDimX, gridDimY * blockDimY, gridDimZ * blockDimZ, + HIP_RETURN(ihipModuleLaunchKernel(f, gridDimX * blockDimX, gridDimY * blockDimY, gridDimZ * blockDimZ, blockDimX, blockDimY, blockDimZ, - sharedMemBytes, hStream, kernelParams, extra, nullptr, nullptr); + sharedMemBytes, hStream, kernelParams, extra, nullptr, nullptr)); } hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, @@ -255,8 +255,8 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, hipEvent_t startEvent, hipEvent_t stopEvent) { - return ihipModuleLaunchKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, - sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent); + HIP_RETURN(ihipModuleLaunchKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, + sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent)); } hipError_t hipModuleLaunchKernelExt(hipFunction_t f, uint32_t gridDimX, @@ -267,8 +267,8 @@ hipError_t hipModuleLaunchKernelExt(hipFunction_t f, uint32_t gridDimX, hipEvent_t startEvent, hipEvent_t stopEvent) { - return ihipModuleLaunchKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, - sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent); + HIP_RETURN(ihipModuleLaunchKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, + sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent)); } diff --git a/projects/clr/hipamd/api/hip/hip_peer.cpp b/projects/clr/hipamd/api/hip/hip_peer.cpp index ad552e94b4..73919a09ea 100644 --- a/projects/clr/hipamd/api/hip/hip_peer.cpp +++ b/projects/clr/hipamd/api/hip/hip_peer.cpp @@ -29,7 +29,7 @@ hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipMemcpyPeer(void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t srcCtx, @@ -38,7 +38,7 @@ hipError_t hipMemcpyPeer(void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t s assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hipCtx_t srcDevice, @@ -47,7 +47,7 @@ hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hi assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId) { @@ -55,7 +55,7 @@ hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDevi assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipDeviceDisablePeerAccess(int peerDeviceId) { @@ -63,7 +63,7 @@ hipError_t hipDeviceDisablePeerAccess(int peerDeviceId) { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) { @@ -71,7 +71,7 @@ hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevice, @@ -80,7 +80,7 @@ hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevic assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice, @@ -89,7 +89,7 @@ hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int src assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) { @@ -97,7 +97,7 @@ hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) { @@ -105,5 +105,5 @@ hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } \ No newline at end of file diff --git a/projects/clr/hipamd/api/hip/hip_platform.cpp b/projects/clr/hipamd/api/hip/hip_platform.cpp index a6bb4e290c..2adf74f7ce 100644 --- a/projects/clr/hipamd/api/hip/hip_platform.cpp +++ b/projects/clr/hipamd/api/hip/hip_platform.cpp @@ -252,7 +252,7 @@ extern "C" hipError_t hipConfigureCall( PlatformState::instance().configureCall(gridDim, blockDim, sharedMem, stream); - return hipSuccess; + HIP_RETURN(hipSuccess); } extern "C" hipError_t hipSetupArgument( @@ -264,7 +264,7 @@ extern "C" hipError_t hipSetupArgument( PlatformState::instance().setupArgument(arg, size, offset); - return hipSuccess; + HIP_RETURN(hipSuccess); } extern "C" hipError_t hipLaunchByPtr(const void *hostFunction) @@ -272,8 +272,9 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction) HIP_INIT_API(hostFunction); hipFunction_t func = PlatformState::instance().getFunc(hostFunction); - if (func == nullptr) - return hipErrorUnknown; + if (func == nullptr) { + HIP_RETURN(hipErrorUnknown); + } ihipExec_t exec; PlatformState::instance().popExec(exec); @@ -284,10 +285,10 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction) HIP_LAUNCH_PARAM_END }; - return hipModuleLaunchKernel(func, + HIP_RETURN(hipModuleLaunchKernel(func, exec.gridDim_.x, exec.gridDim_.y, exec.gridDim_.z, exec.blockDim_.x, exec.blockDim_.y, exec.blockDim_.z, - exec.sharedMem_, exec.hStream_, nullptr, extra); + exec.sharedMem_, exec.hStream_, nullptr, extra)); } #if defined(ATI_OS_LINUX) @@ -514,7 +515,7 @@ static inline std::uint32_t f32_as_u32(float f) { union { float f; std::uint32_t static inline float u32_as_f32(std::uint32_t u) { union { float f; std::uint32_t u; } v; v.u = u; return v.f; } static inline int clamp_int(int i, int l, int h) { return std::min(std::max(i, l), h); } -// half � float, the f16 is in the low 16 bits of the input argument �a� +// half float, the f16 is in the low 16 bits of the input argument static inline float __convert_half_to_float(std::uint32_t a) noexcept { std::uint32_t u = ((a << 13) + 0x70000000U) & 0x8fffe000U; std::uint32_t v = f32_as_u32(u32_as_f32(u) * 0x1.0p+112f) + 0x38000000U; @@ -522,7 +523,7 @@ static inline float __convert_half_to_float(std::uint32_t a) noexcept { return u32_as_f32(u) * 0x1.0p-112f; } -// float � half with nearest even rounding +// float half with nearest even rounding // The lower 16 bits of the result is the bit pattern for the f16 static inline std::uint32_t __convert_float_to_half(float a) noexcept { std::uint32_t u = f32_as_u32(a); diff --git a/projects/clr/hipamd/api/hip/hip_profile.cpp b/projects/clr/hipamd/api/hip/hip_profile.cpp index d53d7ffd46..8b3d42cbbc 100644 --- a/projects/clr/hipamd/api/hip/hip_profile.cpp +++ b/projects/clr/hipamd/api/hip/hip_profile.cpp @@ -29,7 +29,7 @@ hipError_t hipProfilerStart() { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } @@ -38,5 +38,5 @@ hipError_t hipProfilerStop() { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } \ No newline at end of file diff --git a/projects/clr/hipamd/api/hip/hip_stream.cpp b/projects/clr/hipamd/api/hip/hip_stream.cpp index c48eca87ae..a7221b0249 100644 --- a/projects/clr/hipamd/api/hip/hip_stream.cpp +++ b/projects/clr/hipamd/api/hip/hip_stream.cpp @@ -68,13 +68,13 @@ static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int fl hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) { HIP_INIT_API(stream, flags); - return ihipStreamCreateWithFlags(stream, flags); + HIP_RETURN(ihipStreamCreateWithFlags(stream, flags)); } hipError_t hipStreamCreate(hipStream_t *stream) { HIP_INIT_API(stream); - return ihipStreamCreateWithFlags(stream, hipStreamDefault); + HIP_RETURN(ihipStreamCreateWithFlags(stream, hipStreamDefault)); } hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) { @@ -86,10 +86,10 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) { if(flags != nullptr) { *flags = (it == streamSet.end()) ? hipStreamNonBlocking : hipStreamDefault; } else { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipStreamSynchronize(hipStream_t stream) { @@ -108,19 +108,19 @@ hipError_t hipStreamSynchronize(hipStream_t stream) { } if (hostQueue == nullptr) { - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hostQueue->finish(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipStreamDestroy(hipStream_t stream) { HIP_INIT_API(stream); if (stream == nullptr) { - return hipErrorInvalidResourceHandle; + HIP_RETURN(hipErrorInvalidResourceHandle); } amd::ScopedLock lock(streamSetLock); @@ -129,14 +129,14 @@ hipError_t hipStreamDestroy(hipStream_t stream) { hostQueue->release(); streamSet.erase(hostQueue); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { HIP_INIT_API(stream, event, flags); if (stream == nullptr || event == nullptr) { - return hipErrorInvalidResourceHandle; + HIP_RETURN(hipErrorInvalidResourceHandle); } amd::HostQueue* hostQueue = as_amd(reinterpret_cast(stream))->asHostQueue(); @@ -146,17 +146,17 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int amd::Command::EventWaitList eventWaitList; cl_int err = amd::clSetEventWaitList(eventWaitList, *hostQueue, 1, &clEvent); if (err != CL_SUCCESS) { - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } amd::Command* command = new amd::Marker(*hostQueue, true, eventWaitList); if (command == NULL) { - return hipErrorOutOfMemory; + HIP_RETURN(hipErrorOutOfMemory); } command->enqueue(); command->release(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipStreamQuery(hipStream_t stream) { @@ -164,7 +164,7 @@ hipError_t hipStreamQuery(hipStream_t stream) { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData, @@ -173,7 +173,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } diff --git a/projects/clr/hipamd/api/hip/hip_surface.cpp b/projects/clr/hipamd/api/hip/hip_surface.cpp index ecbd9e60b9..8cf7e8f35d 100644 --- a/projects/clr/hipamd/api/hip/hip_surface.cpp +++ b/projects/clr/hipamd/api/hip/hip_surface.cpp @@ -36,7 +36,7 @@ hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } @@ -45,5 +45,5 @@ hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) { assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } \ No newline at end of file diff --git a/projects/clr/hipamd/api/hip/hip_texture.cpp b/projects/clr/hipamd/api/hip/hip_texture.cpp index f8bdbe1c89..b3b5c16c10 100644 --- a/projects/clr/hipamd/api/hip/hip_texture.cpp +++ b/projects/clr/hipamd/api/hip/hip_texture.cpp @@ -145,7 +145,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou amd::Device* device = hip::getCurrentContext()->devices()[0]; if (!device->info().imageSupport_) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } amd::Image* image = nullptr; @@ -203,11 +203,11 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou pResDesc->res.pitch2D.width, pResDesc->res.pitch2D.height, 1, pResDesc->res.pitch2D.pitchInBytes, 0); break; - default: return hipErrorInvalidValue; + default: HIP_RETURN(hipErrorInvalidValue); } *pTexObject = reinterpret_cast(as_cl(image)); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) { @@ -215,7 +215,7 @@ hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) { as_amd(reinterpret_cast(textureObject))->release(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, @@ -224,7 +224,7 @@ hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc, @@ -233,7 +233,7 @@ hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, @@ -242,7 +242,7 @@ hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t ihipBindTexture(cl_mem_object_type type, @@ -282,14 +282,14 @@ hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* dev HIP_INIT_API(offset, tex, devPtr, desc, size); if (desc == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } cl_image_format image_format; getChannelOrderAndType(*desc, hipReadModeElementType, &image_format.image_channel_order, &image_format.image_channel_data_type); const amd::Image::Format imageFormat(image_format); - return ihipBindTexture(CL_MEM_OBJECT_IMAGE1D, offset, tex, devPtr, desc, size / imageFormat.getElementSize(), 1, size); + HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE1D, offset, tex, devPtr, desc, size / imageFormat.getElementSize(), 1, size)); } hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* devPtr, @@ -297,7 +297,7 @@ hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* d size_t pitch) { HIP_INIT_API(offset, tex, devPtr, desc, width, height, pitch); - return ihipBindTexture(CL_MEM_OBJECT_IMAGE2D, offset, tex, devPtr, desc, width, height, pitch); + HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE2D, offset, tex, devPtr, desc, width, height, pitch)); } hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array, @@ -306,7 +306,7 @@ hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array, assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t ihipBindTextureToArrayImpl(int dim, enum hipTextureReadMode readMode, @@ -325,7 +325,7 @@ hipError_t hipBindTextureToMipmappedArray(textureReference* tex, assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipUnbindTexture(const textureReference* tex) { @@ -333,7 +333,7 @@ hipError_t hipUnbindTexture(const textureReference* tex) { as_amd(reinterpret_cast(tex->textureObject))->release(); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array) { @@ -341,7 +341,7 @@ hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array) assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* tex) { @@ -349,7 +349,7 @@ hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipGetTextureReference(const textureReference** tex, const void* symbol) { @@ -357,56 +357,56 @@ hipError_t hipGetTextureReference(const textureReference** tex, const void* symb assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipTexRefSetFormat(textureReference* tex, hipArray_Format fmt, int NumPackedComponents) { HIP_INIT_API(tex, fmt, NumPackedComponents); if (tex == nullptr) { - return hipErrorInvalidImage; + HIP_RETURN(hipErrorInvalidImage); } tex->format = fmt; tex->numChannels = NumPackedComponents; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipTexRefSetFlags(textureReference* tex, unsigned int flags) { HIP_INIT_API(tex, flags); if (tex == nullptr) { - return hipErrorInvalidImage; + HIP_RETURN(hipErrorInvalidImage); } tex->normalized = flags; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipTexRefSetFilterMode(textureReference* tex, hipTextureFilterMode fm) { HIP_INIT_API(tex, fm); if (tex == nullptr) { - return hipErrorInvalidImage; + HIP_RETURN(hipErrorInvalidImage); } tex->filterMode = fm; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipTexRefSetAddressMode(textureReference* tex, int dim, hipTextureAddressMode am) { HIP_INIT_API(tex, dim, am); if (tex == nullptr) { - return hipErrorInvalidImage; + HIP_RETURN(hipErrorInvalidImage); } tex->addressMode[dim] = am; - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsigned int flags) { @@ -414,7 +414,7 @@ hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsi assert(0 && "Unimplemented"); - return hipErrorUnknown; + HIP_RETURN(hipErrorUnknown); } hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDeviceptr_t devPtr, @@ -422,7 +422,7 @@ hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDevicep HIP_INIT_API(offset, tex, devPtr, size); if (tex == nullptr) { - return hipErrorInvalidImage; + HIP_RETURN(hipErrorInvalidImage); } cl_image_format image_format; @@ -430,7 +430,7 @@ hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDevicep &image_format.image_channel_order, &image_format.image_channel_data_type); const amd::Image::Format imageFormat(image_format); - return ihipBindTexture(CL_MEM_OBJECT_IMAGE1D, offset, tex, devPtr, nullptr, size / imageFormat.getElementSize(), 1, size); + HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE1D, offset, tex, devPtr, nullptr, size / imageFormat.getElementSize(), 1, size)); } hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPTOR* desc, @@ -438,9 +438,9 @@ hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPT HIP_INIT_API(tex, desc, devPtr, pitch); if (desc == nullptr) { - return hipErrorInvalidValue; + HIP_RETURN(hipErrorInvalidValue); } size_t offset; - return ihipBindTexture(CL_MEM_OBJECT_IMAGE2D, &offset, tex, devPtr, nullptr, desc->width, desc->height, pitch); + HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE2D, &offset, tex, devPtr, nullptr, desc->width, desc->height, pitch)); }