Add rocr namespace to core header and impl files

Change-Id: I1e1b33f9bba1078d049bc19797889988c3e43360

runtime/hsa-runtime/core/common/shared.cpp

@@ -42,7 +42,9 @@
 
 #include "core/common/shared.h"
 
+namespace rocr {
 namespace core {
 std::function<void*(size_t, size_t, uint32_t)> BaseShared::allocate_ = nullptr;
 std::function<void(void*)> BaseShared::free_ = nullptr;
-}
+}   // namespace core
+}   // namespace rocr

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 <typename T, size_t Align> class SharedArray final : private BaseShared
 };
 
 }  // namespace core
+}  // namespace rocr
 #endif  // header guard

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

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

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

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<uint64_t, true, -1, false> BlitSdmaV4;
 typedef BlitSdma<uint64_t, true, -1, true> BlitSdmaV5;
 
 }  // namespace amd
+}  // namespace rocr
 
 #endif  // header guard

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

runtime/hsa-runtime/core/inc/amd_elf_image.hpp

@@ -87,8 +87,9 @@ mes of its
 #include <vector>
 #include <memory>
 
+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_

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_

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

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

runtime/hsa-runtime/core/inc/amd_hsa_code.hpp

@@ -92,6 +92,7 @@ mes of its
 #include <cassert>
 #include <unordered_map>
 
+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_

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

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

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<MemoryRegion*>(region.handle);
   }
@@ -204,6 +205,7 @@ class MemoryRegion : public core::MemoryRegion {
   mutable SimpleHeap<BlockAllocator> fragment_allocator_;
 };
 
-}  // namespace
+}  // namespace amd
+}  // namespace rocr
 
 #endif  // header guard

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

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

runtime/hsa-runtime/core/inc/cache.h

@@ -49,6 +49,7 @@
 #include <utility>
 #include <string>
 
+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

runtime/hsa-runtime/core/inc/checked.h

@@ -46,6 +46,7 @@
 #include <stdint.h>
 #include <stdlib.h>
 
+namespace rocr {
 namespace core {
 
 /// @brief Compares type codes and pointers to check object validity.  Used for cast validation.
@@ -105,4 +106,6 @@ template <uint64_t code> class Checked {
 };
 
 }  // namespace core
+}  // namespace rocr
+
 #endif  // header guard

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

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 R, class... Args> class callback_t<R (*)(Args...)> {
   func_t function;
 };
 
-}  // namespace AMD
+}  // namespace amd
+}  // namespace rocr
 
 #endif  // HSA_RUNTIME_CORE_INC_EXCEPTIONS_H

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

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

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

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

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

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_

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

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_

runtime/hsa-runtime/core/inc/isa.h

@@ -50,6 +50,7 @@
 #include <unordered_map>
 #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_

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

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

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

runtime/hsa-runtime/core/inc/sdma_registers.h

@@ -46,7 +46,8 @@
 #include <stddef.h>
 #include <stdint.h>
 
-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_

runtime/hsa-runtime/core/inc/signal.h

@@ -75,6 +75,7 @@ template <> struct less<hsa_signal_t> {
 };
 }
 
+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<core::Signal, SignalDeleter>;
 
 }  // namespace core
+}  // namespace rocr
 #endif  // header guard

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<const amd::MemoryRegion*>(region);
+    const MemoryRegion* amdregion = static_cast<const AMD::MemoryRegion*>(region);
     uint64_t base = amdregion->GetBaseAddress();
 
     if (amdregion->IsLDS()) {
@@ -1119,3 +1120,4 @@ void AqlQueue::InitScratchSRD() {
   return;
 }
 }  // namespace amd
