From e76231a330eeaafa095c38b1895b9ff52bdee198 Mon Sep 17 00:00:00 2001 From: foreman Date: Wed, 30 Oct 2019 13:37:03 -0400 Subject: [PATCH] P4 to Git Change 2021977 by kjayapra@1_HIPWS_LNX1_PAL on 2019/10/30 13:27:19 SWDEV-209747 - Improve HIP Error codes and debug messages. Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#24 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_device.cpp#25 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_device_runtime.cpp#20 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.cpp#18 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#80 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#45 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_peer.cpp#8 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#47 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_profile.cpp#4 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_surface.cpp#6 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#25 edit [ROCm/hip commit: e833ad571f2590b52ee73926dd0768f69e5a8b2e] --- projects/hip/api/hip/hip_context.cpp | 14 ++++++------- projects/hip/api/hip/hip_device.cpp | 4 ++-- projects/hip/api/hip/hip_device_runtime.cpp | 10 ++++----- projects/hip/api/hip/hip_event.cpp | 2 +- projects/hip/api/hip/hip_memory.cpp | 20 +++++++++--------- projects/hip/api/hip/hip_module.cpp | 23 +++++++++++---------- projects/hip/api/hip/hip_peer.cpp | 6 +++--- projects/hip/api/hip/hip_platform.cpp | 14 ++++++------- projects/hip/api/hip/hip_profile.cpp | 4 ++-- projects/hip/api/hip/hip_surface.cpp | 2 +- projects/hip/api/hip/hip_texture.cpp | 20 +++++++++--------- 11 files changed, 60 insertions(+), 59 deletions(-) diff --git a/projects/hip/api/hip/hip_context.cpp b/projects/hip/api/hip/hip_context.cpp index ed741c4883..6ff36dc987 100644 --- a/projects/hip/api/hip/hip_context.cpp +++ b/projects/hip/api/hip/hip_context.cpp @@ -268,7 +268,7 @@ hipError_t hipCtxGetDevice(hipDevice_t* device) { HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidContext); } hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) { @@ -276,7 +276,7 @@ hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) { assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) { @@ -284,7 +284,7 @@ hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) { assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) { @@ -292,7 +292,7 @@ hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) { assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) { @@ -300,7 +300,7 @@ hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) { assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipCtxSynchronize(void) { @@ -308,7 +308,7 @@ hipError_t hipCtxSynchronize(void) { assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipCtxGetFlags(unsigned int* flags) { @@ -316,7 +316,7 @@ hipError_t hipCtxGetFlags(unsigned int* flags) { assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int* active) { diff --git a/projects/hip/api/hip/hip_device.cpp b/projects/hip/api/hip/hip_device.cpp index 39c02aa0c7..16d14b2b12 100644 --- a/projects/hip/api/hip/hip_device.cpp +++ b/projects/hip/api/hip/hip_device.cpp @@ -222,7 +222,7 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) { assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av) { @@ -230,5 +230,5 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } diff --git a/projects/hip/api/hip/hip_device_runtime.cpp b/projects/hip/api/hip/hip_device_runtime.cpp index b16be0980d..636ec612c1 100644 --- a/projects/hip/api/hip/hip_device_runtime.cpp +++ b/projects/hip/api/hip/hip_device_runtime.cpp @@ -405,7 +405,7 @@ hipError_t hipDeviceSetCacheConfig ( hipFuncCache_t cacheConfig ) { } hipError_t hipDeviceSetLimit ( hipLimit_t limit, size_t value ) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) { @@ -461,7 +461,7 @@ hipError_t hipGetDeviceCount ( int* count ) { } hipError_t hipGetDeviceFlags ( unsigned int* flags ) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipIpcGetEventHandle ( hipIpcEventHandle_t* handle, hipEvent_t event ) { @@ -469,7 +469,7 @@ hipError_t hipIpcGetEventHandle ( hipIpcEventHandle_t* handle, hipEvent_t event assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipIpcOpenEventHandle ( hipEvent_t* event, hipIpcEventHandle_t handle ) { @@ -477,7 +477,7 @@ hipError_t hipIpcOpenEventHandle ( hipEvent_t* event, hipIpcEventHandle_t handle assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipSetDevice ( int device ) { @@ -512,7 +512,7 @@ hipError_t hipSetValidDevices ( int* device_arr, int len ) { assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount) { diff --git a/projects/hip/api/hip/hip_event.cpp b/projects/hip/api/hip/hip_event.cpp index be2c698f46..f283aa6c5d 100644 --- a/projects/hip/api/hip/hip_event.cpp +++ b/projects/hip/api/hip/hip_event.cpp @@ -110,7 +110,7 @@ hipError_t Event::streamWait(amd::HostQueue* hostQueue, uint flags) { } if (!event_->notifyCmdQueue()) { - return hipErrorUnknown; + return hipErrorLaunchOutOfResources; } amd::Command::EventWaitList eventWaitList; eventWaitList.push_back(event_); diff --git a/projects/hip/api/hip/hip_memory.cpp b/projects/hip/api/hip/hip_memory.cpp index a2ae036d29..6e07ddd30b 100644 --- a/projects/hip/api/hip/hip_memory.cpp +++ b/projects/hip/api/hip/hip_memory.cpp @@ -589,12 +589,12 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t cou /* Get address and size for the global symbol */ if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), &device_ptr, &sym_size)) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidSymbol); } /* Size Check to make sure offset is correct */ if ((offset + count) != sym_size) { - return HIP_RETURN(hipErrorUnknown); + return HIP_RETURN(hipErrorInvalidDevicePointer); } device_ptr = reinterpret_cast
(device_ptr) + offset; @@ -613,12 +613,12 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, /* Get address and size for the global symbol */ if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), &device_ptr, &sym_size)) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidSymbol); } /* Size Check to make sure offset is correct */ if ((offset + count) != sym_size) { - return HIP_RETURN(hipErrorUnknown); + return HIP_RETURN(hipErrorInvalidDevicePointer); } device_ptr = reinterpret_cast
(device_ptr) + offset; @@ -637,12 +637,12 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_ /* Get address and size for the global symbol */ if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), &device_ptr, &sym_size)) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidSymbol); } /* Size Check to make sure offset is correct */ if ((offset + count) != sym_size) { - return HIP_RETURN(hipErrorUnknown); + return HIP_RETURN(hipErrorInvalidDevicePointer); } device_ptr = reinterpret_cast
(device_ptr) + offset; @@ -661,12 +661,12 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co /* Get address and size for the global symbol */ if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), &device_ptr, &sym_size)) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidSymbol); } /* Size Check to make sure offset is correct */ if ((offset + count) != sym_size) { - return HIP_RETURN(hipErrorUnknown); + return HIP_RETURN(hipErrorInvalidDevicePointer); } device_ptr = reinterpret_cast
(device_ptr) + offset; @@ -856,7 +856,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con HIP_INIT_API(hipMemcpy2DToArray, dst, wOffset, hOffset, src, spitch, width, height, kind); if (dst->data == nullptr) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidValue); } hip::syncStreams(); @@ -883,7 +883,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con } if ((wOffset + width > (dpitch)) || width > spitch) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidDevicePointer); } // Create buffer rectangle info structure diff --git a/projects/hip/api/hip/hip_module.cpp b/projects/hip/api/hip/hip_module.cpp index 20a5fbdef1..77a2359614 100644 --- a/projects/hip/api/hip/hip_module.cpp +++ b/projects/hip/api/hip/hip_module.cpp @@ -89,13 +89,13 @@ hipError_t hipModuleUnload(hipModule_t hmod) HIP_INIT_API(hipModuleUnload, hmod); if (hmod == nullptr) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidValue); } amd::Program* program = as_amd(reinterpret_cast(hmod)); if(!ihipModuleUnregisterGlobal(hmod)) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidSymbol); } program->release(); @@ -183,21 +183,21 @@ hipError_t ihipModuleLoadData(hipModule_t *module, const void *image) program->setVarInfoCallBack(&getSvarInfo); if (CL_SUCCESS != program->addDeviceProgram(*hip::getCurrentContext()->devices()[0], image, ElfSize(image))) { - return hipErrorUnknown; + return hipErrorInvalidKernelFile; } - *module = reinterpret_cast(as_cl(program)); + *module = reinterpret_cast(as_cl(program)); if (!ihipModuleRegisterGlobal(program, module)) { - return hipErrorUnknown; + return hipErrorSharedObjectSymbolNotFound; } if (!ihipModuleRegisterUndefined(program, module)) { - return hipErrorUnknown; + return hipErrorSharedObjectSymbolNotFound; } if(CL_SUCCESS != program->build(hip::getCurrentContext()->devices(), nullptr, nullptr, nullptr)) { - return hipErrorUnknown; + return hipErrorSharedObjectInitFailed; } return hipSuccess; @@ -232,7 +232,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h /* Get address and size for the global symbol */ if (!PlatformState::instance().getGlobalVar(name, ihipGetDevice(), dptr, bytes)) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotFound); } HIP_RETURN(hipSuccess); @@ -243,7 +243,7 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) HIP_INIT_API(hipFuncGetAttributes, attr, func); if (!PlatformState::instance().getFuncAttr(func, attr)) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidDeviceFunction); } HIP_RETURN(hipSuccess); @@ -425,7 +425,7 @@ hipError_t hipLaunchCooperativeKernel(const void* f, int deviceId = ihipGetDevice(); hipFunction_t func = PlatformState::instance().getFunc(f, deviceId); if (func == nullptr) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidDeviceFunction); } HIP_RETURN(ihipModuleLaunchKernel(func, gridDim.x * blockDim.x, gridDim.y * blockDim.y, gridDim.z * blockDim.z, @@ -469,6 +469,7 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL } } if (func == nullptr) { + result = hipErrorInvalidDeviceFunction; HIP_RETURN(result); } @@ -515,7 +516,7 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const /* Get address and size for the global symbol */ if (!PlatformState::instance().getTexRef(name, texRef)) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotFound); } HIP_RETURN(hipSuccess); diff --git a/projects/hip/api/hip/hip_peer.cpp b/projects/hip/api/hip/hip_peer.cpp index 95cf206e9c..65d92de363 100644 --- a/projects/hip/api/hip/hip_peer.cpp +++ b/projects/hip/api/hip/hip_peer.cpp @@ -29,7 +29,7 @@ hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } 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"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } 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"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId) { diff --git a/projects/hip/api/hip/hip_platform.cpp b/projects/hip/api/hip/hip_platform.cpp index 777a710b11..bd6cd55317 100644 --- a/projects/hip/api/hip/hip_platform.cpp +++ b/projects/hip/api/hip/hip_platform.cpp @@ -480,7 +480,7 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction) int deviceId = ihipGetDevice(); hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId); if (func == nullptr) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidDeviceFunction); } ihipExec_t exec; @@ -512,7 +512,7 @@ extern "C" hipError_t hipLaunchKernel(const void *hostFunction, int deviceId = ihipGetDevice(); hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId); if (func == nullptr) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidDeviceFunction); } HIP_RETURN(hipModuleLaunchKernel(func, gridDim.x, gridDim.y, gridDim.z, @@ -524,7 +524,7 @@ extern "C" hipError_t hipLaunchKernel(const void *hostFunction, hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { size_t size = 0; if(!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), devPtr, &size)) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidSymbol); } HIP_RETURN(hipSuccess); } @@ -532,7 +532,7 @@ hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { hipError_t hipGetSymbolSize(size_t* sizePtr, const void* symbolName) { hipDeviceptr_t devPtr = nullptr; if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), &devPtr, sizePtr)) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidSymbol); } HIP_RETURN(hipSuccess); } @@ -549,11 +549,11 @@ hipError_t ihipCreateGlobalVarObj(const char* name, hipModule_t hmod, amd::Memor dev_program = program->getDeviceProgram(*hip::getCurrentContext()->devices()[0]); if (dev_program == nullptr) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidDeviceFunction); } /* Find the global Symbols */ if(!dev_program->createGlobalVarObj(amd_mem_obj, dptr, bytes, name)) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidSymbol); } HIP_RETURN(hipSuccess); @@ -571,7 +571,7 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks, int deviceId = ihipGetDevice(); hipFunction_t func = PlatformState::instance().getFunc(f, deviceId); if (func == nullptr) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidDeviceFunction); } hip::Function* function = hip::Function::asFunction(func); diff --git a/projects/hip/api/hip/hip_profile.cpp b/projects/hip/api/hip/hip_profile.cpp index b8c8f28e94..7ff93445c4 100644 --- a/projects/hip/api/hip/hip_profile.cpp +++ b/projects/hip/api/hip/hip_profile.cpp @@ -29,7 +29,7 @@ hipError_t hipProfilerStart() { assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } @@ -38,5 +38,5 @@ hipError_t hipProfilerStop() { assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } \ No newline at end of file diff --git a/projects/hip/api/hip/hip_surface.cpp b/projects/hip/api/hip/hip_surface.cpp index 42741add7a..2a7b58d191 100644 --- a/projects/hip/api/hip/hip_surface.cpp +++ b/projects/hip/api/hip/hip_surface.cpp @@ -92,5 +92,5 @@ hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) { HIP_RETURN(hipFree(reinterpret_cast(surfaceObject))); } - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidValue); } diff --git a/projects/hip/api/hip/hip_texture.cpp b/projects/hip/api/hip/hip_texture.cpp index b7e6293b95..2634d799c1 100644 --- a/projects/hip/api/hip/hip_texture.cpp +++ b/projects/hip/api/hip/hip_texture.cpp @@ -334,7 +334,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou if (!image->create()) { delete image; - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorMemoryAllocation); } amd::Sampler* sampler = fillSamplerDescriptor(pTexDesc->addressMode[0], pTexDesc->filterMode, pTexDesc->normalizedCoords); @@ -371,7 +371,7 @@ hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, memcpy(pResDesc, &(texture->resDesc), sizeof(hipResourceDesc)); } - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidValue); } hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc, @@ -380,7 +380,7 @@ hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, @@ -389,7 +389,7 @@ hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t ihipBindTexture(cl_mem_object_type type, @@ -416,7 +416,7 @@ hipError_t ihipBindTexture(cl_mem_object_type type, if (!image->create()) { delete image; - return hipErrorUnknown; + return hipErrorMemoryAllocation; } *offset = 0; @@ -456,7 +456,7 @@ hipError_t ihipBindTexture(cl_mem_object_type type, return hipSuccess; } - return hipErrorUnknown; + return hipErrorInvalidValue; } hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr, @@ -488,7 +488,7 @@ hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array, assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t ihipBindTextureImpl(TlsData* tls, int dim, enum hipTextureReadMode readMode, size_t* offset, @@ -532,7 +532,7 @@ hipError_t hipBindTextureToMipmappedArray(textureReference* tex, assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipUnbindTexture(const textureReference* tex) { @@ -570,7 +570,7 @@ hipError_t hipGetTextureReference(const textureReference** tex, const void* symb assert(0 && "Unimplemented"); - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorNotSupported); } hipError_t hipTexRefSetFormat(textureReference* tex, hipArray_Format fmt, int NumPackedComponents) { @@ -649,7 +649,7 @@ hipError_t hipTexRefGetArray(hipArray_t* array, textureReference tex) { } if (texture->resDesc.res.array.array == nullptr) { - HIP_RETURN(hipErrorUnknown); + HIP_RETURN(hipErrorInvalidValue); } **array = *(texture->resDesc.res.array.array);