From 0e2eaa1aa1ad27812a3f6dc2c9fbff6ecfc2e445 Mon Sep 17 00:00:00 2001 From: foreman Date: Tue, 1 May 2018 18:10:09 -0400 Subject: [PATCH] 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: 61378a359c8dcec55b853b56574cf24508357abf] --- projects/hip/api/hip/hip_context.cpp | 50 ++++++++--- projects/hip/api/hip/hip_device_runtime.cpp | 4 +- projects/hip/api/hip/hip_event.cpp | 26 ++++-- projects/hip/api/hip/hip_event.hpp | 40 +++++++++ projects/hip/api/hip/hip_internal.hpp | 13 ++- projects/hip/api/hip/hip_memory.cpp | 98 ++++++--------------- projects/hip/api/hip/hip_module.cpp | 16 ++-- projects/hip/api/hip/hip_platform.cpp | 10 +-- projects/hip/api/hip/hip_stream.cpp | 17 +++- projects/hip/api/hip/hip_texture.cpp | 29 +++--- 10 files changed, 177 insertions(+), 126 deletions(-) create mode 100644 projects/hip/api/hip/hip_event.hpp diff --git a/projects/hip/api/hip/hip_context.cpp b/projects/hip/api/hip/hip_context.cpp index e25a87bde8..46e4c864ff 100644 --- a/projects/hip/api/hip/hip_context.cpp +++ b/projects/hip/api/hip/hip_context.cpp @@ -27,13 +27,17 @@ THE SOFTWARE. #include #include -thread_local amd::Context* g_context = nullptr; -thread_local std::stack g_ctxtStack; - std::vector g_devices; std::once_flag g_ihipInitialized; -void ihipInit() { +namespace hip { + +thread_local amd::Context* g_context = nullptr; +thread_local std::stack g_ctxtStack; + +std::map 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(indexdevices()[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(as_amd(ctx)); + hip::g_context = reinterpret_cast(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(g_context); + *ctx = reinterpret_cast(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(i); return hipSuccess; } diff --git a/projects/hip/api/hip/hip_device_runtime.cpp b/projects/hip/api/hip/hip_device_runtime.cpp index effce5974f..77d90d2cb0 100644 --- a/projects/hip/api/hip/hip_device_runtime.cpp +++ b/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(device) < g_devices.size()) { - g_context = g_devices[device]; + hip::setCurrentContext(device); return hipSuccess; } diff --git a/projects/hip/api/hip/hip_event.cpp b/projects/hip/api/hip/hip_event.cpp index 1fe7be9e2f..b9930636bb 100644 --- a/projects/hip/api/hip/hip_event.cpp +++ b/projects/hip/api/hip/hip_event.cpp @@ -22,30 +22,42 @@ THE SOFTWARE. #include -#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(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(e); + + return hipSuccess; } hipError_t hipEventDestroy(hipEvent_t event) { HIP_INIT_API(event); - assert(0 && "Unimplemented"); + delete reinterpret_cast(event); - return hipErrorUnknown; + return hipSuccess; } hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) { diff --git a/projects/hip/api/hip/hip_event.hpp b/projects/hip/api/hip/hip_event.hpp new file mode 100644 index 0000000000..3ac1ea8bfe --- /dev/null +++ b/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 diff --git a/projects/hip/api/hip/hip_internal.hpp b/projects/hip/api/hip/hip_internal.hpp index 2512e35c98..489b8b620f 100644 --- a/projects/hip/api/hip/hip_internal.hpp +++ b/projects/hip/api/hip/hip_internal.hpp @@ -28,7 +28,7 @@ THE SOFTWARE. #include #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 g_devices; extern hipError_t ihipDeviceGetCount(int* count); -extern void ihipInit(); #endif // HIP_SRC_HIP_INTERNAL_H diff --git a/projects/hip/api/hip/hip_memory.cpp b/projects/hip/api/hip/hip_memory.cpp index 830865fb93..3a4af23dee 100644 --- a/projects/hip/api/hip/hip_memory.cpp +++ b/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) { diff --git a/projects/hip/api/hip/hip_module.cpp b/projects/hip/api/hip/hip_module.cpp index 97723d52dc..0a5675114c 100644 --- a/projects/hip/api/hip/hip_module.cpp +++ b/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(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(hStream))->asHostQueue(); } @@ -200,10 +198,6 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, command->awaitCompletion(); command->release(); - if (hStream == nullptr) { - queue->release(); - } - return hipSuccess; } diff --git a/projects/hip/api/hip/hip_platform.cpp b/projects/hip/api/hip/hip_platform.cpp index 7152458d38..9f7ccbe658 100644 --- a/projects/hip/api/hip/hip_platform.cpp +++ b/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( reinterpret_cast(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& 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( reinterpret_cast(obheader) + desc->offset))) diff --git a/projects/hip/api/hip/hip_stream.cpp b/projects/hip/api/hip/hip_stream.cpp index 6d0da6adfc..6c2c29fb51 100644 --- a/projects/hip/api/hip/hip_stream.cpp +++ b/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(stream))->asHostQueue(); + amd::HostQueue* hostQueue; + + if (stream == nullptr) { + hostQueue = hip::getNullStream(); + } else { + hostQueue = as_amd(reinterpret_cast(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(stream))->release(); return hipSuccess; diff --git a/projects/hip/api/hip/hip_texture.cpp b/projects/hip/api/hip/hip_texture.cpp index 220332a57b..e619065f02 100644 --- a/projects/hip/api/hip/hip_texture.cpp +++ b/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) {