diff --git a/runtime/hsa-runtime/core/common/shared.cpp b/runtime/hsa-runtime/core/common/shared.cpp index eae6886bd2..2adfda93bd 100644 --- a/runtime/hsa-runtime/core/common/shared.cpp +++ b/runtime/hsa-runtime/core/common/shared.cpp @@ -42,7 +42,9 @@ #include "core/common/shared.h" +namespace rocr { namespace core { std::function BaseShared::allocate_ = nullptr; std::function BaseShared::free_ = nullptr; -} +} // namespace core +} // namespace rocr diff --git a/runtime/hsa-runtime/core/common/shared.h b/runtime/hsa-runtime/core/common/shared.h index 6802db666a..5b1298a4fd 100644 --- a/runtime/hsa-runtime/core/common/shared.h +++ b/runtime/hsa-runtime/core/common/shared.h @@ -50,6 +50,7 @@ #include "core/util/utils.h" +namespace rocr { namespace core { /// @brief Base class encapsulating the allocator and deallocator for /// shared shared object. As used this will allocate GPU visible host @@ -235,4 +236,5 @@ template class SharedArray final : private BaseShared }; } // namespace core +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/agent.h b/runtime/hsa-runtime/core/inc/agent.h index a78851d989..c76064ac00 100644 --- a/runtime/hsa-runtime/core/inc/agent.h +++ b/runtime/hsa-runtime/core/inc/agent.h @@ -54,6 +54,7 @@ #include "core/inc/memory_region.h" #include "core/util/utils.h" +namespace rocr { namespace core { class Signal; @@ -298,5 +299,6 @@ class Agent : public Checked<0xF6BC25EB17E6F917> { DISALLOW_COPY_AND_ASSIGN(Agent); }; } // namespace core +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/amd_aql_queue.h b/runtime/hsa-runtime/core/inc/amd_aql_queue.h index f75017aaad..744bade58b 100644 --- a/runtime/hsa-runtime/core/inc/amd_aql_queue.h +++ b/runtime/hsa-runtime/core/inc/amd_aql_queue.h @@ -49,7 +49,8 @@ #include "core/inc/amd_gpu_agent.h" #include "core/util/locks.h" -namespace amd { +namespace rocr { +namespace AMD { /// @brief Encapsulates HW Aql Command Processor functionality. It /// provide the interface for things such as Doorbell register, read, /// write pointers and a buffer. @@ -282,5 +283,6 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo }; } // namespace amd +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/amd_blit_kernel.h b/runtime/hsa-runtime/core/inc/amd_blit_kernel.h index 819117ae2b..f261ffc131 100644 --- a/runtime/hsa-runtime/core/inc/amd_blit_kernel.h +++ b/runtime/hsa-runtime/core/inc/amd_blit_kernel.h @@ -49,7 +49,8 @@ #include "core/inc/blit.h" -namespace amd { +namespace rocr { +namespace AMD { class BlitKernel : public core::Blit { public: explicit BlitKernel(core::Queue* queue); @@ -190,5 +191,6 @@ class BlitKernel : public core::Blit { int num_cus_; }; } // namespace amd +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/amd_blit_sdma.h b/runtime/hsa-runtime/core/inc/amd_blit_sdma.h index 2471435689..21dd55273f 100644 --- a/runtime/hsa-runtime/core/inc/amd_blit_sdma.h +++ b/runtime/hsa-runtime/core/inc/amd_blit_sdma.h @@ -55,7 +55,8 @@ #include "core/inc/signal.h" #include "core/util/utils.h" -namespace amd { +namespace rocr { +namespace AMD { class BlitSdmaBase : public core::Blit { public: @@ -290,5 +291,6 @@ typedef BlitSdma BlitSdmaV4; typedef BlitSdma BlitSdmaV5; } // namespace amd +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/amd_cpu_agent.h b/runtime/hsa-runtime/core/inc/amd_cpu_agent.h index 902f08545c..9bbed49c9e 100644 --- a/runtime/hsa-runtime/core/inc/amd_cpu_agent.h +++ b/runtime/hsa-runtime/core/inc/amd_cpu_agent.h @@ -54,7 +54,8 @@ #include "core/inc/queue.h" #include "core/inc/cache.h" -namespace amd { +namespace rocr { +namespace AMD { // @brief Class to represent a CPU device. class CpuAgent : public core::Agent { public: @@ -161,5 +162,6 @@ class CpuAgent : public core::Agent { }; } // namespace amd +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/amd_elf_image.hpp b/runtime/hsa-runtime/core/inc/amd_elf_image.hpp index 2a5daa5380..34bf943224 100644 --- a/runtime/hsa-runtime/core/inc/amd_elf_image.hpp +++ b/runtime/hsa-runtime/core/inc/amd_elf_image.hpp @@ -87,8 +87,9 @@ mes of its #include #include +namespace rocr { namespace amd { - namespace elf { +namespace elf { class Symbol; class SymbolTable; class Section; @@ -258,7 +259,8 @@ namespace amd { std::string GetNoteString(uint32_t s_size, const char* s); - } -} +} // namespace elf +} // namespace amd +} // namespace rocr #endif // AMD_ELF_IMAGE_HPP_ diff --git a/runtime/hsa-runtime/core/inc/amd_filter_device.h b/runtime/hsa-runtime/core/inc/amd_filter_device.h index ff511007a1..bfae8526ab 100644 --- a/runtime/hsa-runtime/core/inc/amd_filter_device.h +++ b/runtime/hsa-runtime/core/inc/amd_filter_device.h @@ -52,7 +52,8 @@ #include "hsakmt.h" -namespace amd { +namespace rocr { +namespace AMD { // ROCr allows users to filter and reorder various Gpu devices that are // present on ROCm system. This ability is made available via environment @@ -209,5 +210,6 @@ class RvdFilter { }; // End of class RvdFilter } // namespace amd +} // namespace rocr #endif // header guard - HSA_RUNTIME_CORE_INC_AMD_FILTER_DEVICE_H_ diff --git a/runtime/hsa-runtime/core/inc/amd_gpu_agent.h b/runtime/hsa-runtime/core/inc/amd_gpu_agent.h index f8e06c26dd..91c1119c95 100644 --- a/runtime/hsa-runtime/core/inc/amd_gpu_agent.h +++ b/runtime/hsa-runtime/core/inc/amd_gpu_agent.h @@ -59,7 +59,8 @@ #include "core/util/locks.h" #include "core/util/lazy_ptr.h" -namespace amd { +namespace rocr { +namespace AMD { class MemoryRegion; // @brief Contains scratch memory information. @@ -80,7 +81,7 @@ class GpuAgentInt : public core::Agent { public: // @brief Constructor GpuAgentInt(uint32_t node_id) - : core::Agent(node_id, core::Agent::DeviceType::kAmdGpuDevice) {} + : core::Agent(node_id,core::Agent::DeviceType::kAmdGpuDevice) {} // @brief Ensure blits are ready (performance hint). virtual void PreloadBlits() {} @@ -261,25 +262,25 @@ class GpuAgent : public GpuAgentInt { // @brief Decrement GWS ref count. void GWSRelease(); - // @brief Override from amd::GpuAgentInt. + // @brief Override from AMD::GpuAgentInt. void AcquireQueueScratch(ScratchInfo& scratch) override; - // @brief Override from amd::GpuAgentInt. + // @brief Override from AMD::GpuAgentInt. void ReleaseQueueScratch(ScratchInfo& scratch) override; - // @brief Override from amd::GpuAgentInt. + // @brief Override from AMD::GpuAgentInt. void TranslateTime(core::Signal* signal, hsa_amd_profiling_dispatch_time_t& time) override; - // @brief Override from amd::GpuAgentInt. + // @brief Override from AMD::GpuAgentInt. void TranslateTime(core::Signal* signal, hsa_amd_profiling_async_copy_time_t& time) override; - // @brief Override from amd::GpuAgentInt. + // @brief Override from AMD::GpuAgentInt. uint64_t TranslateTime(uint64_t tick) override; - // @brief Override from amd::GpuAgentInt. + // @brief Override from AMD::GpuAgentInt. void InvalidateCodeCaches() override; - // @brief Override from amd::GpuAgentInt. + // @brief Override from AMD::GpuAgentInt. bool current_coherency_type(hsa_amd_coherency_type_t type) override; hsa_amd_coherency_type_t current_coherency_type() const override { @@ -314,18 +315,18 @@ class GpuAgent : public GpuAgentInt { // @brief Override from core::Agent. const core::Isa* isa() const override { return isa_; } - // @brief Override from amd::GpuAgentInt. + // @brief Override from AMD::GpuAgentInt. __forceinline bool is_kv_device() const override { return is_kv_device_; } - // @brief Override from amd::GpuAgentInt. + // @brief Override from AMD::GpuAgentInt. __forceinline hsa_profile_t profile() const override { return profile_; } - // @brief Override from amd::GpuAgentInt. + // @brief Override from AMD::GpuAgentInt. __forceinline uint32_t memory_bus_width() const override { return memory_bus_width_; } - // @brief Override from amd::GpuAgentInt. + // @brief Override from AMD::GpuAgentInt. __forceinline uint32_t memory_max_frequency() const override { return memory_max_frequency_; } @@ -518,6 +519,7 @@ class GpuAgent : public GpuAgentInt { DISALLOW_COPY_AND_ASSIGN(GpuAgent); }; -} // namespace +} // namespace amd +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/amd_gpu_shaders.h b/runtime/hsa-runtime/core/inc/amd_gpu_shaders.h index af0bcf385f..2f15c47ca8 100644 --- a/runtime/hsa-runtime/core/inc/amd_gpu_shaders.h +++ b/runtime/hsa-runtime/core/inc/amd_gpu_shaders.h @@ -43,7 +43,8 @@ #ifndef HSA_RUNTIME_CORE_INC_AMD_GPU_SHADERS_H_ #define HSA_RUNTIME_CORE_INC_AMD_GPU_SHADERS_H_ -namespace amd { +namespace rocr { +namespace AMD { static const unsigned int kCodeCopyAligned7[] = { 0xC0820100, 0xC0840104, 0xC0860108, 0xC088010C, 0xC08A0110, 0xC00C0114, @@ -522,5 +523,6 @@ static const unsigned int kCodeTrapHandler10[] = { }; } // namespace amd +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/amd_hsa_code.hpp b/runtime/hsa-runtime/core/inc/amd_hsa_code.hpp index 637698f0d2..3bd4b7acac 100644 --- a/runtime/hsa-runtime/core/inc/amd_hsa_code.hpp +++ b/runtime/hsa-runtime/core/inc/amd_hsa_code.hpp @@ -92,6 +92,7 @@ mes of its #include #include +namespace rocr { namespace amd { namespace hsa { namespace common { @@ -137,7 +138,7 @@ class_type* ObjectAt(uint64_t address) return (class_type*)address; } -} +} // namespace common namespace code { @@ -423,8 +424,9 @@ namespace code { uint64_t SectionOffset() const override { return elfsym->value() - elfsym->section()->addr(); } uint64_t VAddr() const override { return elfsym->value(); } }; -} -} -} +} // namespace code +} // namespace hsa +} // namespace amd +} // namespace rocr #endif // AMD_HSA_CODE_HPP_ diff --git a/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp b/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp index 66c6a9388e..ef356f0a51 100644 --- a/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp +++ b/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp @@ -102,6 +102,7 @@ enum amd_loaded_segment_info_t { AMD_LOADED_SEGMENT_INFO_SIZE = 3 }; +namespace rocr { namespace amd { namespace hsa { namespace loader { @@ -511,5 +512,6 @@ private: } // namespace loader } // namespace hsa } // namespace amd +} // namespace rocr #endif // AMD_HSA_LOADER_HPP diff --git a/runtime/hsa-runtime/core/inc/amd_loader_context.hpp b/runtime/hsa-runtime/core/inc/amd_loader_context.hpp index 96d73fece1..c109476f35 100644 --- a/runtime/hsa-runtime/core/inc/amd_loader_context.hpp +++ b/runtime/hsa-runtime/core/inc/amd_loader_context.hpp @@ -45,11 +45,12 @@ #include "core/inc/amd_hsa_loader.hpp" +namespace rocr { namespace amd { -class LoaderContext final: public hsa::loader::Context { +class LoaderContext final: public amd::hsa::loader::Context { public: - LoaderContext(): hsa::loader::Context() {} + LoaderContext(): amd::hsa::loader::Context() {} ~LoaderContext() {} @@ -89,5 +90,6 @@ private: }; } // namespace amd +} // namespace rocr #endif // HSA_RUNTIME_CORE_INC_AMD_LOADER_CONTEXT_HPP diff --git a/runtime/hsa-runtime/core/inc/amd_memory_region.h b/runtime/hsa-runtime/core/inc/amd_memory_region.h index e0025026e5..f5342610fd 100644 --- a/runtime/hsa-runtime/core/inc/amd_memory_region.h +++ b/runtime/hsa-runtime/core/inc/amd_memory_region.h @@ -55,7 +55,8 @@ #include "inc/hsa_ext_amd.h" -namespace amd { +namespace rocr { +namespace AMD { class MemoryRegion : public core::MemoryRegion { public: /// @brief Convert this object into hsa_region_t. @@ -71,7 +72,7 @@ class MemoryRegion : public core::MemoryRegion { return region_handle; } - /// @brief Convert hsa_region_t into amd::MemoryRegion *. + /// @brief Convert hsa_region_t into AMD::MemoryRegion *. static __forceinline MemoryRegion* Convert(hsa_region_t region) { return reinterpret_cast(region.handle); } @@ -204,6 +205,7 @@ class MemoryRegion : public core::MemoryRegion { mutable SimpleHeap fragment_allocator_; }; -} // namespace +} // namespace amd +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/amd_topology.h b/runtime/hsa-runtime/core/inc/amd_topology.h index 27ff534d8c..a18e0c3a52 100644 --- a/runtime/hsa-runtime/core/inc/amd_topology.h +++ b/runtime/hsa-runtime/core/inc/amd_topology.h @@ -43,7 +43,8 @@ #ifndef HSA_RUNTIME_CORE_INC_AMD_TOPOLOGY_H_ #define HSA_RUNTIME_CORE_INC_AMD_TOPOLOGY_H_ -namespace amd { +namespace rocr { +namespace AMD { /// @brief Initializes the runtime. /// Should not be called directly, must be called only from Runtime::Acquire() bool Load(); @@ -51,6 +52,7 @@ bool Load(); /// @brief Shutdown/cleanup of runtime. /// Should not be called directly, must be called only from Runtime::Release() bool Unload(); -} // namespace +} // namespace amd +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/blit.h b/runtime/hsa-runtime/core/inc/blit.h index 2d7eaff0a9..ad663c36dc 100644 --- a/runtime/hsa-runtime/core/inc/blit.h +++ b/runtime/hsa-runtime/core/inc/blit.h @@ -47,6 +47,7 @@ #include "core/inc/agent.h" +namespace rocr { namespace core { class Blit { public: @@ -111,5 +112,6 @@ class Blit { virtual bool isSDMA() const { return false; } }; } // namespace core +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/cache.h b/runtime/hsa-runtime/core/inc/cache.h index 886ab58763..5316fa4f8b 100644 --- a/runtime/hsa-runtime/core/inc/cache.h +++ b/runtime/hsa-runtime/core/inc/cache.h @@ -49,6 +49,7 @@ #include #include +namespace rocr { namespace core { class Cache : public Checked<0x39A6C7AD3F135B06> { @@ -77,6 +78,8 @@ class Cache : public Checked<0x39A6C7AD3F135B06> { // Forbid copying and moving of this object DISALLOW_COPY_AND_ASSIGN(Cache); }; -} + +} // namespace core +} // namespace rocr #endif // HSA_RUNTIME_CORE_INC_CACHE_H diff --git a/runtime/hsa-runtime/core/inc/checked.h b/runtime/hsa-runtime/core/inc/checked.h index a7b57dcad5..93793bcc78 100644 --- a/runtime/hsa-runtime/core/inc/checked.h +++ b/runtime/hsa-runtime/core/inc/checked.h @@ -46,6 +46,7 @@ #include #include +namespace rocr { namespace core { /// @brief Compares type codes and pointers to check object validity. Used for cast validation. @@ -105,4 +106,6 @@ template class Checked { }; } // namespace core +} // namespace rocr + #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/default_signal.h b/runtime/hsa-runtime/core/inc/default_signal.h index 3237f691b8..0dd1ba26b6 100644 --- a/runtime/hsa-runtime/core/inc/default_signal.h +++ b/runtime/hsa-runtime/core/inc/default_signal.h @@ -49,6 +49,7 @@ #include "core/inc/signal.h" #include "core/util/utils.h" +namespace rocr { namespace core { /// @brief Operations for a simple pure memory based signal. @@ -185,4 +186,5 @@ class DefaultSignal : private LocalSignal, public BusyWaitSignal { }; } // namespace core +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/exceptions.h b/runtime/hsa-runtime/core/inc/exceptions.h index 4b328be8e6..a7acd6bf19 100644 --- a/runtime/hsa-runtime/core/inc/exceptions.h +++ b/runtime/hsa-runtime/core/inc/exceptions.h @@ -48,6 +48,7 @@ #include "core/inc/hsa_internal.h" +namespace rocr { namespace AMD { /// @brief Exception type which carries an error code to return to the user. @@ -95,6 +96,7 @@ template class callback_t { func_t function; }; -} // namespace AMD +} // namespace amd +} // namespace rocr #endif // HSA_RUNTIME_CORE_INC_EXCEPTIONS_H diff --git a/runtime/hsa-runtime/core/inc/host_queue.h b/runtime/hsa-runtime/core/inc/host_queue.h index 2a61816889..3393ede73a 100644 --- a/runtime/hsa-runtime/core/inc/host_queue.h +++ b/runtime/hsa-runtime/core/inc/host_queue.h @@ -48,6 +48,7 @@ #include "core/inc/runtime.h" #include "core/inc/signal.h" +namespace rocr { namespace core { class HostQueue : public Queue { public: @@ -177,4 +178,5 @@ class HostQueue : public Queue { DISALLOW_COPY_AND_ASSIGN(HostQueue); }; } // namespace core +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/hsa_api_trace_int.h b/runtime/hsa-runtime/core/inc/hsa_api_trace_int.h index aed02f10e3..c82abdf965 100644 --- a/runtime/hsa-runtime/core/inc/hsa_api_trace_int.h +++ b/runtime/hsa-runtime/core/inc/hsa_api_trace_int.h @@ -46,6 +46,7 @@ #include "inc/hsa_api_trace.h" #include "core/inc/hsa_internal.h" +namespace rocr { namespace core { struct HsaApiTable { @@ -70,6 +71,7 @@ namespace core { extern HsaApiTable hsa_api_table_; extern HsaApiTable hsa_internal_api_table_; -} +} // namespace core +} // namespace rocr #endif diff --git a/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h b/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h index a02b4fd24a..2c5a40d2f1 100644 --- a/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h +++ b/runtime/hsa-runtime/core/inc/hsa_ext_amd_impl.h @@ -50,6 +50,7 @@ #include "inc/hsa_ext_amd.h" // Wrap internal implementation inside AMD namespace +namespace rocr { namespace AMD { // Mirrors Amd Extension Apis @@ -241,6 +242,7 @@ hsa_status_t HSA_API hsa_amd_register_deallocation_callback( hsa_status_t HSA_API hsa_amd_deregister_deallocation_callback( void* ptr, hsa_amd_deallocation_callback_t callback); -} // end of AMD namespace +} // namespace amd +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/hsa_ext_interface.h b/runtime/hsa-runtime/core/inc/hsa_ext_interface.h index 519c0eedae..20a51759f9 100644 --- a/runtime/hsa-runtime/core/inc/hsa_ext_interface.h +++ b/runtime/hsa-runtime/core/inc/hsa_ext_interface.h @@ -51,6 +51,7 @@ #include "core/util/os.h" #include "core/util/utils.h" +namespace rocr { namespace core { struct ImageExtTableInternal : public ImageExtTable { decltype(::hsa_amd_image_get_info_max_dim)* hsa_amd_image_get_info_max_dim_fn; @@ -95,7 +96,8 @@ class ExtensionEntryPoints { void UpdateAmdExtTable(decltype(::hsa_amd_image_create)* func_ptr); DISALLOW_COPY_AND_ASSIGN(ExtensionEntryPoints); -}; -} +}; +} // namespace core +} // namespace rocr #endif diff --git a/runtime/hsa-runtime/core/inc/hsa_internal.h b/runtime/hsa-runtime/core/inc/hsa_internal.h index 1c7a6d9068..55078aa32c 100644 --- a/runtime/hsa-runtime/core/inc/hsa_internal.h +++ b/runtime/hsa-runtime/core/inc/hsa_internal.h @@ -45,8 +45,8 @@ #include "inc/hsa.h" -namespace HSA -{ +namespace rocr { +namespace HSA { // Define core namespace interfaces - copy of function declarations in hsa.h hsa_status_t HSA_API hsa_init(); @@ -412,12 +412,14 @@ namespace HSA hsa_status_t HSA_API hsa_status_string( hsa_status_t status, const char **status_string); -} + +} // namespace HSA +} // namespace rocr #ifdef BUILDING_HSA_CORE_RUNTIME //This using declaration is deliberate! //We want unqualified name resolution to fail when building the runtime. This is a guard against accidental use of the intercept layer in the runtime. -//using namespace HSA; +//using namespace rocr::HSA; #endif #endif diff --git a/runtime/hsa-runtime/core/inc/intercept_queue.h b/runtime/hsa-runtime/core/inc/intercept_queue.h index c5fea439f8..4ab897fae9 100644 --- a/runtime/hsa-runtime/core/inc/intercept_queue.h +++ b/runtime/hsa-runtime/core/inc/intercept_queue.h @@ -54,6 +54,7 @@ #include "core/inc/exceptions.h" #include "core/util/locks.h" +namespace rocr { namespace core { // @brief Generic container to forward Queue interfaces into Queue* member. @@ -264,5 +265,6 @@ class InterceptQueue : public QueueProxy, private LocalSignal, public DoorbellSi }; } // namespace core +} // namespace rocr #endif // HSA_RUNTIME_CORE_INC_INTERCEPT_QUEUE_H_ diff --git a/runtime/hsa-runtime/core/inc/interrupt_signal.h b/runtime/hsa-runtime/core/inc/interrupt_signal.h index b3e6138daa..a47f853ce2 100644 --- a/runtime/hsa-runtime/core/inc/interrupt_signal.h +++ b/runtime/hsa-runtime/core/inc/interrupt_signal.h @@ -53,6 +53,7 @@ #include "core/inc/signal.h" #include "core/util/utils.h" +namespace rocr { namespace core { /// @brief A Signal implementation using interrupts versus plain memory based. @@ -213,4 +214,5 @@ class InterruptSignal : private LocalSignal, public Signal { }; } // namespace core +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/ipc_signal.h b/runtime/hsa-runtime/core/inc/ipc_signal.h index 0ee739af5a..87858ef18f 100644 --- a/runtime/hsa-runtime/core/inc/ipc_signal.h +++ b/runtime/hsa-runtime/core/inc/ipc_signal.h @@ -50,6 +50,7 @@ #include "core/inc/default_signal.h" #include "core/util/locks.h" +namespace rocr { namespace core { /// @brief Container for ipc shared memory. @@ -108,5 +109,6 @@ class IPCSignal : private SharedMemorySignal, public BusyWaitSignal { }; } // namespace core +} // namespace rocr #endif // HSA_RUNTME_CORE_INC_IPC_SIGNAL_H_ diff --git a/runtime/hsa-runtime/core/inc/isa.h b/runtime/hsa-runtime/core/inc/isa.h index 4d57f7bb29..9386c24f65 100644 --- a/runtime/hsa-runtime/core/inc/isa.h +++ b/runtime/hsa-runtime/core/inc/isa.h @@ -50,6 +50,7 @@ #include #include "core/inc/amd_hsa_code.hpp" +namespace rocr { namespace core { /// @class Wavefront. @@ -235,5 +236,6 @@ class IsaRegistry final { }; // class IsaRegistry } // namespace core +} // namespace rocr #endif // HSA_RUNTIME_CORE_ISA_HPP_ diff --git a/runtime/hsa-runtime/core/inc/memory_region.h b/runtime/hsa-runtime/core/inc/memory_region.h index 671fb6c81c..4940833557 100644 --- a/runtime/hsa-runtime/core/inc/memory_region.h +++ b/runtime/hsa-runtime/core/inc/memory_region.h @@ -50,6 +50,7 @@ #include "core/inc/agent.h" #include "core/inc/checked.h" +namespace rocr { namespace core { class Agent; @@ -118,5 +119,6 @@ class MemoryRegion : public Checked<0x9C961F19EE175BB3> { core::Agent* owner_; }; } // namespace core +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/queue.h b/runtime/hsa-runtime/core/inc/queue.h index 1853df69ff..cb2612ffbe 100644 --- a/runtime/hsa-runtime/core/inc/queue.h +++ b/runtime/hsa-runtime/core/inc/queue.h @@ -56,6 +56,7 @@ #include "hsakmt.h" +namespace rocr { namespace core { struct AqlPacket { @@ -332,6 +333,7 @@ class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue { private: DISALLOW_COPY_AND_ASSIGN(Queue); }; -} +} // namespace core +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/runtime.h b/runtime/hsa-runtime/core/inc/runtime.h index 268b46bb3a..3e1112de13 100644 --- a/runtime/hsa-runtime/core/inc/runtime.h +++ b/runtime/hsa-runtime/core/inc/runtime.h @@ -77,9 +77,10 @@ #define HSA_PACKET_ALIGN_BYTES 64 //Avoids include -namespace amd { +namespace rocr { +namespace AMD { class MemoryRegion; -} +} // namespace amd namespace core { extern bool g_use_interrupt_wait; @@ -95,7 +96,7 @@ extern bool g_use_interrupt_wait; /// - maintain loader state. /// - monitor asynchronous event from agent. class Runtime { - friend class amd::MemoryRegion; + friend class AMD::MemoryRegion; public: /// @brief Structure to describe connectivity between agents. struct LinkInfo { @@ -542,4 +543,5 @@ class Runtime { }; } // namespace core +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/sdma_registers.h b/runtime/hsa-runtime/core/inc/sdma_registers.h index 22be75a061..254744f6a8 100644 --- a/runtime/hsa-runtime/core/inc/sdma_registers.h +++ b/runtime/hsa-runtime/core/inc/sdma_registers.h @@ -46,7 +46,8 @@ #include #include -namespace amd { +namespace rocr { +namespace AMD { // SDMA packet for VI device. // Reference: http://people.freedesktop.org/~agd5f/dma_packets.txt @@ -565,5 +566,6 @@ typedef struct SDMA_PKT_GCR_TAG { } SDMA_PKT_GCR; } // namespace amd +} // namespace rocr #endif // HSA_RUNTIME_CORE_INC_SDMA_REGISTERS_H_ diff --git a/runtime/hsa-runtime/core/inc/signal.h b/runtime/hsa-runtime/core/inc/signal.h index 656ac54886..43411c9f6b 100644 --- a/runtime/hsa-runtime/core/inc/signal.h +++ b/runtime/hsa-runtime/core/inc/signal.h @@ -75,6 +75,7 @@ template <> struct less { }; } +namespace rocr { namespace core { class Agent; class Signal; @@ -627,7 +628,8 @@ class SignalDeleter { public: void operator()(Signal* ptr) { ptr->DestroySignal(); } }; -using unique_signal_ptr = ::std::unique_ptr<::core::Signal, SignalDeleter>; +using unique_signal_ptr = ::std::unique_ptr; } // namespace core +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 7b131a0b65..76961b5597 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -68,7 +68,8 @@ #include "core/inc/hsa_ext_amd_impl.h" #include "core/inc/amd_gpu_pm4.h" -namespace amd { +namespace rocr { +namespace AMD { // Queue::amd_queue_ is cache-aligned for performance. const uint32_t kAmdQueueAlignBytes = 0x40; @@ -182,7 +183,7 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr auto& regions = agent->regions(); for (auto region : regions) { - const MemoryRegion* amdregion = static_cast(region); + const MemoryRegion* amdregion = static_cast(region); uint64_t base = amdregion->GetBaseAddress(); if (amdregion->IsLDS()) { @@ -1119,3 +1120,4 @@ void AqlQueue::InitScratchSRD() { return; } } // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp b/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp index 3bbf11f140..77ea4c7352 100644 --- a/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp @@ -50,7 +50,8 @@ #include "core/inc/hsa_internal.h" #include "core/util/utils.h" -namespace amd { +namespace rocr { +namespace AMD { static const uint16_t kInvalidPacketHeader = HSA_PACKET_TYPE_INVALID; static std::string kBlitKernelSource(R"( @@ -543,7 +544,7 @@ hsa_status_t BlitKernel::Initialize(const core::Agent& agent) { kernarg_async_mask_ = queue_->public_handle()->size - 1; // Obtain the number of compute units in the underlying agent. - const GpuAgent& gpuAgent = static_cast(agent); + const AMD::GpuAgent& gpuAgent = static_cast(agent); num_cus_ = gpuAgent.properties().NumFComputeCores / 4; // Assemble shaders to AQL code objects. @@ -554,7 +555,7 @@ hsa_status_t BlitKernel::Initialize(const core::Agent& agent) { for (auto kernel_name : kernel_names) { KernelCode& kernel = kernels_[kernel_name.first]; - gpuAgent.AssembleShader(kernel_name.second, GpuAgent::AssembleTarget::AQL, kernel.code_buf_, + gpuAgent.AssembleShader(kernel_name.second, AMD::GpuAgent::AssembleTarget::AQL, kernel.code_buf_, kernel.code_buf_size_); } @@ -568,7 +569,7 @@ hsa_status_t BlitKernel::Initialize(const core::Agent& agent) { hsa_status_t BlitKernel::Destroy(const core::Agent& agent) { std::lock_guard guard(lock_); - const GpuAgent& gpuAgent = static_cast(agent); + const AMD::GpuAgent& gpuAgent = static_cast(agent); for (auto kernel_pair : kernels_) { gpuAgent.ReleaseShader(kernel_pair.second.code_buf_, @@ -843,3 +844,4 @@ BlitKernel::KernelArgs* BlitKernel::ObtainAsyncKernelCopyArg() { } } // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp b/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp index 02f6485f08..be390bed52 100644 --- a/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp @@ -55,7 +55,8 @@ #include "core/inc/signal.h" #include "core/inc/interrupt_signal.h" -namespace amd { +namespace rocr { +namespace AMD { inline uint32_t ptrlow32(const void* p) { return static_cast(reinterpret_cast(p)); @@ -138,7 +139,7 @@ hsa_status_t BlitSdma: return HSA_STATUS_ERROR; } - agent_ = reinterpret_cast(&const_cast(agent)); + agent_ = reinterpret_cast(&const_cast(agent)); if (HSA_PROFILE_FULL == agent_->profile()) { assert(false && "Only support SDMA for dgpu currently"); @@ -953,3 +954,4 @@ template class BlitSdma; template class BlitSdma; } // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/amd_cpu_agent.cpp b/runtime/hsa-runtime/core/runtime/amd_cpu_agent.cpp index 70c54c18a3..576f66369c 100644 --- a/runtime/hsa-runtime/core/runtime/amd_cpu_agent.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_cpu_agent.cpp @@ -50,7 +50,8 @@ #include "inc/hsa_ext_image.h" -namespace amd { +namespace rocr { +namespace AMD { CpuAgent::CpuAgent(HSAuint32 node, const HsaNodeProperties& node_props) : core::Agent(node, kAmdCpuDevice), properties_(node_props) { InitRegionList(); @@ -383,3 +384,4 @@ hsa_status_t CpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type, } } // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/amd_filter_device.cpp b/runtime/hsa-runtime/core/runtime/amd_filter_device.cpp index 0fd0263053..51b38a619f 100644 --- a/runtime/hsa-runtime/core/runtime/amd_filter_device.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_filter_device.cpp @@ -60,7 +60,8 @@ #include "core/inc/amd_gpu_agent.h" #include "core/inc/amd_memory_region.h" -namespace amd { +namespace rocr { +namespace AMD { bool RvdFilter::FilterDevices() { return core::Runtime::runtime_singleton_->flag().filter_visible_gpus(); @@ -270,3 +271,4 @@ void RvdFilter::PrintRvdTokenList() { #endif } // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index 760b85f1d1..a5df502917 100644 --- a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -71,9 +71,12 @@ #define MAX_WAVE_SCRATCH 8387584 // See COMPUTE_TMPRING_SIZE.WAVESIZE #define MAX_NUM_DOORBELLS 0x400 -extern core::HsaApiTable hsa_internal_api_table_; +namespace rocr { +namespace core { +extern HsaApiTable hsa_internal_api_table_; +} // namespace core -namespace amd { +namespace AMD { GpuAgent::GpuAgent(HSAuint32 node, const HsaNodeProperties& node_props) : GpuAgentInt(node), properties_(node_props), @@ -458,8 +461,8 @@ hsa_status_t GpuAgent::VisitRegion( void* data) const { AMD::callback_t call(callback); for (const core::MemoryRegion* region : regions) { - const amd::MemoryRegion* amd_region = - reinterpret_cast(region); + const AMD::MemoryRegion* amd_region = + reinterpret_cast(region); // Only expose system, local, and LDS memory. if (amd_region->IsSystem() || amd_region->IsLocalMemory() || @@ -486,7 +489,7 @@ core::Queue* GpuAgent::CreateInterceptibleQueue() { } core::Blit* GpuAgent::CreateBlitSdma(bool use_xgmi) { - amd::BlitSdmaBase* sdma; + AMD::BlitSdmaBase* sdma; switch (isa_->GetMajorVersion()) { case 7: @@ -514,7 +517,7 @@ core::Blit* GpuAgent::CreateBlitSdma(bool use_xgmi) { } core::Blit* GpuAgent::CreateBlitKernel(core::Queue* queue) { - BlitKernel* kernl = new BlitKernel(queue); + AMD::BlitKernel* kernl = new AMD::BlitKernel(queue); if (kernl->Initialize(*this) != HSA_STATUS_SUCCESS) { kernl->Destroy(*this); @@ -1402,4 +1405,5 @@ lazy_ptr& GpuAgent::GetBlitObject(const core::Agent& dst_agent, return GetXgmiBlit(dst_agent); } -} // namespace +} // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp b/runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp index d963be703d..d081726cf6 100644 --- a/runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp @@ -76,6 +76,7 @@ std::string EncodePathname(const char *Pathname) { } // namespace +namespace rocr { namespace amd { namespace hsa { namespace loader { @@ -175,3 +176,4 @@ std::string CodeObjectReaderWrapper::GetUriFromMemory( } // namespace loader } // namespace hsa } // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/amd_loader_context.cpp b/runtime/hsa-runtime/core/runtime/amd_loader_context.cpp index 635ca9475f..abcca0d132 100644 --- a/runtime/hsa-runtime/core/runtime/amd_loader_context.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_loader_context.cpp @@ -62,11 +62,12 @@ #include #endif +namespace rocr { namespace { bool IsLocalRegion(const core::MemoryRegion *region) { - const amd::MemoryRegion *amd_region = (amd::MemoryRegion*)region; + const AMD::MemoryRegion *amd_region = (AMD::MemoryRegion*)region; if (nullptr == amd_region || !amd_region->IsLocalMemory()) { return false; } @@ -292,7 +293,7 @@ private: hsa_region_t RegionMemory::AgentLocal(hsa_agent_t agent) { hsa_region_t invalid_region; invalid_region.handle = 0; - amd::GpuAgent *amd_agent = (amd::GpuAgent*)core::Agent::Convert(agent); + AMD::GpuAgent *amd_agent = (AMD::GpuAgent*)core::Agent::Convert(agent); if (nullptr == amd_agent) { return invalid_region; } @@ -357,7 +358,7 @@ void RegionMemory::Free() bool RegionMemory::Freeze() { assert(this->Allocated() && nullptr != host_ptr_); - core::Agent* agent = reinterpret_cast( + core::Agent* agent = reinterpret_cast( core::MemoryRegion::Convert(region_))->owner(); if (agent != NULL && agent->device_type() == core::Agent::kAmdGpuDevice) { if (HSA_STATUS_SUCCESS != agent->DmaCopy(ptr_, host_ptr_, size_)) { @@ -475,14 +476,14 @@ void* LoaderContext::SegmentAlloc(amdgpu_hsa_elf_segment_t segment, RegionMemory::AgentLocal(agent)); break; case HSA_PROFILE_FULL: - mem = new (std::nothrow) MappedMemory(((GpuAgentInt*)core::Agent::Convert(agent))->is_kv_device()); + mem = new (std::nothrow) MappedMemory(((AMD::GpuAgentInt*)core::Agent::Convert(agent))->is_kv_device()); break; default: assert(false); } // Invalidate agent caches which may hold lines of the new allocation. - ((GpuAgentInt*)core::Agent::Convert(agent))->InvalidateCodeCaches(); + ((AMD::GpuAgentInt*)core::Agent::Convert(agent))->InvalidateCodeCaches(); break; } @@ -613,3 +614,4 @@ hsa_status_t LoaderContext::SamplerDestroy(hsa_agent_t agent, } } // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp b/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp index 9f94a0d131..43581e0923 100644 --- a/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp @@ -51,7 +51,8 @@ #include "core/util/utils.h" #include "core/inc/exceptions.h" -namespace amd { +namespace rocr { +namespace AMD { // Tracks aggregate size of system memory available on platform size_t MemoryRegion::max_sysmem_alloc_size_ = 0; @@ -571,7 +572,7 @@ hsa_status_t MemoryRegion::AllowAccess(uint32_t num_agents, assert(cpu_in_list); // This is a system region and only CPU agents in the whitelist. // Remove old mappings. - amd::MemoryRegion::MakeKfdMemoryUnresident(ptr); + AMD::MemoryRegion::MakeKfdMemoryUnresident(ptr); return HSA_STATUS_SUCCESS; } @@ -590,7 +591,7 @@ hsa_status_t MemoryRegion::AllowAccess(uint32_t num_agents, { ScopedAcquire lock(&core::Runtime::runtime_singleton_->memory_lock_); uint64_t alternate_va = 0; - if (!amd::MemoryRegion::MakeKfdMemoryResident( + if (!AMD::MemoryRegion::MakeKfdMemoryResident( whitelist_nodes.size(), &whitelist_nodes[0], ptr, size, &alternate_va, map_flag)) { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; @@ -670,7 +671,7 @@ hsa_status_t MemoryRegion::Lock(uint32_t num_agents, const hsa_agent_t* agents, return HSA_STATUS_SUCCESS; } - amd::MemoryRegion::DeregisterMemory(host_ptr); + AMD::MemoryRegion::DeregisterMemory(host_ptr); return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } @@ -706,11 +707,12 @@ void* MemoryRegion::BlockAllocator::alloc(size_t request_size, size_t& allocated hsa_status_t err = region_.Allocate( bsize, core::MemoryRegion::AllocateRestrict | core::MemoryRegion::AllocateDirect, &ret); if (err != HSA_STATUS_SUCCESS) - throw ::AMD::hsa_exception(err, "MemoryRegion::BlockAllocator::alloc failed."); + throw AMD::hsa_exception(err, "MemoryRegion::BlockAllocator::alloc failed."); assert(ret != nullptr && "Region returned nullptr on success."); allocated_size = block_size(); return ret; } -} // namespace +} // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/amd_topology.cpp b/runtime/hsa-runtime/core/runtime/amd_topology.cpp index 96e2161907..e6619fec9b 100644 --- a/runtime/hsa-runtime/core/runtime/amd_topology.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_topology.cpp @@ -62,7 +62,8 @@ #include "core/inc/amd_memory_region.h" #include "core/util/utils.h" -namespace amd { +namespace rocr { +namespace AMD { // Minimum acceptable KFD version numbers static const uint kKfdVersionMajor = 0; static const uint kKfdVersionMinor = 99; @@ -287,3 +288,4 @@ bool Unload() { return true; } } // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/cache.cpp b/runtime/hsa-runtime/core/runtime/cache.cpp index ab8d656dfe..a81df3bd57 100644 --- a/runtime/hsa-runtime/core/runtime/cache.cpp +++ b/runtime/hsa-runtime/core/runtime/cache.cpp @@ -43,6 +43,7 @@ #include "core/inc/cache.h" #include "assert.h" +namespace rocr { namespace core { hsa_status_t Cache::GetInfo(hsa_cache_info_t attribute, void* value) { @@ -64,4 +65,5 @@ hsa_status_t Cache::GetInfo(hsa_cache_info_t attribute, void* value) { } return HSA_STATUS_SUCCESS; } -} +} // namespace core +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/default_signal.cpp b/runtime/hsa-runtime/core/runtime/default_signal.cpp index 06766be897..46d3deb08a 100644 --- a/runtime/hsa-runtime/core/runtime/default_signal.cpp +++ b/runtime/hsa-runtime/core/runtime/default_signal.cpp @@ -43,6 +43,7 @@ #include "core/inc/default_signal.h" #include "core/util/timer.h" +namespace rocr { namespace core { int DefaultSignal::rtti_id_ = 0; @@ -275,3 +276,4 @@ hsa_signal_value_t BusyWaitSignal::CasAcqRel(hsa_signal_value_t expected, } } // namespace core +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/host_queue.cpp b/runtime/hsa-runtime/core/runtime/host_queue.cpp index 95c8f61f1d..2eab97b3b2 100644 --- a/runtime/hsa-runtime/core/runtime/host_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/host_queue.cpp @@ -45,6 +45,7 @@ #include "core/inc/runtime.h" #include "core/util/utils.h" +namespace rocr { namespace core { int HostQueue::rtti_id_ = 0; @@ -100,3 +101,4 @@ HostQueue::~HostQueue() { } } // namespace core +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/hsa.cpp b/runtime/hsa-runtime/core/runtime/hsa.cpp index 323cb5b75e..48425212b6 100644 --- a/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -60,6 +60,8 @@ #include "inc/hsa_ven_amd_aqlprofile.h" #include "core/inc/hsa_ext_amd_impl.h" +namespace rocr { + using namespace amd::hsa; template @@ -158,7 +160,7 @@ template static __forceinline T handleExceptionT() { abort(); return T(); } -} +} // namespace amd #define TRY try { #define CATCH } catch(...) { return AMD::handleException(); } @@ -1680,9 +1682,6 @@ hsa_status_t hsa_isa_compatible( //===--- Code Objects (deprecated) ----------------------------------------===// -using code::AmdHsaCode; -using code::AmdHsaCodeManager; - namespace { hsa_status_t IsCodeObjectAllocRegion( @@ -1734,7 +1733,7 @@ hsa_status_t FindCodeObjectAllocRegion( return HSA::hsa_iterate_agents(FindCodeObjectAllocRegionForAgent, data); } -AmdHsaCodeManager *GetCodeManager() { +amd::hsa::code::AmdHsaCodeManager *GetCodeManager() { return core::Runtime::runtime_singleton_->code_manager(); } @@ -1756,7 +1755,7 @@ hsa_status_t hsa_code_object_serialize( IS_BAD_PTR(serialized_code_object); IS_BAD_PTR(serialized_code_object_size); - AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); + amd::hsa::code::AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); if (!code) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } @@ -1917,7 +1916,7 @@ hsa_status_t hsa_code_object_get_info( IS_OPEN(); IS_BAD_PTR(value); - AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); + amd::hsa::code::AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); if (!code) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } @@ -1974,7 +1973,7 @@ hsa_status_t hsa_code_object_get_symbol( IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); - AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); + amd::hsa::code::AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); if (!code) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } @@ -1994,7 +1993,7 @@ hsa_status_t hsa_code_object_get_symbol_from_name( IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); - AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); + amd::hsa::code::AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); if (!code) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } @@ -2032,7 +2031,7 @@ hsa_status_t hsa_code_object_iterate_symbols( IS_OPEN(); IS_BAD_PTR(callback); - AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); + amd::hsa::code::AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); if (!code) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } @@ -2043,10 +2042,10 @@ hsa_status_t hsa_code_object_iterate_symbols( //===--- Executable -------------------------------------------------------===// -using common::Signed; -using loader::CodeObjectReaderWrapper; -using loader::Executable; -using loader::Loader; +using amd::hsa::common::Signed; +using amd::hsa::loader::Loader; +using amd::hsa::loader::Executable; +using amd::hsa::loader::CodeObjectReaderWrapper; namespace { @@ -2528,7 +2527,7 @@ hsa_status_t hsa_executable_iterate_program_symbols( IS_OPEN(); IS_BAD_PTR(callback); - Executable *exec = Executable::Object(executable); + amd::hsa::loader::Executable *exec = amd::hsa::loader::Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } @@ -2748,4 +2747,5 @@ hsa_status_t hsa_status_string( return HSA_STATUS_SUCCESS; } -} // end of namespace HSA +} // namespace HSA +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp b/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp index 93a5b042f9..7bd829d970 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp @@ -48,6 +48,7 @@ #include // Tools only APIs. +namespace rocr { namespace AMD { hsa_status_t hsa_amd_queue_intercept_register(hsa_queue_t* queue, hsa_amd_queue_intercept_handler callback, @@ -59,7 +60,7 @@ hsa_status_t hsa_amd_queue_intercept_create( hsa_status_t hsa_amd_runtime_queue_create_register(hsa_amd_runtime_queue_notifier callback, void* user_data); -} +} // namespace amd namespace core { @@ -396,4 +397,5 @@ class Init { Init() { hsa_table_interface_init(&hsa_api_table_.hsa_api); } }; static Init LinkAtLoad; -} +} // namespace core +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp b/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp index 0a9833d20f..95f9e82d0e 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp @@ -61,6 +61,8 @@ #include "core/inc/intercept_queue.h" #include "core/inc/exceptions.h" +namespace rocr { + template struct ValidityError; template <> @@ -79,7 +81,7 @@ struct ValidityError { }; template <> -struct ValidityError { +struct ValidityError { enum { value = HSA_STATUS_ERROR_INVALID_REGION }; }; @@ -172,8 +174,8 @@ hsa_status_t hsa_amd_coherency_get_type(hsa_agent_t agent_handle, hsa_amd_cohere return HSA_STATUS_ERROR_INVALID_AGENT; } - const amd::GpuAgentInt* gpu_agent = - static_cast(agent); + const AMD::GpuAgentInt* gpu_agent = + static_cast(agent); *type = gpu_agent->current_coherency_type(); @@ -199,7 +201,7 @@ hsa_status_t hsa_amd_coherency_set_type(hsa_agent_t agent_handle, return HSA_STATUS_ERROR_INVALID_AGENT; } - amd::GpuAgent* gpu_agent = static_cast(agent); + AMD::GpuAgent* gpu_agent = static_cast(agent); if (!gpu_agent->current_coherency_type(type)) { return HSA_STATUS_ERROR; @@ -289,7 +291,7 @@ hsa_status_t HSA_API hsa_amd_memory_async_copy_rect( IS_VALID(base_agent); if (base_agent->device_type() != core::Agent::DeviceType::kAmdGpuDevice) return HSA_STATUS_ERROR_INVALID_AGENT; - amd::GpuAgent* agent = static_cast(base_agent); + AMD::GpuAgent* agent = static_cast(base_agent); std::vector dep_signal_list(num_dep_signals); if (num_dep_signals > 0) { @@ -360,7 +362,7 @@ hsa_status_t hsa_amd_profiling_get_dispatch_time( return HSA_STATUS_ERROR_INVALID_AGENT; } - amd::GpuAgentInt* gpu_agent = static_cast(agent); + AMD::GpuAgentInt* gpu_agent = static_cast(agent); // Translate timestamp from GPU to system domain. gpu_agent->TranslateTime(signal, *time); @@ -388,7 +390,7 @@ hsa_status_t hsa_amd_profiling_get_async_copy_time( if (agent->device_type() == core::Agent::DeviceType::kAmdGpuDevice) { // Translate timestamp from GPU to system domain. - static_cast(agent)->TranslateTime(signal, *time); + static_cast(agent)->TranslateTime(signal, *time); return HSA_STATUS_SUCCESS; } @@ -415,7 +417,7 @@ hsa_status_t hsa_amd_profiling_convert_tick_to_system_domain(hsa_agent_t agent_h return HSA_STATUS_ERROR_INVALID_AGENT; } - amd::GpuAgentInt* gpu_agent = static_cast(agent); + AMD::GpuAgentInt* gpu_agent = static_cast(agent); *system_tick = gpu_agent->TranslateTime(agent_tick); @@ -550,7 +552,7 @@ hsa_status_t hsa_amd_memory_lock(void* host_ptr, size_t size, return HSA_STATUS_SUCCESS; } - const amd::MemoryRegion* system_region = static_cast( + const AMD::MemoryRegion* system_region = static_cast( core::Runtime::runtime_singleton_->system_regions_coarse()[0]); return system_region->Lock(num_agent, agents, host_ptr, size, agent_ptr); @@ -573,7 +575,7 @@ hsa_status_t hsa_amd_memory_lock_to_pool(void* host_ptr, size_t size, hsa_agent_ } hsa_region_t region = {pool.handle}; - const amd::MemoryRegion* mem_region = amd::MemoryRegion::Convert(region); + const AMD::MemoryRegion* mem_region = AMD::MemoryRegion::Convert(region); if (mem_region == nullptr) { return (hsa_status_t)HSA_STATUS_ERROR_INVALID_MEMORY_POOL; } @@ -588,8 +590,8 @@ hsa_status_t hsa_amd_memory_unlock(void* host_ptr) { TRY; IS_OPEN(); - const amd::MemoryRegion* system_region = - reinterpret_cast( + const AMD::MemoryRegion* system_region = + reinterpret_cast( core::Runtime::runtime_singleton_->system_regions_fine()[0]); return system_region->Unlock(host_ptr); @@ -603,7 +605,7 @@ hsa_status_t hsa_amd_memory_pool_get_info(hsa_amd_memory_pool_t memory_pool, IS_BAD_PTR(value); hsa_region_t region = {memory_pool.handle}; - const amd::MemoryRegion* mem_region = amd::MemoryRegion::Convert(region); + const AMD::MemoryRegion* mem_region = AMD::MemoryRegion::Convert(region); if (mem_region == NULL) { return (hsa_status_t)HSA_STATUS_ERROR_INVALID_MEMORY_POOL; } @@ -623,13 +625,13 @@ hsa_status_t hsa_amd_agent_iterate_memory_pools( IS_VALID(agent); if (agent->device_type() == core::Agent::kAmdCpuDevice) { - return reinterpret_cast(agent)->VisitRegion( + return reinterpret_cast(agent)->VisitRegion( false, reinterpret_cast(callback), data); } - return reinterpret_cast(agent)->VisitRegion( + return reinterpret_cast(agent)->VisitRegion( false, reinterpret_cast( callback), @@ -686,16 +688,16 @@ hsa_status_t hsa_amd_memory_pool_can_migrate(hsa_amd_memory_pool_t src_memory_po } hsa_region_t src_region_handle = {src_memory_pool.handle}; - const amd::MemoryRegion* src_mem_region = - amd::MemoryRegion::Convert(src_region_handle); + const AMD::MemoryRegion* src_mem_region = + AMD::MemoryRegion::Convert(src_region_handle); if (src_mem_region == NULL || !src_mem_region->IsValid()) { return static_cast(HSA_STATUS_ERROR_INVALID_MEMORY_POOL); } hsa_region_t dst_region_handle = {dst_memory_pool.handle}; - const amd::MemoryRegion* dst_mem_region = - amd::MemoryRegion::Convert(dst_region_handle); + const AMD::MemoryRegion* dst_mem_region = + AMD::MemoryRegion::Convert(dst_region_handle); if (dst_mem_region == NULL || !dst_mem_region->IsValid()) { return static_cast(HSA_STATUS_ERROR_INVALID_MEMORY_POOL); @@ -716,8 +718,8 @@ hsa_status_t hsa_amd_memory_migrate(const void* ptr, } hsa_region_t dst_region_handle = {memory_pool.handle}; - const amd::MemoryRegion* dst_mem_region = - amd::MemoryRegion::Convert(dst_region_handle); + const AMD::MemoryRegion* dst_mem_region = + AMD::MemoryRegion::Convert(dst_region_handle); if (dst_mem_region == NULL || !dst_mem_region->IsValid()) { return static_cast(HSA_STATUS_ERROR_INVALID_MEMORY_POOL); @@ -741,8 +743,8 @@ hsa_status_t hsa_amd_agent_memory_pool_get_info( IS_VALID(agent); hsa_region_t region_handle = {memory_pool.handle}; - const amd::MemoryRegion* mem_region = - amd::MemoryRegion::Convert(region_handle); + const AMD::MemoryRegion* mem_region = + AMD::MemoryRegion::Convert(region_handle); if (mem_region == NULL || !mem_region->IsValid()) { return static_cast(HSA_STATUS_ERROR_INVALID_MEMORY_POOL); @@ -987,4 +989,5 @@ hsa_status_t hsa_amd_runtime_queue_create_register(hsa_amd_runtime_queue_notifie CATCH; } -} // end of AMD namespace +} // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp b/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp index c5a037698b..2931b2b54f 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp @@ -46,6 +46,7 @@ #include +namespace rocr { // Implementations for missing / unsupported extensions template static R hsa_ext_null(ARGS...) { return HSA_STATUS_ERROR_NOT_INITIALIZED; @@ -259,6 +260,7 @@ bool ExtensionEntryPoints::LoadFinalizer(std::string library_name) { } } // namespace core +} // namespace rocr //---------------------------------------------------------------------------// // Exported extension stub functions @@ -268,19 +270,19 @@ hsa_status_t hsa_ext_program_create( hsa_machine_model_t machine_model, hsa_profile_t profile, hsa_default_float_rounding_mode_t default_float_rounding_mode, const char* options, hsa_ext_program_t* program) { - return core::Runtime::runtime_singleton_->extensions_.finalizer_api + return rocr::core::Runtime::runtime_singleton_->extensions_.finalizer_api .hsa_ext_program_create_fn(machine_model, profile, default_float_rounding_mode, options, program); } hsa_status_t hsa_ext_program_destroy(hsa_ext_program_t program) { - return core::Runtime::runtime_singleton_->extensions_.finalizer_api + return rocr::core::Runtime::runtime_singleton_->extensions_.finalizer_api .hsa_ext_program_destroy_fn(program); } hsa_status_t hsa_ext_program_add_module(hsa_ext_program_t program, hsa_ext_module_t module) { - return core::Runtime::runtime_singleton_->extensions_.finalizer_api + return rocr::core::Runtime::runtime_singleton_->extensions_.finalizer_api .hsa_ext_program_add_module_fn(program, module); } @@ -289,14 +291,14 @@ hsa_status_t hsa_ext_program_iterate_modules( hsa_status_t (*callback)(hsa_ext_program_t program, hsa_ext_module_t module, void* data), void* data) { - return core::Runtime::runtime_singleton_->extensions_.finalizer_api + return rocr::core::Runtime::runtime_singleton_->extensions_.finalizer_api .hsa_ext_program_iterate_modules_fn(program, callback, data); } hsa_status_t hsa_ext_program_get_info(hsa_ext_program_t program, hsa_ext_program_info_t attribute, void* value) { - return core::Runtime::runtime_singleton_->extensions_.finalizer_api + return rocr::core::Runtime::runtime_singleton_->extensions_.finalizer_api .hsa_ext_program_get_info_fn(program, attribute, value); } @@ -304,7 +306,7 @@ hsa_status_t hsa_ext_program_finalize( hsa_ext_program_t program, hsa_isa_t isa, int32_t call_convention, hsa_ext_control_directives_t control_directives, const char* options, hsa_code_object_type_t code_object_type, hsa_code_object_t* code_object) { - return core::Runtime::runtime_singleton_->extensions_.finalizer_api + return rocr::core::Runtime::runtime_singleton_->extensions_.finalizer_api .hsa_ext_program_finalize_fn(program, isa, call_convention, control_directives, options, code_object_type, code_object); @@ -313,7 +315,7 @@ hsa_status_t hsa_ext_program_finalize( hsa_status_t hsa_ext_image_get_capability( hsa_agent_t agent, hsa_ext_image_geometry_t geometry, const hsa_ext_image_format_t* image_format, uint32_t* capability_mask) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_image_get_capability_fn(agent, geometry, image_format, capability_mask); } @@ -322,7 +324,7 @@ hsa_status_t hsa_ext_image_data_get_info( hsa_agent_t agent, const hsa_ext_image_descriptor_t* image_descriptor, hsa_access_permission_t access_permission, hsa_ext_image_data_info_t* image_data_info) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_image_data_get_info_fn(agent, image_descriptor, access_permission, image_data_info); } @@ -331,7 +333,7 @@ hsa_status_t hsa_ext_image_create( hsa_agent_t agent, const hsa_ext_image_descriptor_t* image_descriptor, const void* image_data, hsa_access_permission_t access_permission, hsa_ext_image_t* image) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_image_create_fn(agent, image_descriptor, image_data, access_permission, image); } @@ -340,7 +342,7 @@ hsa_status_t hsa_ext_image_import(hsa_agent_t agent, const void* src_memory, size_t src_row_pitch, size_t src_slice_pitch, hsa_ext_image_t dst_image, const hsa_ext_image_region_t* image_region) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_image_import_fn(agent, src_memory, src_row_pitch, src_slice_pitch, dst_image, image_region); } @@ -349,7 +351,7 @@ hsa_status_t hsa_ext_image_export(hsa_agent_t agent, hsa_ext_image_t src_image, void* dst_memory, size_t dst_row_pitch, size_t dst_slice_pitch, const hsa_ext_image_region_t* image_region) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_image_export_fn(agent, src_image, dst_memory, dst_row_pitch, dst_slice_pitch, image_region); } @@ -359,7 +361,7 @@ hsa_status_t hsa_ext_image_copy(hsa_agent_t agent, hsa_ext_image_t src_image, hsa_ext_image_t dst_image, const hsa_dim3_t* dst_offset, const hsa_dim3_t* range) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_image_copy_fn(agent, src_image, src_offset, dst_image, dst_offset, range); } @@ -367,25 +369,25 @@ hsa_status_t hsa_ext_image_copy(hsa_agent_t agent, hsa_ext_image_t src_image, hsa_status_t hsa_ext_image_clear(hsa_agent_t agent, hsa_ext_image_t image, const void* data, const hsa_ext_image_region_t* image_region) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_image_clear_fn(agent, image, data, image_region); } hsa_status_t hsa_ext_image_destroy(hsa_agent_t agent, hsa_ext_image_t image) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_image_destroy_fn(agent, image); } hsa_status_t hsa_ext_sampler_create( hsa_agent_t agent, const hsa_ext_sampler_descriptor_t* sampler_descriptor, hsa_ext_sampler_t* sampler) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_sampler_create_fn(agent, sampler_descriptor, sampler); } hsa_status_t hsa_ext_sampler_destroy(hsa_agent_t agent, hsa_ext_sampler_t sampler) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_sampler_destroy_fn(agent, sampler); } @@ -394,7 +396,7 @@ hsa_status_t hsa_ext_image_get_capability_with_layout( const hsa_ext_image_format_t* image_format, hsa_ext_image_data_layout_t image_data_layout, uint32_t* capability_mask) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_image_get_capability_with_layout_fn(agent, geometry, image_format, image_data_layout, capability_mask); } @@ -406,7 +408,7 @@ hsa_status_t hsa_ext_image_data_get_info_with_layout( size_t image_data_row_pitch, size_t image_data_slice_pitch, hsa_ext_image_data_info_t* image_data_info) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_image_data_get_info_with_layout_fn(agent, image_descriptor, access_permission, image_data_layout, image_data_row_pitch, image_data_slice_pitch, @@ -420,7 +422,7 @@ hsa_status_t hsa_ext_image_create_with_layout( size_t image_data_row_pitch, size_t image_data_slice_pitch, hsa_ext_image_t* image) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_ext_image_create_with_layout_fn(agent, image_descriptor, image_data, access_permission, image_data_layout, image_data_row_pitch, image_data_slice_pitch, @@ -435,6 +437,6 @@ hsa_status_t hsa_ext_image_create_with_layout( hsa_status_t hsa_amd_image_get_info_max_dim(hsa_agent_t component, hsa_agent_info_t attribute, void* value) { - return core::Runtime::runtime_singleton_->extensions_.image_api + return rocr::core::Runtime::runtime_singleton_->extensions_.image_api .hsa_amd_image_get_info_max_dim_fn(component, attribute, value); } diff --git a/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp b/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp index 0f603bd454..fafaf400eb 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp @@ -45,8 +45,10 @@ #include "core/inc/amd_hsa_loader.hpp" #include "core/inc/runtime.h" -using namespace amd::hsa; -using namespace core; +//TODO SPK: These functions are not exports. They must be wrapped under rocr. + +using namespace rocr::amd::hsa; +using namespace rocr::core; using loader::CodeObjectReaderWrapper; using loader::Executable; @@ -55,7 +57,7 @@ using loader::LoadedCodeObject; hsa_status_t hsa_ven_amd_loader_query_host_address( const void *device_address, const void **host_address) { - if (false == core::Runtime::runtime_singleton_->IsOpen()) { + if (false == Runtime::runtime_singleton_->IsOpen()) { return HSA_STATUS_ERROR_NOT_INITIALIZED; } if (nullptr == device_address) { @@ -78,7 +80,7 @@ hsa_status_t hsa_ven_amd_loader_query_host_address( hsa_status_t hsa_ven_amd_loader_query_segment_descriptors( hsa_ven_amd_loader_segment_descriptor_t *segment_descriptors, size_t *num_segment_descriptors) { - if (false == core::Runtime::runtime_singleton_->IsOpen()) { + if (false == Runtime::runtime_singleton_->IsOpen()) { return HSA_STATUS_ERROR_NOT_INITIALIZED; } @@ -90,7 +92,7 @@ hsa_status_t hsa_ven_amd_loader_query_executable( const void *device_address, hsa_executable_t *executable) { - if (false == core::Runtime::runtime_singleton_->IsOpen()) { + if (false == Runtime::runtime_singleton_->IsOpen()) { return HSA_STATUS_ERROR_NOT_INITIALIZED; } if ((nullptr == device_address) || (nullptr == executable)) { @@ -114,7 +116,7 @@ hsa_status_t hsa_ven_amd_loader_executable_iterate_loaded_code_objects( hsa_loaded_code_object_t loaded_code_object, void *data), void *data) { - if (false == core::Runtime::runtime_singleton_->IsOpen()) { + if (false == Runtime::runtime_singleton_->IsOpen()) { return HSA_STATUS_ERROR_NOT_INITIALIZED; } if (nullptr == callback) { @@ -133,7 +135,7 @@ hsa_status_t hsa_ven_amd_loader_loaded_code_object_get_info( hsa_loaded_code_object_t loaded_code_object, hsa_ven_amd_loader_loaded_code_object_info_t attribute, void *value) { - if (false == core::Runtime::runtime_singleton_->IsOpen()) { + if (false == Runtime::runtime_singleton_->IsOpen()) { return HSA_STATUS_ERROR_NOT_INITIALIZED; } if (nullptr == value) { @@ -228,7 +230,7 @@ hsa_ven_amd_loader_code_object_reader_create_from_file_with_offset_size( size_t offset, size_t size, hsa_code_object_reader_t *code_object_reader) { - if (false == core::Runtime::runtime_singleton_->IsOpen()) { + if (false == Runtime::runtime_singleton_->IsOpen()) { return HSA_STATUS_ERROR_NOT_INITIALIZED; } if (nullptr == code_object_reader) { diff --git a/runtime/hsa-runtime/core/runtime/intercept_queue.cpp b/runtime/hsa-runtime/core/runtime/intercept_queue.cpp index 47656521c0..3ad278920c 100644 --- a/runtime/hsa-runtime/core/runtime/intercept_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/intercept_queue.cpp @@ -43,6 +43,7 @@ #include "core/inc/intercept_queue.h" #include "core/util/utils.h" +namespace rocr { namespace core { struct InterceptFrame { @@ -239,3 +240,4 @@ void InterceptQueue::StoreRelaxed(hsa_signal_value_t value) { } } // namespace core +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp b/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp index 3a3c9152d1..0a9ff803c5 100644 --- a/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp +++ b/runtime/hsa-runtime/core/runtime/interrupt_signal.cpp @@ -45,6 +45,7 @@ #include "core/util/timer.h" #include "core/util/locks.h" +namespace rocr { namespace core { HsaEvent* InterruptSignal::EventPool::alloc() { @@ -386,3 +387,4 @@ hsa_signal_value_t InterruptSignal::CasAcqRel(hsa_signal_value_t expected, } } // namespace core +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/ipc_signal.cpp b/runtime/hsa-runtime/core/runtime/ipc_signal.cpp index faab99e890..485a5c915a 100644 --- a/runtime/hsa-runtime/core/runtime/ipc_signal.cpp +++ b/runtime/hsa-runtime/core/runtime/ipc_signal.cpp @@ -47,6 +47,7 @@ #include "core/inc/runtime.h" #include "core/inc/exceptions.h" +namespace rocr { namespace core { int IPCSignal::rtti_id_ = 0; @@ -92,3 +93,4 @@ Signal* IPCSignal::Attach(const hsa_amd_ipc_signal_t* ipc_signal_handle) { } } // namespace core +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/isa.cpp b/runtime/hsa-runtime/core/runtime/isa.cpp index bd9918a01e..68a3fb2085 100755 --- a/runtime/hsa-runtime/core/runtime/isa.cpp +++ b/runtime/hsa-runtime/core/runtime/isa.cpp @@ -46,6 +46,7 @@ #include #include +namespace rocr { namespace core { bool Wavefront::GetInfo( @@ -240,3 +241,4 @@ const IsaRegistry::IsaMap IsaRegistry::GetSupportedIsas() { } } // namespace core +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/queue.cpp b/runtime/hsa-runtime/core/runtime/queue.cpp index 63eb8f6eb8..6c8a661839 100644 --- a/runtime/hsa-runtime/core/runtime/queue.cpp +++ b/runtime/hsa-runtime/core/runtime/queue.cpp @@ -43,6 +43,7 @@ #include "core/inc/queue.h" #include "core/inc/runtime.h" +namespace rocr { namespace core { void Queue::DefaultErrorHandler(hsa_status_t status, hsa_queue_t* source, void* data) { @@ -53,4 +54,5 @@ void Queue::DefaultErrorHandler(hsa_status_t status, hsa_queue_t* source, void* } } -} +} // namespace core +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/runtime.cpp b/runtime/hsa-runtime/core/runtime/runtime.cpp index e60bfe238b..63b2d1660c 100644 --- a/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/runtime/hsa-runtime/core/runtime/runtime.cpp @@ -67,6 +67,7 @@ const char rocrbuildid[] __attribute__((used)) = "ROCR BUILD ID: " STRING(ROCR_BUILD_ID); +namespace rocr { namespace core { bool g_use_interrupt_wait = true; @@ -420,8 +421,8 @@ hsa_status_t Runtime::CopyMemory(void* dst, const void* src, size_t size) { // GPU-CPU // Must ensure that system memory is visible to the GPU during the copy. - const amd::MemoryRegion* system_region = - static_cast(system_regions_fine_[0]); + const AMD::MemoryRegion* system_region = + static_cast(system_regions_fine_[0]); void* gpuPtr = nullptr; const auto& locked_copy = [&](void*& ptr, core::Agent* locking_agent) { @@ -548,7 +549,7 @@ hsa_status_t Runtime::FillMemory(void* ptr, uint32_t value, size_t count) { hsa_status_t Runtime::AllowAccess(uint32_t num_agents, const hsa_agent_t* agents, const void* ptr) { - const amd::MemoryRegion* amd_region = NULL; + const AMD::MemoryRegion* amd_region = NULL; size_t alloc_size = 0; { @@ -560,7 +561,7 @@ hsa_status_t Runtime::AllowAccess(uint32_t num_agents, return HSA_STATUS_ERROR; } - amd_region = reinterpret_cast(it->second.region); + amd_region = reinterpret_cast(it->second.region); alloc_size = it->second.size; } @@ -1229,7 +1230,7 @@ void Runtime::PrintMemoryMapNear(void* ptr) { if (it == runtime_singleton_->allocation_map_.end()) break; std::string kind = "Non-HSA"; if (it->second.region != nullptr) { - const amd::MemoryRegion* region = static_cast(it->second.region); + const AMD::MemoryRegion* region = static_cast(it->second.region); if (region->IsSystem()) kind = "System"; else if (region->IsLocalMemory()) @@ -1279,7 +1280,7 @@ hsa_status_t Runtime::Load() { g_use_interrupt_wait = flag_.enable_interrupt(); - if (!amd::Load()) { + if (!AMD::Load()) { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } @@ -1301,7 +1302,7 @@ hsa_status_t Runtime::Load() { // Initialize per GPU scratch, blits, and trap handler for (core::Agent* agent : gpu_agents_) { hsa_status_t status = - reinterpret_cast(agent)->PostToolsInit(); + reinterpret_cast(agent)->PostToolsInit(); if (status != HSA_STATUS_SUCCESS) { return status; @@ -1341,7 +1342,7 @@ void Runtime::Unload() { CloseTools(); - amd::Unload(); + AMD::Unload(); } void Runtime::LoadExtensions() { @@ -1578,3 +1579,4 @@ void Runtime::InternalQueueCreateNotify(const hsa_queue_t* queue, hsa_agent_t ag } } // namespace core +} // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/signal.cpp b/runtime/hsa-runtime/core/runtime/signal.cpp index 5df921c020..d424e29a06 100644 --- a/runtime/hsa-runtime/core/runtime/signal.cpp +++ b/runtime/hsa-runtime/core/runtime/signal.cpp @@ -49,6 +49,7 @@ #include "core/util/timer.h" #include "core/inc/runtime.h" +namespace rocr { namespace core { KernelMutex Signal::ipcLock_; @@ -324,5 +325,6 @@ SignalGroup::SignalGroup(uint32_t num_signals, const hsa_signal_t* hsa_signals) } } // namespace core +} // namespace rocr #endif // header guard diff --git a/runtime/hsa-runtime/core/util/atomic_helpers.h b/runtime/hsa-runtime/core/util/atomic_helpers.h index b9bcb84c23..8844dc41dd 100644 --- a/runtime/hsa-runtime/core/util/atomic_helpers.h +++ b/runtime/hsa-runtime/core/util/atomic_helpers.h @@ -64,6 +64,7 @@ #endif #endif +namespace rocr { namespace atomic { static constexpr int c11ToBuiltInFlags(std::memory_order order) @@ -504,7 +505,8 @@ static __forceinline T PostFence(order); return ret; } -} +} // namespace atomic +} // namespace rocr #ifdef X64_ORDER_WC #undef X64_ORDER_WC diff --git a/runtime/hsa-runtime/core/util/flag.h b/runtime/hsa-runtime/core/util/flag.h index 4d9ac84293..e0df01c91b 100644 --- a/runtime/hsa-runtime/core/util/flag.h +++ b/runtime/hsa-runtime/core/util/flag.h @@ -50,6 +50,8 @@ #include "core/util/os.h" #include "core/util/utils.h" +namespace rocr { + class Flag { public: explicit Flag() { Refresh(); } @@ -188,4 +190,6 @@ class Flag { DISALLOW_COPY_AND_ASSIGN(Flag); }; +} // namespace rocr + #endif // header guard diff --git a/runtime/hsa-runtime/core/util/lazy_ptr.h b/runtime/hsa-runtime/core/util/lazy_ptr.h index c3bf59e5ba..e2a847b5c4 100644 --- a/runtime/hsa-runtime/core/util/lazy_ptr.h +++ b/runtime/hsa-runtime/core/util/lazy_ptr.h @@ -50,6 +50,8 @@ #include "core/util/locks.h" #include "core/util/utils.h" +namespace rocr { + /* * Wrapper for a std::unique_ptr that initializes its object at first use. */ @@ -148,4 +150,6 @@ template class lazy_ptr { }; +} // namespace rocr + #endif // HSA_RUNTIME_CORE_UTIL_LAZY_PTR_H_ diff --git a/runtime/hsa-runtime/core/util/lnx/os_linux.cpp b/runtime/hsa-runtime/core/util/lnx/os_linux.cpp index 29038ff0d7..e60036929a 100644 --- a/runtime/hsa-runtime/core/util/lnx/os_linux.cpp +++ b/runtime/hsa-runtime/core/util/lnx/os_linux.cpp @@ -60,6 +60,7 @@ #include #include +namespace rocr { namespace os { struct ThreadArgs { @@ -447,6 +448,7 @@ uint64_t AccurateClockFrequency() { if (invPeriod == 0.0) invPeriod = 1.0 / double(time.tv_nsec); return 1000000000ull / uint64_t(time.tv_nsec); } -} +} // namespace os +} // namespace rocr #endif diff --git a/runtime/hsa-runtime/core/util/locks.h b/runtime/hsa-runtime/core/util/locks.h index 24dd718a4a..49ad6906fe 100644 --- a/runtime/hsa-runtime/core/util/locks.h +++ b/runtime/hsa-runtime/core/util/locks.h @@ -48,6 +48,8 @@ #include "utils.h" #include "os.h" +namespace rocr { + /// @brief: A class behaves as a lock in a scope. When trying to enter into the /// critical section, creat a object of this class. After the control path goes /// out of the scope, it will release the lock automatically. @@ -142,4 +144,6 @@ class KernelEvent { DISALLOW_COPY_AND_ASSIGN(KernelEvent); }; +} // namespace rocr + #endif // HSA_RUNTIME_CORE_SUTIL_LOCKS_H_ diff --git a/runtime/hsa-runtime/core/util/os.h b/runtime/hsa-runtime/core/util/os.h index 7743dc9c87..ec9f54a422 100644 --- a/runtime/hsa-runtime/core/util/os.h +++ b/runtime/hsa-runtime/core/util/os.h @@ -48,6 +48,7 @@ #include #include "utils.h" +namespace rocr { namespace os { typedef void* LibHandle; typedef void* Mutex; @@ -222,6 +223,7 @@ uint64_t ReadAccurateClock(); /// seconds. This frequency does not change at runtime. /// @return returns the frequency uint64_t AccurateClockFrequency(); -} +} // namespace os +} // namespace rocr #endif // HSA_RUNTIME_CORE_UTIL_OS_H_ diff --git a/runtime/hsa-runtime/core/util/simple_heap.h b/runtime/hsa-runtime/core/util/simple_heap.h index 9da0db8382..a5a5c1a859 100644 --- a/runtime/hsa-runtime/core/util/simple_heap.h +++ b/runtime/hsa-runtime/core/util/simple_heap.h @@ -53,6 +53,8 @@ #include "core/util/utils.h" +namespace rocr { + template class SimpleHeap { private: struct Fragment_T { @@ -298,4 +300,6 @@ template class SimpleHeap { } }; +} // namespace rocr + #endif // HSA_RUNTME_CORE_UTIL_SIMPLE_HEAP_H_ diff --git a/runtime/hsa-runtime/core/util/small_heap.cpp b/runtime/hsa-runtime/core/util/small_heap.cpp index 9e69222e45..9fe5da5fab 100644 --- a/runtime/hsa-runtime/core/util/small_heap.cpp +++ b/runtime/hsa-runtime/core/util/small_heap.cpp @@ -42,6 +42,8 @@ #include "small_heap.h" +namespace rocr { + // Inserts node into freelist after place. // Assumes node will not be an end of the list (list has guard nodes). void SmallHeap::insertafter(SmallHeap::iterator_t place, SmallHeap::iterator_t node) { @@ -179,3 +181,5 @@ void* SmallHeap::alloc_high(size_t bytes) { // Can't service the request due to fragmentation return nullptr; } + +} // namespace rocr diff --git a/runtime/hsa-runtime/core/util/small_heap.h b/runtime/hsa-runtime/core/util/small_heap.h index f06e667e57..e1f5d7bdeb 100644 --- a/runtime/hsa-runtime/core/util/small_heap.h +++ b/runtime/hsa-runtime/core/util/small_heap.h @@ -51,7 +51,9 @@ #include #include "utils.h" - + +namespace rocr { + class SmallHeap { private: struct Node; @@ -124,4 +126,6 @@ class SmallHeap { void* high_split() const { return *high.begin(); } }; +} // namespace rocr + #endif diff --git a/runtime/hsa-runtime/core/util/timer.cpp b/runtime/hsa-runtime/core/util/timer.cpp index 338bbe0720..d9257eb32b 100644 --- a/runtime/hsa-runtime/core/util/timer.cpp +++ b/runtime/hsa-runtime/core/util/timer.cpp @@ -42,6 +42,7 @@ #include "core/util/timer.h" +namespace rocr { namespace timer { accurate_clock::init::init() { @@ -104,4 +105,5 @@ accurate_clock::init accurate_clock::accurate_clock_init; double fast_clock::period_ps; fast_clock::raw_frequency fast_clock::freq; fast_clock::init fast_clock::fast_clock_init; -} +} // namespace timer +} // namespace rocr diff --git a/runtime/hsa-runtime/core/util/timer.h b/runtime/hsa-runtime/core/util/timer.h index ac6224ca19..95ae24953a 100644 --- a/runtime/hsa-runtime/core/util/timer.h +++ b/runtime/hsa-runtime/core/util/timer.h @@ -49,6 +49,7 @@ #include #include +namespace rocr { namespace timer { // Needed to patch around a mixed arithmetic bug in MSVC's duration_cast as of @@ -166,6 +167,7 @@ class fast_clock { }; static init fast_clock_init; }; -} +} // namespace timer +} // namespace rocr #endif diff --git a/runtime/hsa-runtime/core/util/utils.h b/runtime/hsa-runtime/core/util/utils.h index 4a65e90eec..29478f4dd2 100644 --- a/runtime/hsa-runtime/core/util/utils.h +++ b/runtime/hsa-runtime/core/util/utils.h @@ -53,6 +53,8 @@ #include #include +namespace rocr { + typedef unsigned int uint; typedef uint64_t uint64; @@ -326,6 +328,8 @@ static __forceinline std::string& rtrim(std::string& s) { static __forceinline std::string& trim(std::string& s) { return ltrim(rtrim(s)); } +} // namespace rocr + template static __forceinline uint32_t BitSelect(T p) { static_assert(sizeof(T) <= sizeof(uintptr_t), "Type out of range."); diff --git a/runtime/hsa-runtime/core/util/win/os_win.cpp b/runtime/hsa-runtime/core/util/win/os_win.cpp index a55b3af49a..05aba21f59 100644 --- a/runtime/hsa-runtime/core/util/win/os_win.cpp +++ b/runtime/hsa-runtime/core/util/win/os_win.cpp @@ -58,6 +58,7 @@ #undef Yield #undef CreateMutex +namespace rocr { namespace os { static_assert(sizeof(LibHandle) == sizeof(HMODULE), @@ -224,6 +225,7 @@ uint64_t AccurateClockFrequency() { QueryPerformanceFrequency((LARGE_INTEGER*)&ret); return ret; } -} +} // namespace os +} // namespace rocr #endif diff --git a/runtime/hsa-runtime/libamdhsacode/amd_elf_image.cpp b/runtime/hsa-runtime/libamdhsacode/amd_elf_image.cpp index 4950b0b189..a6b9021eab 100644 --- a/runtime/hsa-runtime/libamdhsacode/amd_elf_image.cpp +++ b/runtime/hsa-runtime/libamdhsacode/amd_elf_image.cpp @@ -99,10 +99,11 @@ #define NOTE_RECORD_ALIGNMENT 4 -using amd::hsa::alignUp; +using rocr::amd::hsa::alignUp; +namespace rocr { namespace amd { - namespace elf { +namespace elf { class FileImage { public: @@ -1711,5 +1712,6 @@ namespace amd { } } - } -} +} // namespace elf +} // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/libamdhsacode/amd_hsa_code.cpp b/runtime/hsa-runtime/libamdhsacode/amd_hsa_code.cpp index 5b1304f560..edf4e04efc 100644 --- a/runtime/hsa-runtime/libamdhsacode/amd_hsa_code.cpp +++ b/runtime/hsa-runtime/libamdhsacode/amd_hsa_code.cpp @@ -61,6 +61,7 @@ #define _alloca alloca #endif +namespace rocr { namespace amd { namespace hsa { namespace code { @@ -1686,6 +1687,7 @@ namespace code { KernelSymbolV2::KernelSymbolV2(amd::elf::Symbol* elfsym_, const amd_kernel_code_t* akc) : KernelSymbol(elfsym_, akc) { } -} -} -} +} // namespace code +} // namespace hsa +} // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/libamdhsacode/amd_hsa_code_util.cpp b/runtime/hsa-runtime/libamdhsacode/amd_hsa_code_util.cpp index 1fa277cea3..150840dd46 100644 --- a/runtime/hsa-runtime/libamdhsacode/amd_hsa_code_util.cpp +++ b/runtime/hsa-runtime/libamdhsacode/amd_hsa_code_util.cpp @@ -77,6 +77,7 @@ std::ostream& attr2(std::ostream& out) } } // namespace anonymous +namespace rocr { namespace amd { namespace hsa { namespace common { @@ -118,7 +119,7 @@ bool IsAccessibleMemoryAddress(uint64_t address) #endif // _WIN32 || _WIN64 } -} +} // namespace common std::string HsaSymbolKindToString(hsa_symbol_kind_t kind) { @@ -1045,5 +1046,6 @@ std::string DumpFileName(const std::string& dir, const char* prefix, const char* } -} -} +} // namespace hsa +} // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/libamdhsacode/amd_hsa_code_util.hpp b/runtime/hsa-runtime/libamdhsacode/amd_hsa_code_util.hpp index 07cc61058c..dc0a863ed0 100644 --- a/runtime/hsa-runtime/libamdhsacode/amd_hsa_code_util.hpp +++ b/runtime/hsa-runtime/libamdhsacode/amd_hsa_code_util.hpp @@ -67,6 +67,7 @@ std::abort(); \ } \ +namespace rocr { namespace amd { namespace hsa { @@ -189,7 +190,8 @@ inline uint32_t alignUp(uint32_t num, uint32_t align) std::string DumpFileName(const std::string& dir, const char* prefix, const char* ext, unsigned n, unsigned i = 0); -} -} +} // namespace hsa +} // namespace amd +} // namespace rocr #endif // AMD_HSA_CODE_UTIL_HPP_ diff --git a/runtime/hsa-runtime/libamdhsacode/amd_hsa_locks.cpp b/runtime/hsa-runtime/libamdhsacode/amd_hsa_locks.cpp index 5658c38a14..30b22e226d 100644 --- a/runtime/hsa-runtime/libamdhsacode/amd_hsa_locks.cpp +++ b/runtime/hsa-runtime/libamdhsacode/amd_hsa_locks.cpp @@ -42,6 +42,7 @@ #include "amd_hsa_locks.hpp" +namespace rocr { namespace amd { namespace hsa { namespace common { @@ -92,3 +93,4 @@ void ReaderWriterLock::WriterUnlock() } // namespace common } // namespace hsa } // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/libamdhsacode/amd_hsa_locks.hpp b/runtime/hsa-runtime/libamdhsacode/amd_hsa_locks.hpp index adb39a239c..d7fbe4680a 100644 --- a/runtime/hsa-runtime/libamdhsacode/amd_hsa_locks.hpp +++ b/runtime/hsa-runtime/libamdhsacode/amd_hsa_locks.hpp @@ -47,6 +47,7 @@ #include #include +namespace rocr { namespace amd { namespace hsa { namespace common { @@ -123,5 +124,6 @@ private: } // namespace common } // namespace hsa } // namespace amd +} // namespace rocr #endif // AMD_HSA_LOCKS_HPP diff --git a/runtime/hsa-runtime/libamdhsacode/amd_options.cpp b/runtime/hsa-runtime/libamdhsacode/amd_options.cpp index fc163d7d31..93b6bee139 100644 --- a/runtime/hsa-runtime/libamdhsacode/amd_options.cpp +++ b/runtime/hsa-runtime/libamdhsacode/amd_options.cpp @@ -54,6 +54,7 @@ #include +namespace rocr { namespace amd { namespace options { @@ -379,3 +380,4 @@ void OptionParser::Reset() { } // namespace options } // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/libamdhsacode/amd_options.hpp b/runtime/hsa-runtime/libamdhsacode/amd_options.hpp index 87fce99202..753a45d575 100644 --- a/runtime/hsa-runtime/libamdhsacode/amd_options.hpp +++ b/runtime/hsa-runtime/libamdhsacode/amd_options.hpp @@ -55,6 +55,7 @@ #include #include +namespace rocr { namespace amd { namespace options { @@ -472,5 +473,6 @@ private: } // namespace options } // namespace amd +} // namespace rocr #endif // AMD_OPTIONS_HPP diff --git a/runtime/hsa-runtime/loader/AMDHSAKernelDescriptor.h b/runtime/hsa-runtime/loader/AMDHSAKernelDescriptor.h index f6ced01ef8..4f64d2482c 100644 --- a/runtime/hsa-runtime/loader/AMDHSAKernelDescriptor.h +++ b/runtime/hsa-runtime/loader/AMDHSAKernelDescriptor.h @@ -72,6 +72,7 @@ DST |= ((VAL << MSK ## _SHIFT) & MSK) #endif // AMDHSA_BITS_SET +namespace rocr { namespace llvm { namespace amdhsa { @@ -208,5 +209,6 @@ static_assert( } // end namespace amdhsa } // end namespace llvm +} // end namespace rocr #endif // LLVM_SUPPORT_AMDHSAKERNELDESCRIPTOR_H diff --git a/runtime/hsa-runtime/loader/executable.cpp b/runtime/hsa-runtime/loader/executable.cpp index e804ef8b7f..6ca46154b2 100644 --- a/runtime/hsa-runtime/loader/executable.cpp +++ b/runtime/hsa-runtime/loader/executable.cpp @@ -63,8 +63,8 @@ #include "AMDHSAKernelDescriptor.h" -using namespace amd::hsa; -using namespace amd::hsa::common; +using namespace rocr::amd::hsa; +using namespace rocr::amd::hsa::common; // Having a side effect prevents call site optimization that allows removal of a noinline function call // with no side effect. @@ -82,6 +82,7 @@ HSA_API r_debug _amdgpu_r_debug = {2, 0}; static link_map* r_debug_tail = nullptr; +namespace rocr { namespace amd { namespace hsa { namespace loader { @@ -1906,3 +1907,4 @@ bool ExecutableImpl::PrintToFile(const std::string& filename) } // namespace loader } // namespace hsa } // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/loader/executable.hpp b/runtime/hsa-runtime/loader/executable.hpp index 727108a12d..c360df3488 100644 --- a/runtime/hsa-runtime/loader/executable.hpp +++ b/runtime/hsa-runtime/loader/executable.hpp @@ -62,6 +62,7 @@ #include "inc/amd_hsa_kernel_code.h" #include "amd_hsa_locks.hpp" +namespace rocr { namespace amd { namespace hsa { namespace loader { @@ -600,5 +601,6 @@ public: } // namespace loader } // namespace hsa } // namespace amd +} // namespace rocr #endif // HSA_RUNTIME_CORE_LOADER_EXECUTABLE_HPP_ diff --git a/runtime/hsa-runtime/loader/loaders.cpp b/runtime/hsa-runtime/loader/loaders.cpp index fdab349d6d..2c8dad0e14 100644 --- a/runtime/hsa-runtime/loader/loaders.cpp +++ b/runtime/hsa-runtime/loader/loaders.cpp @@ -44,6 +44,7 @@ #include #include "loaders.hpp" +namespace rocr { namespace amd { namespace hsa { namespace loader { @@ -253,6 +254,7 @@ namespace loader { return HSA_STATUS_SUCCESS; } -} -} -} +} // namespace loader +} // namespace hsa +} // namespace amd +} // namespace rocr diff --git a/runtime/hsa-runtime/loader/loaders.hpp b/runtime/hsa-runtime/loader/loaders.hpp index 9f26f82cb7..728d6d0832 100644 --- a/runtime/hsa-runtime/loader/loaders.hpp +++ b/runtime/hsa-runtime/loader/loaders.hpp @@ -47,6 +47,7 @@ #include #include +namespace rocr { namespace amd { namespace hsa { namespace loader { @@ -100,8 +101,9 @@ namespace loader { hsa_status_t SamplerDestroy( hsa_agent_t agent, hsa_ext_sampler_t sampler_handle) override; }; -} -} -} +} // namespace loader +} // namespace hsa +} // namespace amd +} // namespace rocr #endif // LOADERS_HPP_