P4 to Git Change 1548476 by cpaquot@cpaquot-ocl-lc-lnx on 2018/05/01 15:50:51

SWDEV-145570 - [HIP]
	Added support for null stream avoiding creating/destroying dummy streams.
	Added basic event class for hipEvent* support.
	Refactored some common functionality: No more direct access to g_context.
	Support hipStreamSynchronize(0).

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#9 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_device_runtime.cpp#7 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.cpp#3 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.hpp#1 add
... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#8 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#15 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#8 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#9 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#4 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#7 edit


[ROCm/hip commit: 61378a359c]

projects/hip/api/hip/hip_context.cpp

@@ -27,13 +27,17 @@ THE SOFTWARE.
 #include <stack>
 #include <thread>
 
-thread_local amd::Context* g_context = nullptr;
-thread_local std::stack<amd::Context*> g_ctxtStack;
-
 std::vector<amd::Context*> g_devices;
 std::once_flag g_ihipInitialized;
 
-void ihipInit() {
+namespace hip {
+
+thread_local amd::Context* g_context = nullptr;
+thread_local std::stack<amd::Context*> g_ctxtStack;
+
+std::map<amd::Context*,amd::HostQueue*> g_nullStreams;
+
+void init() {
   if (!amd::Runtime::initialized()) {
     amd::Runtime::init();
   }
@@ -54,6 +58,32 @@ void ihipInit() {
   }
 }
 
+amd::Context* getCurrentContext() {
+  return g_context;
+}
+
+void setCurrentContext(unsigned int index) {
+  assert(index<g_devices.size());
+  g_context = g_devices[index];
+}
+
+amd::HostQueue* getNullStream() {
+  auto stream = g_nullStreams.find(getCurrentContext());
+  if (stream == g_nullStreams.end()) {
+    amd::Device* device = getCurrentContext()->devices()[0];
+    amd::HostQueue* queue = new amd::HostQueue(*hip::getCurrentContext(), *device, 0,
+                                               amd::CommandQueue::RealTimeDisabled,
+                                               amd::CommandQueue::Priority::Normal);
+    g_nullStreams[getCurrentContext()] = queue;
+    return queue;
+  }
+  return stream->second;
+}
+
+};
+
+using namespace hip;
+
 hipError_t hipInit(unsigned int flags) {
   HIP_INIT_API(flags);
 
@@ -84,11 +114,11 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) {
       g_ctxtStack.pop();
     }
   } else {
-    g_context = reinterpret_cast<amd::Context*>(as_amd(ctx));
+    hip::g_context = reinterpret_cast<amd::Context*>(as_amd(ctx));
     if(!g_ctxtStack.empty()) {
       g_ctxtStack.pop();
     }
-    g_ctxtStack.push(g_context);
+    g_ctxtStack.push(hip::getCurrentContext());
   }
 
   return hipSuccess;
@@ -97,7 +127,7 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) {
 hipError_t hipCtxGetCurrent(hipCtx_t* ctx) {
   HIP_INIT_API(ctx);
 
-  *ctx = reinterpret_cast<hipCtx_t>(g_context);
+  *ctx = reinterpret_cast<hipCtx_t>(hip::getCurrentContext());
 
   return hipSuccess;
 }
@@ -164,8 +194,8 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) {
     return hipErrorInvalidContext;
   }
 
-  g_context = amdContext;
-  g_ctxtStack.push(g_context);
+  hip::g_context = amdContext;
+  g_ctxtStack.push(hip::getCurrentContext());
 
   return hipSuccess;
 }