+}  // namespace rocr

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<const GpuAgent&>(agent);
+  const AMD::GpuAgent& gpuAgent = static_cast<const AMD::GpuAgent&>(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<std::mutex> guard(lock_);
 
-  const GpuAgent& gpuAgent = static_cast<const GpuAgent&>(agent);
+  const AMD::GpuAgent& gpuAgent = static_cast<const AMD::GpuAgent&>(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

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<uint32_t>(reinterpret_cast<uintptr_t>(p));
@@ -138,7 +139,7 @@ hsa_status_t BlitSdma<RingIndexTy, HwIndexMonotonic, SizeToCountOffset, useGCR>:
     return HSA_STATUS_ERROR;
   }
 
-  agent_ = reinterpret_cast<amd::GpuAgent*>(&const_cast<core::Agent&>(agent));
+  agent_ = reinterpret_cast<AMD::GpuAgent*>(&const_cast<core::Agent&>(agent));
 
   if (HSA_PROFILE_FULL == agent_->profile()) {
     assert(false && "Only support SDMA for dgpu currently");
@@ -953,3 +954,4 @@ template class BlitSdma<uint64_t, true, -1, false>;
 template class BlitSdma<uint64_t, true, -1, true>;
 
 }  // namespace amd
+}  // namespace rocr

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

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

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<decltype(callback)> call(callback);
   for (const core::MemoryRegion* region : regions) {
-    const amd::MemoryRegion* amd_region =
-        reinterpret_cast<const amd::MemoryRegion*>(region);
+    const AMD::MemoryRegion* amd_region =
+        reinterpret_cast<const AMD::MemoryRegion*>(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<core::Blit>& GpuAgent::GetBlitObject(const core::Agent& dst_agent,
   return GetXgmiBlit(dst_agent);
 }
 
-}  // namespace
+}  // namespace amd
+}  // namespace rocr

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

runtime/hsa-runtime/core/runtime/amd_loader_context.cpp

@@ -62,11 +62,12 @@
 #include <sys/mman.h>
 #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<amd::MemoryRegion*>(
