AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-310292 - checking hip initialization

Browse Source 
Change-Id: I83c9a0d511699c150abe61c3053cc4de2f7f1205
This commit is contained in:
pghafari
2021-11-18 16:19:55 -05:00
committed by Payam Ghafari
parent d3c829aec2
commit 24e984fd57
4 changed files with 46 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hip_context.cpp
+14 -12
View File
@@ -32,29 +32,30 @@ namespace hip {
thread_local Device* g_device = nullptr;
thread_local std::stack<Device*> 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<amd::Device*>& devices = amd::Device::getDevices(CL_DEVICE_TYPE_GPU, false);
for (unsigned int i=0; i<devices.size(); i++) {
const std::vector<amd::Device*> 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() {
hipamd/src/hip_internal.hpp
+27 -10
View File
@@ -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();
hipamd/src/hip_platform.cpp
+4 -4
View File
@@ -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<hip::Stream*>(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<void*>(function_address),
numBlocks, dimBlocks, kernarg, sharedMemBytes, stream);
hipamd/src/hiprtc_internal.hpp
+1 -1
View File
@@ -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; \