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: e833ad571f]

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) {

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);
 }

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) {

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_);

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<address>(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<address>(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<address>(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<address>(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

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<cl_program>(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<hipModule_t>(as_cl(program));
+  *module = reinterpret_cast<hipModule_t>(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);

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) {

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);

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);
 }

projects/hip/api/hip/hip_surface.cpp

@@ -92,5 +92,5 @@ hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) {
     HIP_RETURN(hipFree(reinterpret_cast<void*>(surfaceObject)));
   }
 
-  HIP_RETURN(hipErrorUnknown);
+  HIP_RETURN(hipErrorInvalidValue);
 }

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);