@@ -191,7 +221,7 @@ hipError_t hipCtxGetDevice(hipDevice_t* device) {
 
   if (device != nullptr) {
     for (unsigned int i = 0; i < g_devices.size(); i++) {
-      if (g_devices[i] == g_context) {
+      if (g_devices[i] == hip::getCurrentContext()) {
         *device = static_cast<hipDevice_t>(i);
         return hipSuccess;
       }

projects/hip/api/hip/hip_device_runtime.cpp

@@ -383,7 +383,7 @@ hipError_t hipGetDevice ( int* deviceId ) {
 
   if (deviceId != nullptr) {
     for (unsigned int i = 0; i < g_devices.size(); i++) {
-      if (g_devices[i] == g_context) {
+      if (g_devices[i] == hip::getCurrentContext()) {
         *deviceId = i;
         return hipSuccess;
       }
@@ -425,7 +425,7 @@ hipError_t hipSetDevice ( int  device ) {
   HIP_INIT_API(device);
 
   if (static_cast<unsigned int>(device) < g_devices.size()) {
-    g_context = g_devices[device];
+    hip::setCurrentContext(device);
 
     return hipSuccess;
   }

projects/hip/api/hip/hip_event.cpp

@@ -22,30 +22,42 @@ THE SOFTWARE.
 
 #include <hip/hip_runtime.h>
 
-#include "hip_internal.hpp"
+#include "hip_event.hpp"
 
 hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) {
   HIP_INIT_API(event, flags);
 
-  assert(0 && "Unimplemented");
+  hip::Event* e = new hip::Event(flags);
 
-  return hipErrorUnknown;
+  if (e == nullptr) {
+    return hipErrorOutOfMemory;
+  }
+
+  *event = reinterpret_cast<hipEvent_t>(e);
+
+  return hipSuccess;
 }
 
 hipError_t hipEventCreate(hipEvent_t* event) {
   HIP_INIT_API(event);
 
-  assert(0 && "Unimplemented");
+  hip::Event* e = new hip::Event(0);
 
-  return hipErrorUnknown;
+  if (e == nullptr) {
+    return hipErrorOutOfMemory;
+  }
+
+  *event = reinterpret_cast<hipEvent_t>(e);
+
+  return hipSuccess;
 }
 
 hipError_t hipEventDestroy(hipEvent_t event) {
   HIP_INIT_API(event);
 
-  assert(0 && "Unimplemented");
+  delete reinterpret_cast<hip::Event*>(event);
 
-  return hipErrorUnknown;
+  return hipSuccess;
 }
 
 hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) {

projects/hip/api/hip/hip_event.hpp

@@ -0,0 +1,40 @@
+/*
+Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_EVENT_H
+#define HIP_EVENT_H
+
+#include "hip_internal.hpp"
+
+namespace hip {
+
+class Event {
+public:
+  Event(unsigned int flags) : flags(flags) {}
+  ~Event() {}
+  unsigned int flags;
+private:
+};
+
+};
+
+#endif // HIP_EVEMT_H

projects/hip/api/hip/hip_internal.hpp

@@ -28,7 +28,7 @@ THE SOFTWARE.
 #include <thread>
 
 #define HIP_INIT() \
-  std::call_once(g_ihipInitialized, ihipInit);
+  std::call_once(g_ihipInitialized, hip::init);
 
 
 // This macro should be called at the beginning of every HIP API.
@@ -46,10 +46,17 @@ class accelerator_view;
 };
 
 extern std::once_flag g_ihipInitialized;
-extern thread_local amd::Context* g_context;
+
+namespace hip {
+  extern void init();
+
+  extern amd::Context* getCurrentContext();
+  extern void setCurrentContext(unsigned int index);
+
+  extern amd::HostQueue* getNullStream();
+};
 extern std::vector<amd::Context*> g_devices;
 
 extern hipError_t ihipDeviceGetCount(int* count);
-extern void ihipInit();
 
 #endif // HIP_SRC_HIP_INTERNAL_H

projects/hip/api/hip/hip_memory.cpp

@@ -39,11 +39,11 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
     return hipErrorInvalidValue;
   }
 
-  if (g_context->devices()[0]->info().maxMemAllocSize_ < sizeBytes) {
+  if (hip::getCurrentContext()->devices()[0]->info().maxMemAllocSize_ < sizeBytes) {
     return hipErrorOutOfMemory;
   }
 
-  *ptr = amd::SvmBuffer::malloc(*g_context, flags, sizeBytes, g_context->devices()[0]->info().memBaseAddrAlign_);
+  *ptr = amd::SvmBuffer::malloc(*hip::getCurrentContext(), flags, sizeBytes, hip::getCurrentContext()->devices()[0]->info().memBaseAddrAlign_);
   if (!*ptr) {
     return hipErrorOutOfMemory;
   }
@@ -65,7 +65,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
 
 hipError_t hipFree(void* ptr) {
   if (amd::SvmBuffer::malloced(ptr)) {
-    amd::SvmBuffer::free(*g_context, ptr);
+    amd::SvmBuffer::free(*hip::getCurrentContext(), ptr);
     return hipSuccess;
   }
   return hipErrorInvalidValue;
@@ -74,11 +74,8 @@ hipError_t hipFree(void* ptr) {
 hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
   HIP_INIT_API(dst, src, sizeBytes, kind);
 
-  amd::Device* device = g_context->devices()[0];
+  amd::HostQueue* queue = hip::getNullStream();
 
-  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
-                                             amd::CommandQueue::RealTimeDisabled,
-                                             amd::CommandQueue::Priority::Normal);
   if (!queue) {
     return hipErrorOutOfMemory;
   }
@@ -119,8 +116,6 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
   command->awaitCompletion();
   command->release();
 
-  queue->release();
-
   return hipSuccess;
 }
 
@@ -135,11 +130,8 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st
 hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
   HIP_INIT_API(dst, value, sizeBytes);
 
-  amd::Device* device = g_context->devices()[0];
+  amd::HostQueue* queue = hip::getNullStream();
 
-  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
-                                             amd::CommandQueue::RealTimeDisabled,
-                                             amd::CommandQueue::Priority::Normal);
   if (!queue) {
     return hipErrorOutOfMemory;
   }
@@ -162,8 +154,6 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
   command->awaitCompletion();
   command->release();
 
-  queue->release();
-
   return hipSuccess;
 }
 
@@ -185,7 +175,7 @@ hipError_t hipHostFree(void* ptr) {
   HIP_INIT_API(ptr);
 
   if (amd::SvmBuffer::malloced(ptr)) {
-    amd::SvmBuffer::free(*g_context, ptr);
+    amd::SvmBuffer::free(*hip::getCurrentContext(), ptr);
     return hipSuccess;
   }
   return hipErrorInvalidValue;
@@ -195,7 +185,7 @@ hipError_t hipFreeArray(hipArray* array) {
   HIP_INIT_API(array);
 
   if (amd::SvmBuffer::malloced(array->data)) {
-    amd::SvmBuffer::free(*g_context, array->data);
+    amd::SvmBuffer::free(*hip::getCurrentContext(), array->data);
     return hipSuccess;
   }
   return hipErrorInvalidValue;
@@ -222,7 +212,7 @@ hipError_t hipMemGetInfo(size_t* free, size_t* total) {
   HIP_INIT_API(free, total);
 
   size_t freeMemory[2];
-  amd::Device* device = g_context->devices()[0];
+  amd::Device* device = hip::getCurrentContext()->devices()[0];
   if(!device) {
     return hipErrorInvalidDevice;
   }
@@ -240,7 +230,7 @@ return hipSuccess;
 hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height, size_t depth,
                            cl_mem_object_type imageType, const cl_image_format* image_format) {
 
-  amd::Device* device = g_context->devices()[0];
+  amd::Device* device = hip::getCurrentContext()->devices()[0];
 
   if ((width == 0) || (height == 0)) {
     *ptr = nullptr;
@@ -251,7 +241,7 @@ hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t heigh
     return hipErrorInvalidValue;
   }
 
-  if (g_context->devices()[0]->info().maxMemAllocSize_ < (width * height)) {
+  if (device->info().maxMemAllocSize_ < (width * height)) {
     return hipErrorOutOfMemory;
   }
 
@@ -260,8 +250,8 @@ hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t heigh
   *pitch = width * imageFormat.getElementSize();
 
   size_t sizeBytes = *pitch * height * depth;
-  *ptr = amd::SvmBuffer::malloc(*g_context, CL_MEM_SVM_FINE_GRAIN_BUFFER, sizeBytes,
-                                g_context->devices()[0]->info().memBaseAddrAlign_);
+  *ptr = amd::SvmBuffer::malloc(*hip::getCurrentContext(), CL_MEM_SVM_FINE_GRAIN_BUFFER, sizeBytes,
+                                device->info().memBaseAddrAlign_);
 
   if (!*ptr) {
     return hipErrorMemoryAllocation;
@@ -559,11 +549,7 @@ hipError_t hipMemcpyToArray(hipArray* dstArray, size_t wOffset, size_t hOffset,
                             size_t count, hipMemcpyKind kind) {
   HIP_INIT_API(dstArray, wOffset, hOffset, src, count, kind);
 
-  amd::Device* device = g_context->devices()[0];
-
-  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
-                                             amd::CommandQueue::RealTimeDisabled,
-                                             amd::CommandQueue::Priority::Normal);
+  amd::HostQueue* queue = hip::getNullStream();
   if (!queue) {
     return hipErrorOutOfMemory;
   }
@@ -597,8 +583,6 @@ hipError_t hipMemcpyToArray(hipArray* dstArray, size_t wOffset, size_t hOffset,
   command->awaitCompletion();
   command->release();
 
-  queue->release();
-
   return hipSuccess;
 }
 
@@ -606,11 +590,7 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs
                               size_t count, hipMemcpyKind kind) {
   HIP_INIT_API(dst, srcArray, wOffset, hOffset, count, kind);
 
-  amd::Device* device = g_context->devices()[0];
-
-  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
-                                             amd::CommandQueue::RealTimeDisabled,
-                                             amd::CommandQueue::Priority::Normal);
+  amd::HostQueue* queue = hip::getNullStream();
   if (!queue) {
     return hipErrorOutOfMemory;
   }
@@ -644,19 +624,13 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs
   command->awaitCompletion();
   command->release();
 
-  queue->release();
-
   return hipSuccess;
 }
 
 hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) {
   HIP_INIT_API(dstArray, dstOffset, srcHost, count);
 
-  amd::Device* device = g_context->devices()[0];
-
-  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
-                                             amd::CommandQueue::RealTimeDisabled,
-                                             amd::CommandQueue::Priority::Normal);
+  amd::HostQueue* queue = hip::getNullStream();
   if (!queue) {
     return hipErrorOutOfMemory;
   }
@@ -674,19 +648,13 @@ hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHo
   command->awaitCompletion();
   command->release();
 
-  queue->release();
-
   return hipSuccess;
 }
 
 hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count) {
   HIP_INIT_API(dst, srcArray, srcOffset, count);
 
-  amd::Device* device = g_context->devices()[0];
-
-  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
-                                             amd::CommandQueue::RealTimeDisabled,
-                                             amd::CommandQueue::Priority::Normal);
+  amd::HostQueue* queue = hip::getNullStream();
   if (!queue) {
     return hipErrorOutOfMemory;
   }
@@ -704,19 +672,13 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t
   command->awaitCompletion();
   command->release();
 
-  queue->release();
-
   return hipSuccess;
 }
 
 hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) {
   HIP_INIT_API(p);
 
-  amd::Device* device = g_context->devices()[0];
-
-  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
-                                             amd::CommandQueue::RealTimeDisabled,
-                                             amd::CommandQueue::Priority::Normal);
+  amd::HostQueue* queue = hip::getNullStream();
   if (!queue) {
     return hipErrorOutOfMemory;
   }
@@ -826,19 +788,13 @@ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) {
   command->awaitCompletion();
   command->release();
 
-  queue->release();
-
   return hipSuccess;
 }
 
 hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) {
   HIP_INIT_API(dst, pitch, value, width, height);
 
-  amd::Device* device = g_context->devices()[0];
-
-  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
-                                             amd::CommandQueue::RealTimeDisabled,
-                                             amd::CommandQueue::Priority::Normal);
+  amd::HostQueue* queue = hip::getNullStream();
   if (!queue) {
     return hipErrorOutOfMemory;
   }
@@ -862,19 +818,13 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t
   command->awaitCompletion();
   command->release();
 
-  queue->release();
-
   return hipSuccess;
 }
 
 hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) {
   HIP_INIT_API(dst, value, sizeBytes);
 
-  amd::Device* device = g_context->devices()[0];
-
-  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
-                                             amd::CommandQueue::RealTimeDisabled,
-                                             amd::CommandQueue::Priority::Normal);
+  amd::HostQueue* queue = hip::getNullStream();
   if (!queue) {
     return hipErrorOutOfMemory;
   }
@@ -896,8 +846,6 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes
   command->awaitCompletion();
   command->release();
 
-  queue->release();
-
   return hipSuccess;
 }
 
@@ -938,9 +886,13 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel
 hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsigned flags) {
   HIP_INIT_API(devicePointer, hostPointer, flags);
 
-  assert(0 && "Unimplemented");
+  if (!amd::SvmBuffer::malloced(hostPointer)) {
+    return hipErrorInvalidValue;
+  }
+  // right now we have SVM
+  *devicePointer = hostPointer;
 
-  return hipErrorUnknown;
+  return hipSuccess;
 }
 
 hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr) {

projects/hip/api/hip/hip_module.cpp

@@ -94,13 +94,13 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image)
 
 hipError_t ihipModuleLoadData(hipModule_t *module, const void *image)
 {
-  amd::Program* program = new amd::Program(*g_context);
+  amd::Program* program = new amd::Program(*hip::getCurrentContext());
   if (program == NULL) {
     return hipErrorOutOfMemory;
   }
 
-  if (CL_SUCCESS != program->addDeviceProgram(*g_context->devices()[0], image, ElfSize(image)) ||
-    CL_SUCCESS != program->build(g_context->devices(), nullptr, nullptr, nullptr)) {
+  if (CL_SUCCESS != program->addDeviceProgram(*hip::getCurrentContext()->devices()[0], image, ElfSize(image)) ||
+    CL_SUCCESS != program->build(hip::getCurrentContext()->devices(), nullptr, nullptr, nullptr)) {
     return hipErrorUnknown;
   }
 
@@ -142,13 +142,11 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
                kernelParams, extra);
 
   amd::Kernel* kernel = as_amd(reinterpret_cast<cl_kernel>(f));
-  amd::Device* device = g_context->devices()[0];
+  amd::Device* device = hip::getCurrentContext()->devices()[0];
 
   amd::HostQueue* queue;
   if (hStream == nullptr) {
-    queue = new amd::HostQueue(*g_context, *device, 0,
-                               amd::CommandQueue::RealTimeDisabled,
-                               amd::CommandQueue::Priority::Normal);
+    queue = hip::getNullStream();
   } else {
     queue = as_amd(reinterpret_cast<cl_command_queue>(hStream))->asHostQueue();
   }
@@ -200,10 +198,6 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
   command->awaitCompletion();
   command->release();
 
-  if (hStream == nullptr) {
-    queue->release();
-  }
-
   return hipSuccess;
 }
 

projects/hip/api/hip/hip_platform.cpp

@@ -68,7 +68,7 @@ extern "C" hipModule_t __hipRegisterFatBinary(const void* data)
     return nullptr;
   }
 
-  amd::Program* program = new amd::Program(*g_context);
+  amd::Program* program = new amd::Program(*hip::getCurrentContext());
   if (!program)
     return nullptr;
 
@@ -84,15 +84,15 @@ extern "C" hipModule_t __hipRegisterFatBinary(const void* data)
 
     std::string target(desc->triple + sizeof(OPENMP_AMDGCN_AMDHSA_TRIPLE),
                        desc->tripleSize - sizeof(OPENMP_AMDGCN_AMDHSA_TRIPLE));
-    if (target.compare(g_context->devices()[0]->info().name_))
+    if (target.compare(hip::getCurrentContext()->devices()[0]->info().name_))
       continue;
 
     const void *image = reinterpret_cast<const void*>(
         reinterpret_cast<uintptr_t>(obheader) + desc->offset);
     size_t size = desc->size;
 
-    if (CL_SUCCESS == program->addDeviceProgram(*g_context->devices()[0], image, size) &&
-        CL_SUCCESS == program->build(g_context->devices(), nullptr, nullptr, nullptr))
+    if (CL_SUCCESS == program->addDeviceProgram(*hip::getCurrentContext()->devices()[0], image, size) &&
+        CL_SUCCESS == program->build(hip::getCurrentContext()->devices(), nullptr, nullptr, nullptr))
       break;
   }
 
@@ -372,7 +372,7 @@ const std::vector<hipModule_t>& modules() {
           std::string target(desc->triple + sizeof(HCC_AMDGCN_AMDHSA_TRIPLE),
                              desc->tripleSize - sizeof(HCC_AMDGCN_AMDHSA_TRIPLE));
 
-          if (!target.compare(g_context->devices()[0]->info().name_)) {
+          if (!target.compare(hip::getCurrentContext()->devices()[0]->info().name_)) {
             hipModule_t module;
             if (hipSuccess == hipModuleLoadData(&module, reinterpret_cast<const void*>(
                 reinterpret_cast<uintptr_t>(obheader) + desc->offset)))

projects/hip/api/hip/hip_stream.cpp

@@ -27,9 +27,9 @@ THE SOFTWARE.
 static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
   assert(flags == 0); // we don't handle flags yet
 
-  amd::Device* device = g_context->devices()[0];
+  amd::Device* device = hip::getCurrentContext()->devices()[0];
 
-  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
+  amd::HostQueue* queue = new amd::HostQueue(*hip::getCurrentContext(), *device, 0,
                                              amd::CommandQueue::RealTimeDisabled,
                                              amd::CommandQueue::Priority::Normal);
 
@@ -68,7 +68,14 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) {
 hipError_t hipStreamSynchronize(hipStream_t stream) {
   HIP_INIT_API(stream);
 
-  amd::HostQueue* hostQueue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
+  amd::HostQueue* hostQueue;
+
+  if (stream == nullptr) {
+    hostQueue = hip::getNullStream();
+  } else {
+    hostQueue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
+  }
+
   if (hostQueue == nullptr) {
     return hipErrorUnknown;
   }
@@ -82,6 +89,10 @@ hipError_t hipStreamSynchronize(hipStream_t stream) {
 hipError_t hipStreamDestroy(hipStream_t stream) {
   HIP_INIT_API(stream);
 
+  if (stream == nullptr) {
+    return hipErrorInvalidResourceHandle;
+  }
+
   as_amd(reinterpret_cast<cl_command_queue>(stream))->release();
 
   return hipSuccess;

projects/hip/api/hip/hip_texture.cpp

@@ -142,7 +142,9 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
                                   const hipResourceViewDesc* pResViewDesc) {
   HIP_INIT_API(pTexObject, pResDesc, pTexDesc, pResViewDesc);
 
-  if (!g_context->devices()[0]->info().imageSupport_) {
+  amd::Device* device = hip::getCurrentContext()->devices()[0];
+
+  if (!device->info().imageSupport_) {
     return hipErrorInvalidValue;
   }
 
@@ -173,8 +175,9 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
           case hipArrayTextureGather:
           case hipArrayDefault:
           default:
-            image = new (*g_context) amd::Image(*memory->asBuffer(), CL_MEM_OBJECT_IMAGE2D, memory->getMemFlags(), imageFormat,
-                                          pResDesc->res.array.array->width, pResDesc->res.array.array->height, 1, 0, 0);
+            image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
+              CL_MEM_OBJECT_IMAGE2D, memory->getMemFlags(), imageFormat,
+              pResDesc->res.array.array->width, pResDesc->res.array.array->height, 1, 0, 0);
             break;
         }
       }
@@ -186,17 +189,19 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
       assert(pResViewDesc == nullptr);
 
       memory = amd::SvmManager::FindSvmBuffer(pResDesc->res.linear.devPtr);
-      image = new (*g_context) amd::Image(*memory->asBuffer(), CL_MEM_OBJECT_IMAGE1D, memory->getMemFlags(), imageFormat,
-                                          pResDesc->res.linear.sizeInBytes / imageFormat.getElementSize(), 1, 1,
-                                          pResDesc->res.linear.sizeInBytes, 0);
+      image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
+        CL_MEM_OBJECT_IMAGE1D, memory->getMemFlags(), imageFormat,
+        pResDesc->res.linear.sizeInBytes / imageFormat.getElementSize(), 1, 1,
+        pResDesc->res.linear.sizeInBytes, 0);
       break;
     case hipResourceTypePitch2D:
       assert(pResViewDesc == nullptr);
 
       memory = amd::SvmManager::FindSvmBuffer(pResDesc->res.pitch2D.devPtr);
-      image = new (*g_context) amd::Image(*memory->asBuffer(), CL_MEM_OBJECT_IMAGE2D, memory->getMemFlags(), imageFormat,
-                                          pResDesc->res.pitch2D.width, pResDesc->res.pitch2D.height, 1,
-                                          pResDesc->res.pitch2D.pitchInBytes, 0);
+      image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
+        CL_MEM_OBJECT_IMAGE2D, memory->getMemFlags(), imageFormat,
+        pResDesc->res.pitch2D.width, pResDesc->res.pitch2D.height, 1,
+        pResDesc->res.pitch2D.pitchInBytes, 0);
       break;
     default: return hipErrorInvalidValue;
   }
@@ -247,7 +252,7 @@ hipError_t ihipBindTexture(cl_mem_object_type type,
   if (tex == nullptr) {
     return hipErrorInvalidImage;
   }
-  if (g_context) {
+  if (hip::getCurrentContext()) {
     cl_image_format image_format;
 
     if (nullptr == desc) {
@@ -260,8 +265,8 @@ hipError_t ihipBindTexture(cl_mem_object_type type,
     const amd::Image::Format imageFormat(image_format);
 
     amd::Memory* memory = amd::SvmManager::FindSvmBuffer(devPtr);
-    amd::Image* image = new (*g_context) amd::Image(*memory->asBuffer(), type, memory->getMemFlags(),
-                                                    imageFormat, width, height, 1, pitch, 0);
+    amd::Image* image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
+      type, memory->getMemFlags(), imageFormat, width, height, 1, pitch, 0);
 
     *offset = 0;
     if (tex->textureObject) {