+  core::Agent* agent = reinterpret_cast<AMD::MemoryRegion*>(
                            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

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<KernelMutex> 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

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

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

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

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

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 <class T>
@@ -158,7 +160,7 @@ template <class T> 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

runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp

@@ -48,6 +48,7 @@
 #include <iostream>
 
 // 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

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 <class T>
 struct ValidityError;
 template <>
@@ -79,7 +81,7 @@ struct ValidityError<core::MemoryRegion*> {
 };
 
 template <>
-struct ValidityError<amd::MemoryRegion*> {
+struct ValidityError<AMD::MemoryRegion*> {
   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<const amd::GpuAgentInt*>(agent);
+  const AMD::GpuAgentInt* gpu_agent =
+      static_cast<const AMD::GpuAgentInt*>(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<amd::GpuAgent*>(agent);
+  AMD::GpuAgent* gpu_agent = static_cast<AMD::GpuAgent*>(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<amd::GpuAgent*>(base_agent);
+  AMD::GpuAgent* agent = static_cast<AMD::GpuAgent*>(base_agent);
 
   std::vector<core::Signal*> 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<amd::GpuAgentInt*>(agent);
+  AMD::GpuAgentInt* gpu_agent = static_cast<AMD::GpuAgentInt*>(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<amd::GpuAgentInt*>(agent)->TranslateTime(signal, *time);
+    static_cast<AMD::GpuAgentInt*>(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<amd::GpuAgentInt*>(agent);
+  AMD::GpuAgentInt* gpu_agent = static_cast<AMD::GpuAgentInt*>(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*>(
+  const AMD::MemoryRegion* system_region = static_cast<const AMD::MemoryRegion*>(
       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*>(
+  const AMD::MemoryRegion* system_region =
+      reinterpret_cast<const AMD::MemoryRegion*>(
           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<const amd::CpuAgent*>(agent)->VisitRegion(
+    return reinterpret_cast<const AMD::CpuAgent*>(agent)->VisitRegion(
         false, reinterpret_cast<hsa_status_t (*)(hsa_region_t memory_pool,
                                                  void* data)>(callback),
         data);
   }
 
-  return reinterpret_cast<const amd::GpuAgentInt*>(agent)->VisitRegion(
+  return reinterpret_cast<const AMD::GpuAgentInt*>(agent)->VisitRegion(
       false,
       reinterpret_cast<hsa_status_t (*)(hsa_region_t memory_pool, void* data)>(
           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_t>(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_t>(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_t>(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_t>(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

runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp

@@ -46,6 +46,7 @@
 
 #include <string>
 
+namespace rocr {
 // Implementations for missing / unsupported extensions
 template <class R, class... ARGS> 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);
 }

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) {

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

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

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

runtime/hsa-runtime/core/runtime/isa.cpp

@@ -46,6 +46,7 @@
 #include <sstream>
 #include <utility>
 
+namespace rocr {
 namespace core {
 
 bool Wavefront::GetInfo(
@@ -240,3 +241,4 @@ const IsaRegistry::IsaMap IsaRegistry::GetSupportedIsas() {
 }
 
 } // namespace core
+} // namespace rocr

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

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<const amd::MemoryRegion*>(system_regions_fine_[0]);
+  const AMD::MemoryRegion* system_region =
+      static_cast<const AMD::MemoryRegion*>(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<const amd::MemoryRegion*>(it->second.region);
+    amd_region = reinterpret_cast<const AMD::MemoryRegion*>(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<const amd::MemoryRegion*>(it->second.region);
+      const AMD::MemoryRegion* region = static_cast<const AMD::MemoryRegion*>(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<amd::GpuAgentInt*>(agent)->PostToolsInit();
+        reinterpret_cast<AMD::GpuAgentInt*>(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

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

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

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

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 <typename T> class lazy_ptr {
 
 };
 
+}  // namespace rocr
+
 #endif  // HSA_RUNTIME_CORE_UTIL_LAZY_PTR_H_

runtime/hsa-runtime/core/util/lnx/os_linux.cpp

@@ -60,6 +60,7 @@
 #include <string>
 #include <utility>
 
+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

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_

runtime/hsa-runtime/core/util/os.h

@@ -48,6 +48,7 @@
 #include <string>
 #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_

runtime/hsa-runtime/core/util/simple_heap.h

@@ -53,6 +53,8 @@
 
 #include "core/util/utils.h"
 
+namespace rocr {
+
 template <typename Allocator> class SimpleHeap {
  private:
   struct Fragment_T {
@@ -298,4 +300,6 @@ template <typename Allocator> class SimpleHeap {
   }
 };
 
+}  // namespace rocr
+
 #endif  // HSA_RUNTME_CORE_UTIL_SIMPLE_HEAP_H_

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

runtime/hsa-runtime/core/util/small_heap.h

@@ -51,7 +51,9 @@
 #include <set>
 
 #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

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

runtime/hsa-runtime/core/util/timer.h

@@ -49,6 +49,7 @@
 #include <time.h>
 #include <type_traits>
 
+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

runtime/hsa-runtime/core/util/utils.h

@@ -53,6 +53,8 @@
 #include <string>
 #include <algorithm>
 
+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 <uint32_t lowBit, uint32_t highBit, typename T>
 static __forceinline uint32_t BitSelect(T p) {
   static_assert(sizeof(T) <= sizeof(uintptr_t), "Type out of range.");

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

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

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

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

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_

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

runtime/hsa-runtime/libamdhsacode/amd_hsa_locks.hpp

@@ -47,6 +47,7 @@
 #include <cstddef>
 #include <mutex>
 
+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

runtime/hsa-runtime/libamdhsacode/amd_options.cpp

@@ -54,6 +54,7 @@
 
 #include <cstddef>
 
+namespace rocr {
 namespace amd {
 namespace options {
 
@@ -379,3 +380,4 @@ void OptionParser::Reset() {
 
 } // namespace options
 } // namespace amd
+} // namespace rocr

runtime/hsa-runtime/libamdhsacode/amd_options.hpp

@@ -55,6 +55,7 @@
 #include <unordered_map>
 #include <unordered_set>
 
+namespace rocr {
 namespace amd {
 namespace options {
 
@@ -472,5 +473,6 @@ private:
 
 } // namespace options
 } // namespace amd
+} // namespace rocr
 
 #endif // AMD_OPTIONS_HPP

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

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

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_

runtime/hsa-runtime/loader/loaders.cpp

@@ -44,6 +44,7 @@
 #include <cassert>
 #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

runtime/hsa-runtime/loader/loaders.hpp

@@ -47,6 +47,7 @@
 #include <set>
 #include <iostream>
 
+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_