Merge remote-tracking branch 'origin/hsa_pool_api' into develop

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 {

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,

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