diff --git a/hipamd/src/hip_context.cpp b/hipamd/src/hip_context.cpp index f01b41575f..d28f9c350f 100644 --- a/hipamd/src/hip_context.cpp +++ b/hipamd/src/hip_context.cpp @@ -32,29 +32,30 @@ namespace hip { thread_local Device* g_device = nullptr; thread_local std::stack g_ctxtStack; thread_local hipError_t g_lastError = hipSuccess; -std::once_flag g_ihipInitialized; Device* host_device = nullptr; -void init() { - if (!amd::Runtime::initialized()) { - amd::IS_HIP = true; - GPU_NUM_MEM_DEPENDENCY = 0; +//init() is only to be called from the HIP_INIT macro only once +bool init() { + amd::IS_HIP = true; + GPU_NUM_MEM_DEPENDENCY = 0; #if DISABLE_DIRECT_DISPATCH - constexpr bool kDirectDispatch = false; + constexpr bool kDirectDispatch = false; #else - constexpr bool kDirectDispatch = IS_LINUX; + constexpr bool kDirectDispatch = IS_LINUX; #endif - AMD_DIRECT_DISPATCH = flagIsDefault(AMD_DIRECT_DISPATCH) ? kDirectDispatch : AMD_DIRECT_DISPATCH; - amd::Runtime::init(); - LogPrintfInfo("Direct Dispatch: %d", AMD_DIRECT_DISPATCH); + AMD_DIRECT_DISPATCH = flagIsDefault(AMD_DIRECT_DISPATCH) ? kDirectDispatch : AMD_DIRECT_DISPATCH; + if (!amd::Runtime::init()) { + return false; } + LogPrintfInfo("Direct Dispatch: %d", AMD_DIRECT_DISPATCH); + const std::vector& devices = amd::Device::getDevices(CL_DEVICE_TYPE_GPU, false); for (unsigned int i=0; i device(1, devices[i]); amd::Context* context = new amd::Context(device, amd::Context::Info()); - if (!context) return; + if (!context) return false; // Enable active wait on the device by default devices[i]->SetActiveWait(true); @@ -67,7 +68,7 @@ void init() { } amd::Context* hContext = new amd::Context(devices, amd::Context::Info()); - if (!hContext) return; + if (!hContext) return false; if (CL_SUCCESS != hContext->create(nullptr)) { hContext->release(); @@ -75,6 +76,7 @@ void init() { host_device = new Device(hContext, -1); PlatformState::instance().init(); + return true; } Device* getCurrentDevice() { diff --git a/hipamd/src/hip_internal.hpp b/hipamd/src/hip_internal.hpp index e8819a7f7d..629e2e6868 100644 --- a/hipamd/src/hip_internal.hpp +++ b/hipamd/src/hip_internal.hpp @@ -72,16 +72,34 @@ typedef struct ihipIpcEventHandle_st { inline int getpid() { return _getpid(); } #endif -#define HIP_INIT() \ - std::call_once(hip::g_ihipInitialized, hip::init); \ - if (hip::g_device == nullptr && g_devices.size() > 0) { \ - hip::g_device = g_devices[0]; \ - amd::Os::setPreferredNumaNode(g_devices[0]->devices()[0]->getPreferredNumaNode()); \ +static amd::Monitor g_hipInitlock{"hipInit lock"}; +#define HIP_INIT() {\ + amd::ScopedLock lock(g_hipInitlock); \ + if (!amd::Runtime::initialized()) { \ + if (!hip::init()) { \ + HIP_RETURN(hipErrorInvalidDevice); \ + } \ + } \ + if (hip::g_device == nullptr && g_devices.size() > 0) { \ + hip::g_device = g_devices[0]; \ + amd::Os::setPreferredNumaNode(g_devices[0]->devices()[0]->getPreferredNumaNode()); \ + } \ } -#define HIP_API_PRINT(...) \ - uint64_t startTimeUs=0 ; HIPPrintDuration(amd::LOG_INFO, amd::LOG_API, \ - &startTimeUs, "%s%s ( %s )%s", KGRN, \ +#define HIP_INIT_VOID() {\ + amd::ScopedLock lock(g_hipInitlock); \ + if (!amd::Runtime::initialized()) { \ + if (hip::init()) {} \ + } \ + if (hip::g_device == nullptr && g_devices.size() > 0) { \ + hip::g_device = g_devices[0]; \ + amd::Os::setPreferredNumaNode(g_devices[0]->devices()[0]->getPreferredNumaNode()); \ + } \ + } + + +#define HIP_API_PRINT(...) \ + uint64_t startTimeUs=0 ; HIPPrintDuration(amd::LOG_INFO, amd::LOG_API, &startTimeUs, "%s%s ( %s )%s", KGRN, \ __func__, ToString( __VA_ARGS__ ).c_str(),KNRM); #define HIP_ERROR_PRINT(err, ...) \ @@ -342,14 +360,13 @@ namespace hip { amd::HostQueue* NullStream(bool skip_alloc = false); }; - extern std::once_flag g_ihipInitialized; /// Current thread's device extern thread_local Device* g_device; extern thread_local hipError_t g_lastError; /// Device representing the host - for pinned memory extern Device* host_device; - extern void init(); + extern bool init(); extern Device* getCurrentDevice(); diff --git a/hipamd/src/hip_platform.cpp b/hipamd/src/hip_platform.cpp index fc7fe37d1d..12e9aa21d0 100644 --- a/hipamd/src/hip_platform.cpp +++ b/hipamd/src/hip_platform.cpp @@ -104,7 +104,7 @@ extern "C" void __hipRegisterFunction( PlatformState::instance().registerStatFunction(hostFunction, func); if (!enable_deferred_loading) { - HIP_INIT(); + HIP_INIT_VOID(); hipFunction_t hfunc = nullptr; hipError_t hip_error = hipSuccess; for (size_t dev_idx = 0; dev_idx < g_devices.size(); ++dev_idx) { @@ -149,7 +149,7 @@ extern "C" void __hipRegisterManagedVar(void *hipModule, // Pointer to hip mod const char *name, // Name of the variable in code object size_t size, unsigned align) { - HIP_INIT(); + HIP_INIT_VOID(); hipError_t status = ihipMallocManaged(pointer, size, align); if( status == hipSuccess) { amd::HostQueue* queue = hip::getNullStream(); @@ -597,7 +597,7 @@ void hipLaunchKernelGGLImpl( hipStream_t stream, void** kernarg) { - HIP_INIT(); + HIP_INIT_VOID(); hip::Stream* s = reinterpret_cast(stream); int deviceId = (s != nullptr)? s->DeviceId() : ihipGetDevice(); @@ -625,7 +625,7 @@ void hipLaunchCooperativeKernelGGLImpl( hipStream_t stream, void** kernarg) { - HIP_INIT(); + HIP_INIT_VOID(); hipLaunchCooperativeKernel(reinterpret_cast(function_address), numBlocks, dimBlocks, kernarg, sharedMemBytes, stream); diff --git a/hipamd/src/hiprtc_internal.hpp b/hipamd/src/hiprtc_internal.hpp index ae1c31c2d6..55fbff8200 100644 --- a/hipamd/src/hiprtc_internal.hpp +++ b/hipamd/src/hiprtc_internal.hpp @@ -53,7 +53,7 @@ extern "C" char * __cxa_demangle(const char *mangled_name, char *output_buffer, if (!VDI_CHECK_THREAD(thread)) { \ HIPRTC_RETURN(HIPRTC_ERROR_INTERNAL_ERROR); \ } \ - HIP_INIT(); + HIP_INIT_VOID(); #define HIPRTC_RETURN(ret) \ hiprtc::g_lastRtcError = ret; \