diff --git a/include/hc.hpp b/include/hc.hpp
index 9e208830fb6..67580b999d0 100644
--- a/include/hc.hpp
+++ b/include/hc.hpp
@@ -264,7 +264,7 @@ class accelerator_view {
      * Returns the maximum size of tile static area available on this
      * accelerator view.
      */
-    size_t get_max_tile_static_size() {
+    size_t get_max_tile_static_size() const {
         return pQueue.get()->getDev()->GetMaxTileStaticSize();
     }
 
@@ -275,7 +275,7 @@ class accelerator_view {
      * The number returned would be immediately obsolete. This functions shall
      * only be used for testing and debugging purpose.
      */
-    int get_pending_async_ops() {
+    int get_pending_async_ops() const {
         return pQueue->getPendingAsyncOps();
     }
 
@@ -285,7 +285,7 @@ class accelerator_view {
      * @return An opaque handle of the underlying HSA queue, if the accelerator
      *         view is based on HSA.  NULL if otherwise.
      */
-    void* get_hsa_queue() {
+    void* get_hsa_queue() const {
         return pQueue->getHSAQueue();
     }
 
@@ -295,7 +295,7 @@ class accelerator_view {
      * @return An opaque handle of the underlying HSA agent, if the accelerator
      *         view is based on HSA.  NULL otherwise.
      */
-    void* get_hsa_agent() {
+    void* get_hsa_agent() const {
         return pQueue->getHSAAgent();
     }
 
@@ -307,7 +307,7 @@ class accelerator_view {
      * @return An opaque handle of the region, if the accelerator is based
      *         on HSA.  NULL otherwise.
      */
-    void* get_hsa_am_region() {
+    void* get_hsa_am_region() const {
         return pQueue->getHSAAMRegion();
     }
 
@@ -320,7 +320,7 @@ class accelerator_view {
      * @return An opaque handle of the region, if the accelerator is based
      *         on HSA.  NULL otherwise.
      */
-    void* get_hsa_am_system_region() {
+    void* get_hsa_am_system_region() const {
         return pQueue->getHSAAMHostRegion();
     }
 
@@ -331,14 +331,14 @@ class accelerator_view {
      * @return An opaque handle of the region, if the accelerator view is based
      *         on HSA.  NULL otherwise.
      */
-    void* get_hsa_kernarg_region() {
+    void* get_hsa_kernarg_region() const {
         return pQueue->getHSAKernargRegion();
     }
 
     /**
      * Returns if the accelerator view is based on HSA.
      */
-    bool is_hsa_accelerator() {
+    bool is_hsa_accelerator() const {
         return pQueue->hasHSAInterOp();
     }
 
@@ -669,14 +669,14 @@ class accelerator
      * Returns the maximum size of tile static area available on this
      * accelerator.
      */
-    size_t get_max_tile_static_size() {
+    size_t get_max_tile_static_size() const {
       return get_default_view().get_max_tile_static_size();
     }
   
     /**
      * Returns a vector of all accelerator_view associated with this accelerator.
      */
-    std::vector<accelerator_view> get_all_views() {
+    std::vector<accelerator_view> get_all_views() const {
         std::vector<accelerator_view> result;
         std::vector< std::shared_ptr<Kalmar::KalmarQueue> > queues = pDev->get_all_queues();
         for (auto q : queues) {
@@ -763,7 +763,7 @@ class accelerator
      * Check if @p other is peer of this accelerator.
      *
      * @return true if other can access this accelerator's device memory pool or false if not.
-     * the acceleratos is its own peer.
+     * the accelerator is its own peer.
      */
     bool get_is_peer(const accelerator& other) const {
         return pDev->is_peer(other.pDev);
@@ -970,7 +970,7 @@ class completion_future {
      * this completion_future object. The method is mostly used for debugging
      * purpose.
      */
-    void* get_native_handle() {
+    void* get_native_handle() const {
       if (__asyncOp != nullptr) {
         return __asyncOp->getNativeHandle();
       } else {
@@ -1012,7 +1012,7 @@ class completion_future {
      * @return An implementation-defined frequency in Hz in case the instance is
      *         created by a kernel dispatch or a barrier packet. 0 otherwise.
      */
-    uint64_t get_tick_frequency() {
+    uint64_t get_tick_frequency() const {
       if (__asyncOp != nullptr) {
         return __asyncOp->getTimestampFrequency();
       } else {
diff --git a/include/kalmar_runtime.h b/include/kalmar_runtime.h
index eb1c74b5d6c..d345e02b886 100644
--- a/include/kalmar_runtime.h
+++ b/include/kalmar_runtime.h
@@ -65,7 +65,7 @@ struct rw_info;
 class KalmarAsyncOp {
 public:
   virtual ~KalmarAsyncOp() {} 
-  virtual std::shared_future<void>* getFuture() { return nullptr; }
+  virtual std::shared_future<void>* getFuture() const { return nullptr; }
   virtual void* getNativeHandle() { return nullptr;}
 
   /**
@@ -87,7 +87,7 @@ class KalmarAsyncOp {
    *
    * @return An implementation-defined frequency for the asynchronous operation.
    */
-  virtual uint64_t getTimestampFrequency() { return 0L; }
+  virtual uint64_t getTimestampFrequency() const { return 0L; }
 
   /**
    * Get if the async operations has been completed.
@@ -149,9 +149,9 @@ class KalmarQueue
   /// push device pointer to kernel argument list
   virtual void Push(void *kernel, int idx, void* device, bool modify) = 0;
 
-  virtual uint32_t GetGroupSegmentSize(void *kernel) { return 0; }
+  virtual uint32_t GetGroupSegmentSize(void *kernel) const { return 0; }
 
-  KalmarDevice* getDev() { return pDev; }
+  KalmarDevice* getDev() const { return pDev; }
   queuing_mode get_mode() const { return mode; }
   void set_mode(queuing_mode mod) { mode = mod; }
 
@@ -175,7 +175,7 @@ class KalmarQueue
   virtual void* getHSAKernargRegion() { return nullptr; }
 
   /// check if the queue is an HSA queue
-  virtual bool hasHSAInterOp() { return false; }
+  virtual bool hasHSAInterOp() const { return false; }
 
   /// enqueue marker
   virtual std::shared_ptr<KalmarAsyncOp> EnqueueMarker() { return nullptr; }
@@ -278,7 +278,7 @@ class KalmarDevice
     }
 
     /// get max tile static area size
-    virtual size_t GetMaxTileStaticSize() { return 0; }
+    virtual size_t GetMaxTileStaticSize() const { return 0; }
 
     /// get all queues associated with this device
     virtual std::vector< std::shared_ptr<KalmarQueue> > get_all_queues() { return std::vector< std::shared_ptr<KalmarQueue> >(); }
@@ -287,13 +287,13 @@ class KalmarDevice
 
     virtual void memcpySymbol(void* symbolAddr, void* hostptr, size_t count, size_t offset = 0, hcMemcpyKind kind = hcMemcpyHostToDevice) {}
 
-    virtual void* getSymbolAddress(const char* symbolName) { return nullptr; }
+    virtual void* getSymbolAddress(const char* symbolName) const { return nullptr; }
 
     /// get underlying native agent handle
     virtual void* getHSAAgent() { return nullptr; }
 
     /// get the profile of the agent
-    virtual hcAgentProfile getProfile() { return hcAgentProfileNone; }
+    virtual hcAgentProfile getProfile() const { return hcAgentProfileNone; }
 
     /// check if @p other can access to this device's device memory, return true if so, false otherwise
     virtual bool is_peer(const KalmarDevice* other) {return false;}
@@ -376,7 +376,7 @@ class KalmarContext
 public:
     virtual ~KalmarContext() {}
 
-    std::vector<KalmarDevice*> getDevices() { return Devices; }
+    std::vector<KalmarDevice*> getDevices() const { return Devices; }
 
     /// set default device by path
     bool set_default(const std::wstring& path) {
@@ -410,10 +410,10 @@ class KalmarContext
     }
 
     /// get system ticks
-    virtual uint64_t getSystemTicks() { return 0L; };
+    virtual uint64_t getSystemTicks() const { return 0L; };
 
     /// get tick frequency
-    virtual uint64_t getSystemTickFrequency() { return 0L; };
+    virtual uint64_t getSystemTickFrequency() const { return 0L; };
 };
 
 KalmarContext *getContext();
@@ -552,7 +552,7 @@ struct rw_info
     /// construct array
     /// According to AMP standard, array should be constructed with
     /// 1. one accelerator_view
-    /// 2. one acceleratir_view, with another staged one
+    /// 2. one accelerator_view, with another staged one
     ///    In this case, master should be cpu device
     ///    If it is not, ignore the stage one, fallback to case 1.
     rw_info(const std::shared_ptr<KalmarQueue>& Queue, const std::shared_ptr<KalmarQueue>& Stage,
diff --git a/lib/hsa/mcwamp_hsa.cpp b/lib/hsa/mcwamp_hsa.cpp
index 1a35d382951..7e87c2646f3 100644
--- a/lib/hsa/mcwamp_hsa.cpp
+++ b/lib/hsa/mcwamp_hsa.cpp
@@ -228,7 +228,7 @@ class HSABarrier : public Kalmar::KalmarAsyncOp {
     Kalmar::HSAQueue* hsaQueue;
 
 public:
-    std::shared_future<void>* getFuture() override { return future; }
+    std::shared_future<void>* getFuture() const override { return future; }
 
     void* getNativeHandle() override { return &signal; }
 
@@ -270,7 +270,7 @@ class HSABarrier : public Kalmar::KalmarAsyncOp {
 
     void dispose();
 
-    uint64_t getTimestampFrequency() override {
+    uint64_t getTimestampFrequency() const override {
         // get system tick frequency
         uint64_t timestamp_frequency_hz = 0L;
         hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &timestamp_frequency_hz);
@@ -312,7 +312,7 @@ class HSADispatch : public Kalmar::KalmarAsyncOp {
     Kalmar::HSAQueue* hsaQueue;
 
 public:
-    std::shared_future<void>* getFuture() override { return future; }
+    std::shared_future<void>* getFuture() const override { return future; }
 
     void* getNativeHandle() override { return &signal; }
 
@@ -371,7 +371,7 @@ class HSADispatch : public Kalmar::KalmarAsyncOp {
 
     hsa_status_t dispatchKernelAsync(Kalmar::HSAQueue*);
 
-    uint32_t getGroupSegmentSize() {
+    uint32_t getGroupSegmentSize() const {
         hsa_status_t status = HSA_STATUS_SUCCESS;
         uint32_t group_segment_size = 0;
         status = hsa_executable_symbol_get_info(kernel->hsaExecutableSymbol,
@@ -389,7 +389,7 @@ class HSADispatch : public Kalmar::KalmarAsyncOp {
 
     void dispose();
 
-    uint64_t getTimestampFrequency() override {
+    uint64_t getTimestampFrequency() const override {
         // get system tick frequency
         uint64_t timestamp_frequency_hz = 0L;
         hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &timestamp_frequency_hz);
@@ -701,7 +701,7 @@ class HSAQueue final : public KalmarQueue
         return sp_dispatch;
     }
 
-    uint32_t GetGroupSegmentSize(void *ker) override {
+    uint32_t GetGroupSegmentSize(void *ker) const override {
         HSADispatch *dispatch = reinterpret_cast<HSADispatch*>(ker);
         return dispatch->getGroupSegmentSize();
     }
@@ -917,7 +917,7 @@ class HSAQueue final : public KalmarQueue
         }
     }
 
-    void* getHSAQueue() override {
+    void* getHSAQueue() const override {
         return static_cast<void*>(commandQueue);
     }
 
@@ -1005,11 +1005,11 @@ class HSADevice final : public KalmarDevice
 
 public:
  
-    uint32_t getWorkgroupMaxSize() {
+    uint32_t getWorkgroupMaxSize() const {
         return workgroup_max_size;
     }
 
-    const uint16_t* getWorkgroupMaxDim() {
+    const uint16_t* getWorkgroupMaxDim() const {
         return &workgroup_max_dim[0];
     }
 
@@ -1497,7 +1497,7 @@ class HSADevice final : public KalmarDevice
         return q;
     }
 
-    size_t GetMaxTileStaticSize() override {
+    size_t GetMaxTileStaticSize() const override {
         return max_tile_static_size;
     }
 
@@ -1544,7 +1544,7 @@ class HSADevice final : public KalmarDevice
       auto self_pool = getHSAAMRegion();
       hsa_amd_memory_pool_access_t access;
 
-      hsa_agent_t* agent = static_cast<hsa_agent_t*>( const_cast<KalmarDevice *> (other)->getHSAAgent());
+      hsa_agent_t* agent = static_cast<hsa_agent_t*>(const_cast<KalmarDevice*>(other)->getHSAAgent());
 
       //TODO: CPU acclerator will return NULL currently, return false.
       if(nullptr == agent)
@@ -1702,7 +1702,7 @@ class HSADevice final : public KalmarDevice
         return std::make_pair(ret, cursor);
     }
 
-    void* getSymbolAddress(const char* symbolName) override {
+    void* getSymbolAddress(const char* symbolName) const override {
         hsa_status_t status;
 
         unsigned long* symbol_ptr = nullptr;
@@ -2209,14 +2209,14 @@ class HSAContext final : public KalmarContext
 #endif
     }
 
-    uint64_t getSystemTicks() override {
+    uint64_t getSystemTicks() const override {
         // get system tick
         uint64_t timestamp = 0L;
         hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP, &timestamp);
         return timestamp;
     }
 
-    uint64_t getSystemTickFrequency() override {
+    uint64_t getSystemTickFrequency() const override {
         // get system tick frequency
         uint64_t timestamp_frequency_hz = 0L;
         hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &timestamp_frequency_hz);