diff --git a/runtime/src/iree-amd-aie/driver/hsa/CMakeLists.txt b/runtime/src/iree-amd-aie/driver/hsa/CMakeLists.txt
index 5a051fd80d..ef64152597 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/CMakeLists.txt
+++ b/runtime/src/iree-amd-aie/driver/hsa/CMakeLists.txt
@@ -25,9 +25,9 @@ iree_cc_library(
   TEXTUAL_HDRS
     "dynamic_symbol_tables.h"
   SRCS
-    "dynamic_symbols.c"
+    "dynamic_symbols.cc"
     "hsa_headers.h"
-    "status_util.c"
+    "status_util.cc"
   DEPS
     hsa-runtime64::hsa-runtime64
     iree::base
@@ -43,26 +43,26 @@ iree_cc_library(
     "api.h"
   SRCS
     "api.h"
-    "event_pool.c"
+    "event_pool.cc"
     "event_pool.h"
-    "event_semaphore.c"
+    "event_semaphore.cc"
     "event_semaphore.h"
-    "hsa_allocator.c"
+    "hsa_allocator.cc"
     "hsa_allocator.h"
-    "hsa_buffer.c"
+    "hsa_buffer.cc"
     "hsa_buffer.h"
-    "hsa_device.c"
+    "hsa_device.cc"
     "hsa_device.h"
-    "hsa_driver.c"
-    "native_executable.c"
+    "hsa_driver.cc"
+    "native_executable.cc"
     "native_executable.h"
-    "nop_executable_cache.c"
+    "nop_executable_cache.cc"
     "nop_executable_cache.h"
-    "pending_queue_actions.c"
+    "pending_queue_actions.cc"
     "pending_queue_actions.h"
-    "queue_command_buffer.c"
+    "queue_command_buffer.cc"
     "queue_command_buffer.h"
-    "timepoint_pool.c"
+    "timepoint_pool.cc"
     "timepoint_pool.h"
   DEPS
     hsa-runtime64::hsa-runtime64
diff --git a/runtime/src/iree-amd-aie/driver/hsa/dynamic_symbol_tables.h b/runtime/src/iree-amd-aie/driver/hsa/dynamic_symbol_tables.h
index 2feed7f097..6c78417350 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/dynamic_symbol_tables.h
+++ b/runtime/src/iree-amd-aie/driver/hsa/dynamic_symbol_tables.h
@@ -9,7 +9,7 @@
 // HSA symbols
 //===----------------------------------------------------------------------===//
 
-#include <stdint.h>
+#include <cstdint>
 
 IREE_HAL_HSA_REQUIRED_PFN_DECL(hsa_init)
 IREE_HAL_HSA_REQUIRED_PFN_DECL(hsa_shut_down)
diff --git a/runtime/src/iree-amd-aie/driver/hsa/dynamic_symbols.c b/runtime/src/iree-amd-aie/driver/hsa/dynamic_symbols.cc
similarity index 93%
rename from runtime/src/iree-amd-aie/driver/hsa/dynamic_symbols.c
rename to runtime/src/iree-amd-aie/driver/hsa/dynamic_symbols.cc
index d7032e4444..04b67274f3 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/dynamic_symbols.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/dynamic_symbols.cc
@@ -7,12 +7,10 @@
 
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 
-#include <string.h>
+#include <cstring>
 
-#include "iree-amd-aie/driver/hsa/status_util.h"
 #include "iree/base/api.h"
 #include "iree/base/internal/dynamic_library.h"
-#include "iree/base/target_platform.h"
 
 //===----------------------------------------------------------------------===//
 // HSA dynamic symbols
@@ -36,7 +34,7 @@ static iree_status_t iree_hal_hsa_dynamic_symbols_resolve_all(
         syms->dylib, name, (void**)&syms->hsa_symbol_name)); \
   }
 
-#include "iree-amd-aie/driver/hsa/dynamic_symbol_tables.h"  // IWYU pragma: keep
+#include "dynamic_symbol_tables.h"  // IWYU pragma: keep
 #undef IREE_HAL_HSA_REQUIRED_PFN_DECL
   return iree_ok_status();
 }
diff --git a/runtime/src/iree-amd-aie/driver/hsa/event_pool.c b/runtime/src/iree-amd-aie/driver/hsa/event_pool.cc
similarity index 96%
rename from runtime/src/iree-amd-aie/driver/hsa/event_pool.c
rename to runtime/src/iree-amd-aie/driver/hsa/event_pool.cc
index f5f3d167b2..687bba0d07 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/event_pool.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/event_pool.cc
@@ -8,16 +8,13 @@
 
 #include "iree-amd-aie/driver/hsa/event_pool.h"
 
-#include <stdbool.h>
-#include <stddef.h>
-#include <string.h>
+#include <cstring>
 
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 #include "iree-amd-aie/driver/hsa/status_util.h"
 #include "iree/base/api.h"
 #include "iree/base/internal/atomics.h"
 #include "iree/base/internal/synchronization.h"
-#include "iree/hal/api.h"
 
 //===----------------------------------------------------------------------===//
 // iree_hal_hsa_event_t
@@ -35,8 +32,8 @@ struct iree_hal_hsa_event_t {
   // The symbols used to create and destroy signals objects.
   const iree_hal_hsa_dynamic_symbols_t* symbols;
 
-  // The event pool that owns this event. This cannot be NULL. We retain it to
-  // make sure the event outlive the pool.
+  // The event pool that owns this event. This cannot be nullptr. We retain it
+  // to make sure the event outlive the pool.
   iree_hal_hsa_event_pool_t* pool;
 
   hsa_signal_t signal;
@@ -65,10 +62,10 @@ static inline iree_status_t iree_hal_hsa_event_create(
   IREE_ASSERT_ARGUMENT(symbols);
   IREE_ASSERT_ARGUMENT(pool);
   IREE_ASSERT_ARGUMENT(out_event);
-  *out_event = NULL;
+  *out_event = nullptr;
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  iree_hal_hsa_event_t* event = NULL;
+  iree_hal_hsa_event_t* event = nullptr;
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0,
       iree_allocator_malloc(host_allocator, sizeof(*event), (void**)&event));
@@ -79,7 +76,7 @@ static inline iree_status_t iree_hal_hsa_event_create(
 
   hsa_signal_value_t signal_value = 1;
   uint32_t num_consumers = 0;
-  const hsa_agent_t* consumers = NULL;
+  const hsa_agent_t* consumers = nullptr;
 
   iree_status_t status = IREE_HSA_RESULT_TO_STATUS(
       symbols,
@@ -153,10 +150,10 @@ iree_status_t iree_hal_hsa_event_pool_allocate(
     iree_hal_hsa_event_pool_t** out_event_pool) {
   IREE_ASSERT_ARGUMENT(symbols);
   IREE_ASSERT_ARGUMENT(out_event_pool);
-  *out_event_pool = NULL;
+  *out_event_pool = nullptr;
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  iree_hal_hsa_event_pool_t* event_pool = NULL;
+  iree_hal_hsa_event_pool_t* event_pool = nullptr;
   iree_host_size_t total_size =
       sizeof(*event_pool) +
       available_capacity * sizeof(*event_pool->available_list);
diff --git a/runtime/src/iree-amd-aie/driver/hsa/event_semaphore.c b/runtime/src/iree-amd-aie/driver/hsa/event_semaphore.cc
similarity index 96%
rename from runtime/src/iree-amd-aie/driver/hsa/event_semaphore.c
rename to runtime/src/iree-amd-aie/driver/hsa/event_semaphore.cc
index 445f177f86..031cb612d4 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/event_semaphore.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/event_semaphore.cc
@@ -8,7 +8,6 @@
 #include "iree-amd-aie/driver/hsa/event_semaphore.h"
 
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
-#include "iree-amd-aie/driver/hsa/status_util.h"
 #include "iree-amd-aie/driver/hsa/timepoint_pool.h"
 #include "iree/base/internal/synchronization.h"
 #include "iree/base/internal/wait_handle.h"
@@ -45,7 +44,9 @@ typedef struct iree_hal_hsa_semaphore_t {
   iree_status_t failure_status IREE_GUARDED_BY(mutex);
 } iree_hal_hsa_semaphore_t;
 
-static const iree_hal_semaphore_vtable_t iree_hal_hsa_semaphore_vtable;
+namespace {
+extern const iree_hal_semaphore_vtable_t iree_hal_hsa_semaphore_vtable;
+}
 
 static iree_hal_hsa_semaphore_t* iree_hal_hsa_semaphore_cast(
     iree_hal_semaphore_t* base_value) {
@@ -64,7 +65,7 @@ iree_status_t iree_hal_hsa_event_semaphore_create(
   IREE_ASSERT_ARGUMENT(out_semaphore);
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  iree_hal_hsa_semaphore_t* semaphore = NULL;
+  iree_hal_hsa_semaphore_t* semaphore = nullptr;
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_allocator_malloc(host_allocator, sizeof(*semaphore),
                                 (void**)&semaphore));
@@ -235,14 +236,14 @@ static iree_status_t iree_hal_hsa_semaphore_acquire_timepoint_host_wait(
 static bool iree_hal_hsa_semaphore_acquire_event_host_wait(
     iree_hal_hsa_semaphore_t* semaphore, uint64_t min_value,
     iree_hal_hsa_event_t** out_event) {
-  *out_event = NULL;
+  *out_event = nullptr;
   IREE_TRACE_ZONE_BEGIN(z0);
 
   // Scan through the timepoint list and try to find a device event signal to
   // wait on. We need to lock with the timepoint list mutex here.
   iree_slim_mutex_lock(&semaphore->base.timepoint_mutex);
   for (iree_hal_semaphore_timepoint_t* tp = semaphore->base.timepoint_list.head;
-       tp != NULL; tp = tp->next) {
+       tp != nullptr; tp = tp->next) {
     iree_hal_hsa_timepoint_t* signal_timepoint = (iree_hal_hsa_timepoint_t*)tp;
     if (signal_timepoint->kind == IREE_HAL_HSA_TIMEPOINT_KIND_DEVICE_SIGNAL &&
         signal_timepoint->base.minimum_value >= min_value) {
@@ -254,7 +255,7 @@ static bool iree_hal_hsa_semaphore_acquire_event_host_wait(
   iree_slim_mutex_unlock(&semaphore->base.timepoint_mutex);
 
   IREE_TRACE_ZONE_END(z0);
-  return *out_event != NULL;
+  return *out_event != nullptr;
 }
 
 static iree_status_t iree_hal_hsa_semaphore_wait(
@@ -290,7 +291,7 @@ static iree_status_t iree_hal_hsa_semaphore_wait(
   // Slow path: try to see if we can have a device signal to wait on. This
   // should happen outside of the lock given that acquiring has its own internal
   // locks. This is faster than waiting on a host timepoint.
-  iree_hal_hsa_event_t* wait_event = NULL;
+  iree_hal_hsa_event_t* wait_event = nullptr;
   if (iree_hal_hsa_semaphore_acquire_event_host_wait(semaphore, value,
                                                      &wait_event)) {
     semaphore->symbols->hsa_signal_wait_scacquire(
@@ -304,7 +305,7 @@ static iree_status_t iree_hal_hsa_semaphore_wait(
 
   // Slow path: acquire a timepoint. This should happen outside of the lock too
   // given that acquiring has its own internal locks.
-  iree_hal_hsa_timepoint_t* timepoint = NULL;
+  iree_hal_hsa_timepoint_t* timepoint = nullptr;
   iree_status_t status = iree_hal_hsa_semaphore_acquire_timepoint_host_wait(
       semaphore, value, timeout, &timepoint);
   if (IREE_UNLIKELY(!iree_status_is_ok(status))) {
@@ -343,13 +344,13 @@ iree_status_t iree_hal_hsa_semaphore_multi_wait(
   // Avoid heap allocations by using the device block pool for the wait set.
   iree_arena_allocator_t arena;
   iree_arena_initialize(block_pool, &arena);
-  iree_wait_set_t* wait_set = NULL;
+  iree_wait_set_t* wait_set = nullptr;
   iree_status_t status = iree_wait_set_allocate(
       semaphore_list.count, iree_arena_allocator(&arena), &wait_set);
 
   // Acquire a host wait handle for each semaphore timepoint we are to wait on.
   iree_host_size_t timepoint_count = 0;
-  iree_hal_hsa_timepoint_t** timepoints = NULL;
+  iree_hal_hsa_timepoint_t** timepoints = nullptr;
   iree_host_size_t total_timepoint_size =
       semaphore_list.count * sizeof(timepoints[0]);
   bool needs_wait = true;
@@ -377,7 +378,7 @@ iree_status_t iree_hal_hsa_semaphore_multi_wait(
         // Slow path: get a native host wait handle for the timepoint. This
         // should happen outside of the lock given that acquiring has its own
         // internal locks.
-        iree_hal_hsa_timepoint_t* timepoint = NULL;
+        iree_hal_hsa_timepoint_t* timepoint = nullptr;
         status = iree_hal_hsa_semaphore_acquire_timepoint_host_wait(
             semaphore, semaphore_list.payload_values[i], timeout, &timepoint);
         if (iree_status_is_ok(status)) {
@@ -393,7 +394,8 @@ iree_status_t iree_hal_hsa_semaphore_multi_wait(
   // Perform the wait.
   if (iree_status_is_ok(status) && needs_wait) {
     if (wait_mode == IREE_HAL_WAIT_MODE_ANY) {
-      status = iree_wait_any(wait_set, deadline_ns, /*out_wake_handle=*/NULL);
+      status =
+          iree_wait_any(wait_set, deadline_ns, /*out_wake_handle=*/nullptr);
     } else {
       status = iree_wait_all(wait_set, deadline_ns);
     }
@@ -439,7 +441,7 @@ iree_status_t iree_hal_hsa_event_semaphore_acquire_timepoint_device_signal(
     hsa_signal_t* out_signal) {
   iree_hal_hsa_semaphore_t* semaphore =
       iree_hal_hsa_semaphore_cast(base_semaphore);
-  iree_hal_hsa_timepoint_t* signal_timepoint = NULL;
+  iree_hal_hsa_timepoint_t* signal_timepoint = nullptr;
   IREE_TRACE_ZONE_BEGIN(z0);
 
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
@@ -462,10 +464,10 @@ iree_status_t iree_hal_hsa_event_semaphore_acquire_timepoint_device_signal(
   // list mutex here.
   iree_slim_mutex_lock(&semaphore->base.timepoint_mutex);
   for (iree_hal_semaphore_timepoint_t* tp = semaphore->base.timepoint_list.head;
-       tp != NULL; tp = tp->next) {
+       tp != nullptr; tp = tp->next) {
     iree_hal_hsa_timepoint_t* wait_timepoint = (iree_hal_hsa_timepoint_t*)tp;
     if (wait_timepoint->kind == IREE_HAL_HSA_TIMEPOINT_KIND_DEVICE_WAIT &&
-        wait_timepoint->timepoint.device_wait == NULL &&
+        wait_timepoint->timepoint.device_wait == nullptr &&
         wait_timepoint->base.minimum_value <= to_value) {
       iree_hal_hsa_event_retain(event);
       wait_timepoint->timepoint.device_wait = event;
@@ -502,7 +504,7 @@ iree_status_t iree_hal_hsa_event_semaphore_acquire_timepoint_device_wait(
     hsa_signal_t* out_signal) {
   iree_hal_hsa_semaphore_t* semaphore =
       iree_hal_hsa_semaphore_cast(base_semaphore);
-  iree_hal_hsa_timepoint_t* wait_timepoint = NULL;
+  iree_hal_hsa_timepoint_t* wait_timepoint = nullptr;
   IREE_TRACE_ZONE_BEGIN(z0);
 
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
@@ -519,7 +521,7 @@ iree_status_t iree_hal_hsa_event_semaphore_acquire_timepoint_device_wait(
       },
       &wait_timepoint->base);
 
-  iree_hal_hsa_event_t* wait_event = NULL;
+  iree_hal_hsa_event_t* wait_event = nullptr;
   if (iree_hal_hsa_semaphore_acquire_event_host_wait(semaphore, min_value,
                                                      &wait_event)) {
     // We've found an existing signal timepoint to wait on; we don't need a
diff --git a/runtime/src/iree-amd-aie/driver/hsa/event_semaphore.h b/runtime/src/iree-amd-aie/driver/hsa/event_semaphore.h
index de1010eece..bc48e46609 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/event_semaphore.h
+++ b/runtime/src/iree-amd-aie/driver/hsa/event_semaphore.h
@@ -8,7 +8,7 @@
 #ifndef IREE_EXPERIMENTAL_HSA_EVENT_SEMAPHORE_H_
 #define IREE_EXPERIMENTAL_HSA_EVENT_SEMAPHORE_H_
 
-#include <stdint.h>
+#include <cstdint>
 
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 #include "iree-amd-aie/driver/hsa/pending_queue_actions.h"
diff --git a/runtime/src/iree-amd-aie/driver/hsa/hsa_allocator.c b/runtime/src/iree-amd-aie/driver/hsa/hsa_allocator.cc
similarity index 98%
rename from runtime/src/iree-amd-aie/driver/hsa/hsa_allocator.c
rename to runtime/src/iree-amd-aie/driver/hsa/hsa_allocator.cc
index a1a05c2560..40ab4e1470 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/hsa_allocator.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/hsa_allocator.cc
@@ -7,8 +7,6 @@
 
 #include "iree-amd-aie/driver/hsa/hsa_allocator.h"
 
-#include <stddef.h>
-
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 #include "iree-amd-aie/driver/hsa/hsa_buffer.h"
 #include "iree-amd-aie/driver/hsa/status_util.h"
@@ -41,7 +39,9 @@ typedef struct iree_hal_hsa_allocator_t {
   IREE_STATISTICS(iree_hal_allocator_statistics_t statistics;)
 } iree_hal_hsa_allocator_t;
 
-static const iree_hal_allocator_vtable_t iree_hal_hsa_allocator_vtable;
+namespace {
+extern const iree_hal_allocator_vtable_t iree_hal_hsa_allocator_vtable;
+}
 
 static iree_hal_hsa_allocator_t* iree_hal_hsa_allocator_cast(
     iree_hal_allocator_t* base_value) {
@@ -111,7 +111,7 @@ static hsa_status_t get_fine_grained_memory_pool(hsa_amd_memory_pool_t pool,
 static hsa_status_t iterate_find_cpu_agent_callback(hsa_agent_t agent,
                                                     void* base_allocator) {
   iree_hal_hsa_allocator_t* allocator =
-      iree_hal_hsa_allocator_cast(base_allocator);
+      iree_hal_hsa_allocator_cast((iree_hal_allocator_t*)base_allocator);
 
   hsa_device_type_t type;
   hsa_status_t status = allocator->symbols->hsa_agent_get_info(
@@ -181,7 +181,7 @@ iree_status_t iree_hal_hsa_allocator_create(
               : "no CONCURRENT_MANAGED_ACCESS (expect slow accesses on "
                 "device-local + host-visible memory)");
 
-  iree_hal_hsa_allocator_t* allocator = NULL;
+  iree_hal_hsa_allocator_t* allocator = nullptr;
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_allocator_malloc(host_allocator, sizeof(*allocator),
                                 (void**)&allocator));
@@ -455,8 +455,8 @@ static iree_status_t iree_hal_hsa_allocator_allocate_buffer(
 
   iree_status_t status = iree_ok_status();
   iree_hal_hsa_buffer_type_t buffer_type = IREE_HAL_HSA_BUFFER_TYPE_DEVICE;
-  void* host_ptr = NULL;
-  hsa_device_pointer_t device_ptr = NULL;
+  void* host_ptr = nullptr;
+  hsa_device_pointer_t device_ptr = nullptr;
   IREE_TRACE_ZONE_BEGIN_NAMED(z0, "iree_hal_hsa_buffer_allocate");
   IREE_TRACE_ZONE_APPEND_VALUE_I64(z0, allocation_size);
 
@@ -510,7 +510,7 @@ static iree_status_t iree_hal_hsa_allocator_allocate_buffer(
   }
   IREE_TRACE_ZONE_END(z0);
 
-  iree_hal_buffer_t* buffer = NULL;
+  iree_hal_buffer_t* buffer = nullptr;
   if (iree_status_is_ok(status)) {
     status = iree_hal_hsa_buffer_wrap(
         base_allocator, compat_params.type, compat_params.access,
@@ -610,8 +610,8 @@ static iree_status_t iree_hal_hsa_allocator_import_buffer(
 
   iree_status_t status = iree_ok_status();
   iree_hal_hsa_buffer_type_t buffer_type = IREE_HAL_HSA_BUFFER_TYPE_DEVICE;
-  void* host_ptr = NULL;
-  hsa_device_pointer_t device_ptr = NULL;
+  void* host_ptr = nullptr;
+  hsa_device_pointer_t device_ptr = nullptr;
 
   switch (external_buffer->type) {
     case IREE_HAL_EXTERNAL_BUFFER_TYPE_HOST_ALLOCATION: {
@@ -621,7 +621,7 @@ static iree_status_t iree_hal_hsa_allocator_import_buffer(
           allocator->symbols,
           hsa_amd_memory_lock_to_pool(host_ptr, external_buffer->size,
                                       &allocator->cpu_agent, num_agents,
-                                      allocator->cpu_pool, flags, device_ptr),
+                                      allocator->cpu_pool, flags, &device_ptr),
           "hsa_amd_memory_lock_to_pool");
 
       break;
@@ -638,7 +638,7 @@ static iree_status_t iree_hal_hsa_allocator_import_buffer(
                               "external buffer type not supported");
   }
 
-  iree_hal_buffer_t* buffer = NULL;
+  iree_hal_buffer_t* buffer = nullptr;
   if (iree_status_is_ok(status)) {
     status = iree_hal_hsa_buffer_wrap(
         base_allocator, compat_params.type, compat_params.access,
diff --git a/runtime/src/iree-amd-aie/driver/hsa/hsa_buffer.c b/runtime/src/iree-amd-aie/driver/hsa/hsa_buffer.cc
similarity index 93%
rename from runtime/src/iree-amd-aie/driver/hsa/hsa_buffer.c
rename to runtime/src/iree-amd-aie/driver/hsa/hsa_buffer.cc
index f4b79fc948..6b6bc0d9a9 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/hsa_buffer.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/hsa_buffer.cc
@@ -7,22 +7,23 @@
 
 #include "iree-amd-aie/driver/hsa/hsa_buffer.h"
 
-#include <stddef.h>
-#include <stdint.h>
-#include <string.h>
+#include <cstdint>
+#include <cstring>
 
 #include "iree/base/api.h"
 #include "iree/base/tracing.h"
 
-typedef struct iree_hal_hsa_buffer_t {
+using iree_hal_hsa_buffer_t = struct iree_hal_hsa_buffer_t {
   iree_hal_buffer_t base;
   iree_hal_hsa_buffer_type_t type;
   void* host_ptr;
   hsa_device_pointer_t device_ptr;
   iree_hal_buffer_release_callback_t release_callback;
-} iree_hal_hsa_buffer_t;
+};
 
-static const iree_hal_buffer_vtable_t iree_hal_hsa_buffer_vtable;
+namespace {
+extern const iree_hal_buffer_vtable_t iree_hal_hsa_buffer_vtable;
+}
 
 static iree_hal_hsa_buffer_t* iree_hal_hsa_buffer_cast(
     iree_hal_buffer_t* base_value) {
@@ -54,7 +55,7 @@ iree_status_t iree_hal_hsa_buffer_wrap(
 
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  iree_hal_hsa_buffer_t* buffer = NULL;
+  iree_hal_hsa_buffer_t* buffer = nullptr;
   iree_status_t status =
       iree_allocator_malloc(host_allocator, sizeof(*buffer), (void**)&buffer);
   if (iree_status_is_ok(status)) {
@@ -157,11 +158,6 @@ void* iree_hal_hsa_buffer_host_pointer(const iree_hal_buffer_t* base_buffer) {
   return buffer->host_ptr;
 }
 
-void iree_hal_hsa_buffer_drop_release_callback(iree_hal_buffer_t* base_buffer) {
-  iree_hal_hsa_buffer_t* buffer = iree_hal_hsa_buffer_cast(base_buffer);
-  buffer->release_callback = iree_hal_buffer_release_callback_null();
-}
-
 static const iree_hal_buffer_vtable_t iree_hal_hsa_buffer_vtable = {
     .recycle = iree_hal_buffer_recycle,
     .destroy = iree_hal_hsa_buffer_destroy,
diff --git a/runtime/src/iree-amd-aie/driver/hsa/hsa_device.c b/runtime/src/iree-amd-aie/driver/hsa/hsa_device.cc
similarity index 96%
rename from runtime/src/iree-amd-aie/driver/hsa/hsa_device.c
rename to runtime/src/iree-amd-aie/driver/hsa/hsa_device.cc
index d98e9ef7e9..737453ec46 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/hsa_device.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/hsa_device.cc
@@ -7,15 +7,14 @@
 
 #include "iree-amd-aie/driver/hsa/hsa_device.h"
 
-#include <stddef.h>
-#include <stdint.h>
-#include <string.h>
+#include <cstddef>
+#include <cstdint>
+#include <cstring>
 
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 #include "iree-amd-aie/driver/hsa/event_pool.h"
 #include "iree-amd-aie/driver/hsa/event_semaphore.h"
 #include "iree-amd-aie/driver/hsa/hsa_allocator.h"
-#include "iree-amd-aie/driver/hsa/hsa_buffer.h"
 #include "iree-amd-aie/driver/hsa/nop_executable_cache.h"
 #include "iree-amd-aie/driver/hsa/pending_queue_actions.h"
 #include "iree-amd-aie/driver/hsa/queue_command_buffer.h"
@@ -77,7 +76,9 @@ typedef struct iree_hal_hsa_device_t {
   iree_hal_allocator_t* device_allocator;
 } iree_hal_hsa_device_t;
 
-static const iree_hal_device_vtable_t iree_hal_hsa_device_vtable;
+namespace {
+extern const iree_hal_device_vtable_t iree_hal_hsa_device_vtable;
+}
 
 static iree_hal_hsa_device_t* iree_hal_hsa_device_cast(
     iree_hal_device_t* base_value) {
@@ -85,11 +86,6 @@ static iree_hal_hsa_device_t* iree_hal_hsa_device_cast(
   return (iree_hal_hsa_device_t*)base_value;
 }
 
-static iree_hal_hsa_device_t* iree_hal_hsa_device_cast_unsafe(
-    iree_hal_device_t* base_value) {
-  return (iree_hal_hsa_device_t*)base_value;
-}
-
 IREE_API_EXPORT void iree_hal_hsa_device_params_initialize(
     iree_hal_hsa_device_params_t* out_params) {
   memset(out_params, 0, sizeof(*out_params));
@@ -117,7 +113,7 @@ static iree_status_t iree_hal_hsa_device_create_internal(
     const iree_hal_hsa_device_params_t* params, hsa_agent_t agent,
     hsa_queue_t* dispatch_queue, const iree_hal_hsa_dynamic_symbols_t* symbols,
     iree_allocator_t host_allocator, iree_hal_device_t** out_device) {
-  iree_hal_hsa_device_t* device = NULL;
+  iree_hal_hsa_device_t* device = nullptr;
   iree_host_size_t total_size = iree_sizeof_struct(*device) + identifier.size;
   IREE_RETURN_IF_ERROR(
       iree_allocator_malloc(host_allocator, total_size, (void**)&device));
@@ -168,8 +164,8 @@ iree_status_t iree_hal_hsa_device_create(
 
   size_t num_queue_packets = 1024;
   hsa_queue_type_t queue_type = HSA_QUEUE_TYPE_MULTI;
-  void* callback = NULL;
-  void* data = NULL;
+  void (*callback)(hsa_status_t, hsa_queue_t*, void*) = nullptr;
+  void* data = nullptr;
   uint32_t private_segment_size = 0;
   uint32_t group_segment_size = 0;
   hsa_queue_t* dispatch_queue;
@@ -185,20 +181,20 @@ iree_status_t iree_hal_hsa_device_create(
                                                agent, dispatch_queue, symbols,
                                                host_allocator, out_device);
 
-  iree_event_pool_t* host_event_pool = NULL;
+  iree_event_pool_t* host_event_pool = nullptr;
   if (iree_status_is_ok(status)) {
     status = iree_event_pool_allocate(params->event_pool_capacity,
                                       host_allocator, &host_event_pool);
   }
 
-  iree_hal_hsa_event_pool_t* device_event_pool = NULL;
+  iree_hal_hsa_event_pool_t* device_event_pool = nullptr;
   if (iree_status_is_ok(status)) {
     status =
         iree_hal_hsa_event_pool_allocate(symbols, params->event_pool_capacity,
                                          host_allocator, &device_event_pool);
   }
 
-  iree_hal_hsa_timepoint_pool_t* timepoint_pool = NULL;
+  iree_hal_hsa_timepoint_pool_t* timepoint_pool = nullptr;
   if (iree_status_is_ok(status)) {
     status = iree_hal_hsa_timepoint_pool_allocate(
         host_event_pool, device_event_pool, params->event_pool_capacity,
@@ -223,12 +219,6 @@ iree_status_t iree_hal_hsa_device_create(
   return status;
 }
 
-const iree_hal_hsa_dynamic_symbols_t* iree_hal_hsa_device_dynamic_symbols(
-    iree_hal_device_t* base_device) {
-  iree_hal_hsa_device_t* device = iree_hal_hsa_device_cast_unsafe(base_device);
-  return device->hsa_symbols;
-}
-
 static void iree_hal_hsa_device_destroy(iree_hal_device_t* base_device) {
   iree_hal_hsa_device_t* device = iree_hal_hsa_device_cast(base_device);
   iree_allocator_t host_allocator = iree_hal_device_host_allocator(base_device);
diff --git a/runtime/src/iree-amd-aie/driver/hsa/hsa_device.h b/runtime/src/iree-amd-aie/driver/hsa/hsa_device.h
index d6988f9238..13c584ec3c 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/hsa_device.h
+++ b/runtime/src/iree-amd-aie/driver/hsa/hsa_device.h
@@ -33,7 +33,7 @@ iree_status_t iree_hal_hsa_device_create_queue_command_buffer(
     iree_hal_command_buffer_t** out_command_buffer);
 
 // Returns the dynamic symbol table from the |device| if it is a HSA device
-// and otherwise returns NULL.
+// and otherwise returns nullptr.
 //
 // WARNING: the symbols are only valid for as long as the device is. Hosting
 // libraries and applications should prefer to either link against HSA
diff --git a/runtime/src/iree-amd-aie/driver/hsa/hsa_driver.c b/runtime/src/iree-amd-aie/driver/hsa/hsa_driver.cc
similarity index 98%
rename from runtime/src/iree-amd-aie/driver/hsa/hsa_driver.c
rename to runtime/src/iree-amd-aie/driver/hsa/hsa_driver.cc
index a4090d3f41..0b8daa8c41 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/hsa_driver.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/hsa_driver.cc
@@ -5,8 +5,8 @@
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
-#include <stdint.h>
-#include <string.h>
+#include <cstdint>
+#include <cstring>
 
 #include "iree-amd-aie/driver/hsa/api.h"
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
@@ -51,9 +51,6 @@ typedef struct iree_hal_hsa_driver_t {
   hsa_agent_t agents[IREE_HAL_HSA_MAX_DEVICES];
 } iree_hal_hsa_driver_t;
 
-typedef struct iree_hal_hsa_device_info_t {
-} iree_hal_hsa_device_info_t;
-
 // A struct encapsulating common variables we need while communicating with HSA
 // callbacks
 typedef struct iree_hal_hsa_callback_package_t {
@@ -62,7 +59,9 @@ typedef struct iree_hal_hsa_callback_package_t {
   void* return_value;
 } iree_hal_hsa_callback_package_t;
 
-static const iree_hal_driver_vtable_t iree_hal_hsa_driver_vtable;
+namespace {
+extern const iree_hal_driver_vtable_t iree_hal_hsa_driver_vtable;
+}
 
 static iree_hal_hsa_driver_t* iree_hal_hsa_driver_cast(
     iree_hal_driver_t* base_value) {
@@ -142,7 +141,7 @@ static iree_status_t iree_hal_hsa_driver_create_internal(
     iree_string_view_t identifier, const iree_hal_hsa_driver_options_t* options,
     const iree_hal_hsa_device_params_t* device_params,
     iree_allocator_t host_allocator, iree_hal_driver_t** out_driver) {
-  iree_hal_hsa_driver_t* driver = NULL;
+  iree_hal_hsa_driver_t* driver = nullptr;
   iree_host_size_t total_size = iree_sizeof_struct(*driver) + identifier.size;
   IREE_RETURN_IF_ERROR(
       iree_allocator_malloc(host_allocator, total_size, (void**)&driver));
@@ -244,7 +243,9 @@ static iree_status_t get_hsa_agent_uuid(iree_hal_hsa_dynamic_symbols_t* syms,
                                         char* out_device_uuid) {
   // `HSA_AMD_AGENT_INFO_UUID` is part of the `hsa_amd_agent_info_t`
   // However, hsa_agent_get_info expects a hsa_agent_info_t.
-  hsa_agent_info_t uuid_info = (int)HSA_AMD_AGENT_INFO_UUID;
+  // TODO(max): this should be updated to use hsa_amd_agent_info_s
+  hsa_agent_info_t uuid_info =
+      static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_UUID);
   IREE_HSA_RETURN_IF_ERROR(
       syms, hsa_agent_get_info(agent, uuid_info, out_device_uuid),
       "hsa_agent_get_info");
@@ -320,7 +321,7 @@ static iree_status_t iree_hal_hsa_driver_query_available_devices(
   int device_count = driver->num_gpu_agents;
 
   // Allocate the return infos and populate with the devices.
-  iree_hal_device_info_t* device_infos = NULL;
+  iree_hal_device_info_t* device_infos = nullptr;
   iree_host_size_t total_size =
       device_count * (sizeof(iree_hal_device_info_t) +
                       IREE_HAL_HSA_MAX_DEVICE_NAME_LENGTH * sizeof(char));
@@ -365,7 +366,7 @@ static iree_status_t iree_hal_hsa_driver_select_default_device(
     iree_hal_driver_t* base_driver, iree_hal_hsa_dynamic_symbols_t* syms,
     int default_device_index, iree_allocator_t host_allocator,
     hsa_agent_t* out_device) {
-  iree_hal_device_info_t* device_infos = NULL;
+  iree_hal_device_info_t* device_infos = nullptr;
   iree_host_size_t device_count = 0;
   IREE_RETURN_IF_ERROR(iree_hal_hsa_driver_query_available_devices(
       base_driver, host_allocator, &device_count, &device_infos));
diff --git a/runtime/src/iree-amd-aie/driver/hsa/native_executable.c b/runtime/src/iree-amd-aie/driver/hsa/native_executable.cc
similarity index 96%
rename from runtime/src/iree-amd-aie/driver/hsa/native_executable.c
rename to runtime/src/iree-amd-aie/driver/hsa/native_executable.cc
index 777a13bd2b..75a0f3b34b 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/native_executable.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/native_executable.cc
@@ -7,7 +7,10 @@
 
 #include "iree-amd-aie/driver/hsa/native_executable.h"
 
-#include <stddef.h>
+#include <cstddef>
+#include <fstream>
+#include <iostream>
+#include <vector>
 
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 #include "iree-amd-aie/driver/hsa/status_util.h"
@@ -38,7 +41,9 @@ typedef struct iree_hal_hsa_native_executable_t {
 } iree_hal_hsa_native_executable_t;
 // + Additional inline allocation for holding entry point information.
 
-static const iree_hal_executable_vtable_t iree_hal_hsa_native_executable_vtable;
+namespace {
+extern const iree_hal_executable_vtable_t iree_hal_hsa_native_executable_vtable;
+}
 
 static iree_hal_hsa_native_executable_t* iree_hal_hsa_native_executable_cast(
     iree_hal_executable_t* base_value) {
@@ -161,8 +166,8 @@ iree_status_t iree_hal_hsa_native_executable_create(
   IREE_ASSERT_ARGUMENT(out_executable);
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  *out_executable = NULL;
-  iree_hal_hsa_native_executable_t* executable = NULL;
+  *out_executable = nullptr;
+  iree_hal_hsa_native_executable_t* executable = nullptr;
 
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_hal_hsa_native_executable_flatbuffer_verify(
@@ -233,19 +238,20 @@ iree_status_t iree_hal_hsa_native_executable_create(
   status = IREE_HSA_RESULT_TO_STATUS(
       symbols, hsa_executable_create_alt(
                    HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
-                   NULL, &hsa_executable));
+                   nullptr, &hsa_executable));
   if (!iree_status_is_ok(status)) {
     return status;
   }
   status = IREE_HSA_RESULT_TO_STATUS(
-      symbols, hsa_executable_load_agent_code_object(
-                   hsa_executable, agent, code_object_reader, NULL, NULL));
+      symbols,
+      hsa_executable_load_agent_code_object(
+          hsa_executable, agent, code_object_reader, nullptr, nullptr));
   if (!iree_status_is_ok(status)) {
     return status;
   }
 
   status = IREE_HSA_RESULT_TO_STATUS(
-      symbols, hsa_executable_freeze(hsa_executable, NULL));
+      symbols, hsa_executable_freeze(hsa_executable, nullptr));
   if (!iree_status_is_ok(status)) {
     return status;
   }
@@ -263,7 +269,7 @@ iree_status_t iree_hal_hsa_native_executable_create(
       iree_string_view_t name_view = iree_make_cstring_view(entry_name);
       iree_string_view_t suffix_view = iree_make_cstring_view(".kd");
       iree_host_size_t total_length = name_view.size + suffix_view.size;
-      char* kd_entry_name = NULL;
+      char* kd_entry_name = nullptr;
       IREE_RETURN_AND_END_ZONE_IF_ERROR(
           z0, iree_allocator_malloc(host_allocator, total_length + 1,
                                     (void**)&kd_entry_name));
diff --git a/runtime/src/iree-amd-aie/driver/hsa/native_executable.h b/runtime/src/iree-amd-aie/driver/hsa/native_executable.h
index afdeb1fdb9..34eecc4a35 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/native_executable.h
+++ b/runtime/src/iree-amd-aie/driver/hsa/native_executable.h
@@ -8,7 +8,7 @@
 #ifndef IREE_EXPERIMENTAL_HSA_NATIVE_EXECUTABLE_H_
 #define IREE_EXPERIMENTAL_HSA_NATIVE_EXECUTABLE_H_
 
-#include <stdint.h>
+#include <cstdint>
 
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 #include "iree-amd-aie/driver/hsa/hsa_headers.h"
diff --git a/runtime/src/iree-amd-aie/driver/hsa/nop_executable_cache.c b/runtime/src/iree-amd-aie/driver/hsa/nop_executable_cache.cc
similarity index 94%
rename from runtime/src/iree-amd-aie/driver/hsa/nop_executable_cache.c
rename to runtime/src/iree-amd-aie/driver/hsa/nop_executable_cache.cc
index 31d5e111de..e0437ff667 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/nop_executable_cache.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/nop_executable_cache.cc
@@ -7,10 +7,6 @@
 
 #include "iree-amd-aie/driver/hsa/nop_executable_cache.h"
 
-#include <stdbool.h>
-#include <stddef.h>
-
-#include "iree-amd-aie/driver/hsa/hsa_allocator.h"
 #include "iree-amd-aie/driver/hsa/native_executable.h"
 #include "iree/base/api.h"
 #include "iree/base/tracing.h"
@@ -30,8 +26,10 @@ typedef struct iree_hal_hsa_nop_executable_cache_t {
 
 } iree_hal_hsa_nop_executable_cache_t;
 
-static const iree_hal_executable_cache_vtable_t
+namespace {
+extern const iree_hal_executable_cache_vtable_t
     iree_hal_hsa_nop_executable_cache_vtable;
+}
 
 static iree_hal_hsa_nop_executable_cache_t*
 iree_hal_hsa_nop_executable_cache_cast(
@@ -49,8 +47,8 @@ iree_status_t iree_hal_hsa_nop_executable_cache_create(
   IREE_ASSERT_ARGUMENT(out_executable_cache);
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  *out_executable_cache = NULL;
-  iree_hal_hsa_nop_executable_cache_t* executable_cache = NULL;
+  *out_executable_cache = nullptr;
+  iree_hal_hsa_nop_executable_cache_t* executable_cache = nullptr;
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_allocator_malloc(host_allocator, sizeof(*executable_cache),
                                 (void**)&executable_cache));
diff --git a/runtime/src/iree-amd-aie/driver/hsa/pending_queue_actions.c b/runtime/src/iree-amd-aie/driver/hsa/pending_queue_actions.cc
similarity index 96%
rename from runtime/src/iree-amd-aie/driver/hsa/pending_queue_actions.c
rename to runtime/src/iree-amd-aie/driver/hsa/pending_queue_actions.cc
index 49a68f92fa..8708712070 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/pending_queue_actions.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/pending_queue_actions.cc
@@ -7,8 +7,7 @@
 
 #include "iree-amd-aie/driver/hsa/pending_queue_actions.h"
 
-#include <stdbool.h>
-#include <stddef.h>
+#include <cstddef>
 
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 #include "iree-amd-aie/driver/hsa/event_semaphore.h"
@@ -63,7 +62,7 @@ typedef struct iree_hal_hsa_queue_action_t {
   // destruction.
   iree_hal_hsa_queue_action_state_t state;
   // The callback to run after completing this action and before freeing
-  // all resources. Can be NULL.
+  // all resources. Can be nullptr.
   iree_hal_hsa_pending_action_cleanup_callback_t cleanup_callback;
   // User data to pass into the callback.
   void* callback_user_data;
@@ -110,7 +109,7 @@ typedef struct iree_hal_hsa_queue_action_list_t {
 // Returns true if the action list is empty.
 static inline bool iree_hal_hsa_queue_action_list_is_empty(
     const iree_hal_hsa_queue_action_list_t* list) {
-  return list->head == NULL;
+  return list->head == nullptr;
 }
 
 // Pushes |action| on to the end of the given action |list|.
@@ -122,7 +121,7 @@ static void iree_hal_hsa_queue_action_list_push_back(
   } else {
     list->head = action;
   }
-  action->next = NULL;
+  action->next = nullptr;
   action->prev = list->tail;
   list->tail = action;
 }
@@ -135,13 +134,13 @@ static void iree_hal_hsa_queue_action_list_erase(
   iree_hal_hsa_queue_action_t* prev = action->prev;
   if (prev) {
     prev->next = next;
-    action->prev = NULL;
+    action->prev = nullptr;
   } else {
     list->head = next;
   }
   if (next) {
     next->prev = prev;
-    action->next = NULL;
+    action->next = nullptr;
   } else {
     list->tail = prev;
   }
@@ -154,14 +153,14 @@ static void iree_hal_hsa_queue_action_list_take_all(
   IREE_ASSERT_NE(available_list, ready_list);
   ready_list->head = available_list->head;
   ready_list->tail = available_list->tail;
-  available_list->head = NULL;
-  available_list->tail = NULL;
+  available_list->head = nullptr;
+  available_list->tail = nullptr;
 }
 
 // Frees all actions in the given |list|.
 static void iree_hal_hsa_queue_action_list_free_actions(
     iree_allocator_t host_allocator, iree_hal_hsa_queue_action_list_t* list) {
-  for (iree_hal_hsa_queue_action_t* action = list->head; action != NULL;) {
+  for (iree_hal_hsa_queue_action_t* action = list->head; action != nullptr;) {
     iree_hal_hsa_queue_action_t* next_action = action->next;
     iree_allocator_free(host_allocator, action);
     action = next_action;
@@ -280,8 +279,10 @@ struct iree_hal_hsa_pending_queue_actions_t {
   iree_hal_hsa_working_area_t working_area;
 };
 
-static const iree_hal_resource_vtable_t
+namespace {
+extern const iree_hal_resource_vtable_t
     iree_hal_hsa_pending_queue_actions_vtable;
+}
 
 iree_status_t iree_hal_hsa_pending_queue_actions_create(
     const iree_hal_hsa_dynamic_symbols_t* symbols,
@@ -292,7 +293,7 @@ iree_status_t iree_hal_hsa_pending_queue_actions_create(
   IREE_ASSERT_ARGUMENT(out_actions);
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  iree_hal_hsa_pending_queue_actions_t* actions = NULL;
+  iree_hal_hsa_pending_queue_actions_t* actions = nullptr;
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_allocator_malloc(host_allocator, sizeof(*actions),
                                 (void**)&actions));
@@ -382,7 +383,7 @@ static iree_status_t iree_hal_hsa_copy_command_buffer_list(
     iree_host_size_t command_buffer_count,
     iree_hal_command_buffer_t* const* in_list, iree_allocator_t host_allocator,
     iree_hal_command_buffer_t*** out_list) {
-  *out_list = NULL;
+  *out_list = nullptr;
   if (!command_buffer_count) return iree_ok_status();
 
   iree_host_size_t total_size = command_buffer_count * sizeof(*in_list);
@@ -440,7 +441,7 @@ iree_status_t iree_hal_hsa_pending_queue_actions_enqueue_execution(
   IREE_ASSERT_ARGUMENT(command_buffer_count == 0 || command_buffers);
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  iree_hal_hsa_queue_action_t* action = NULL;
+  iree_hal_hsa_queue_action_t* action = nullptr;
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_allocator_malloc(actions->host_allocator, sizeof(*action),
                                 (void**)&action));
@@ -458,7 +459,7 @@ iree_status_t iree_hal_hsa_pending_queue_actions_enqueue_execution(
   action->is_pending = true;
 
   // Retain all command buffers and semaphores.
-  iree_hal_resource_set_t* resource_set = NULL;
+  iree_hal_resource_set_t* resource_set = nullptr;
   iree_status_t status =
       iree_hal_resource_set_allocate(actions->block_pool, &resource_set);
   if (IREE_LIKELY(iree_status_is_ok(status))) {
@@ -552,7 +553,9 @@ static bool iree_hal_hsa_execution_device_signal_host_callback(
   iree_hal_hsa_worker_state_t prev_state =
       IREE_HAL_HSA_WORKER_STATE_IDLE_WAITING;
   iree_atomic_compare_exchange_strong_int32(
-      &actions->working_area.worker_state, /*expected=*/&prev_state,
+      &actions->working_area.worker_state,
+      // TODO(max): shouldn't these by int64 since they're ptrs?
+      /*expected=*/reinterpret_cast<int32_t*>(&prev_state),
       /*desired=*/IREE_HAL_HSA_WORKER_STATE_WORKLOAD_PENDING,
       /*order_succ=*/iree_memory_order_acq_rel,
       /*order_fail=*/iree_memory_order_acquire);
@@ -592,7 +595,7 @@ static iree_status_t iree_hal_hsa_pending_queue_actions_issue_execution(
   for (iree_host_size_t i = 0; i < action->payload.command_buffers.count; ++i) {
     iree_hal_command_buffer_t* command_buffer =
         action->payload.command_buffers.ptr[i];
-    iree_hal_command_buffer_t* queue_command_buffer = NULL;
+    iree_hal_command_buffer_t* queue_command_buffer = nullptr;
     iree_hal_command_buffer_mode_t mode =
         IREE_HAL_COMMAND_BUFFER_MODE_ONE_SHOT |
         IREE_HAL_COMMAND_BUFFER_MODE_ALLOW_INLINE_EXECUTION |
@@ -700,8 +703,8 @@ iree_status_t iree_hal_hsa_pending_queue_actions_issue(
     iree_hal_hsa_pending_queue_actions_t* actions) {
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  iree_hal_hsa_queue_action_list_t pending_list = {NULL, NULL};
-  iree_hal_hsa_queue_action_list_t ready_list = {NULL, NULL};
+  iree_hal_hsa_queue_action_list_t pending_list = {nullptr, nullptr};
+  iree_hal_hsa_queue_action_list_t ready_list = {nullptr, nullptr};
 
   iree_slim_mutex_lock(&actions->action_mutex);
 
@@ -716,7 +719,7 @@ iree_status_t iree_hal_hsa_pending_queue_actions_issue(
   iree_hal_hsa_queue_action_t* action = actions->action_list.head;
   while (action) {
     iree_hal_hsa_queue_action_t* next_action = action->next;
-    action->next = NULL;
+    action->next = nullptr;
 
     iree_host_size_t semaphore_count = action->wait_semaphore_list.count;
     iree_hal_semaphore_t** semaphores = action->wait_semaphore_list.semaphores;
@@ -779,13 +782,13 @@ iree_status_t iree_hal_hsa_pending_queue_actions_issue(
 
   iree_slim_mutex_unlock(&actions->action_mutex);
 
-  if (ready_list.head == NULL) {
+  if (ready_list.head == nullptr) {
     // Nothing ready yet. Just return.
     IREE_TRACE_ZONE_END(z0);
     return status;
   }
 
-  iree_hal_hsa_atomic_slist_entry_t* entry = NULL;
+  iree_hal_hsa_atomic_slist_entry_t* entry = nullptr;
   // TODO: avoid host allocator malloc; use some pool for the allocation.
   if (iree_status_is_ok(status)) {
     status = iree_allocator_malloc(actions->host_allocator, sizeof(*entry),
@@ -813,7 +816,8 @@ iree_status_t iree_hal_hsa_pending_queue_actions_issue(
   iree_hal_hsa_worker_state_t prev_state =
       IREE_HAL_HSA_WORKER_STATE_IDLE_WAITING;
   iree_atomic_compare_exchange_strong_int32(
-      &actions->working_area.worker_state, /*expected=*/&prev_state,
+      &actions->working_area.worker_state,
+      /*expected=*/reinterpret_cast<int32_t*>(&prev_state),
       /*desired=*/IREE_HAL_HSA_WORKER_STATE_WORKLOAD_PENDING,
       /*order_succ=*/iree_memory_order_acq_rel,
       /*order_fail=*/iree_memory_order_acquire);
@@ -822,8 +826,9 @@ iree_status_t iree_hal_hsa_pending_queue_actions_issue(
 
   // Handle potential error cases from the worker thread.
   if (prev_state == IREE_HAL_HSA_WORKER_STATE_EXIT_ERROR) {
-    iree_status_code_t code = iree_atomic_load_int32(
-        &actions->working_area.error_code, iree_memory_order_acquire);
+    iree_status_code_t code =
+        static_cast<iree_status_code_t>(iree_atomic_load_int32(
+            &actions->working_area.error_code, iree_memory_order_acquire));
     status = iree_status_from_code(code);
   }
 
@@ -837,8 +842,9 @@ iree_status_t iree_hal_hsa_pending_queue_actions_issue(
 
 static bool iree_hal_hsa_worker_has_incoming_request(
     iree_hal_hsa_working_area_t* working_area) {
-  iree_hal_hsa_worker_state_t value = iree_atomic_load_int32(
-      &working_area->worker_state, iree_memory_order_acquire);
+  iree_hal_hsa_worker_state_t value =
+      static_cast<iree_hal_hsa_worker_state_t>(iree_atomic_load_int32(
+          &working_area->worker_state, iree_memory_order_acquire));
   // These are the only two possible states that set from the main thread to
   // the worker thread.
   return value == IREE_HAL_HSA_WORKER_STATE_WORKLOAD_PENDING ||
@@ -868,7 +874,7 @@ static iree_status_t iree_hal_hsa_worker_process_ready_list(
     iree_hal_hsa_queue_action_t* action = entry->ready_list_head;
     while (action) {
       iree_hal_hsa_queue_action_t* next_action = action->next;
-      action->next = NULL;
+      action->next = nullptr;
 
       switch (action->state) {
         case IREE_HAL_HSA_QUEUE_ACTION_STATE_ALIVE:
@@ -912,7 +918,8 @@ static int iree_hal_hsa_worker_execute(
     iree_hal_hsa_worker_state_t prev_state =
         IREE_HAL_HSA_WORKER_STATE_WORKLOAD_PENDING;
     iree_atomic_compare_exchange_strong_int32(
-        &working_area->worker_state, /*expected=*/&prev_state,
+        &working_area->worker_state,
+        /*expected=*/reinterpret_cast<int32_t*>(&prev_state),
         /*desired=*/IREE_HAL_HSA_WORKER_STATE_IDLE_WAITING,
         /*order_succ=*/iree_memory_order_acq_rel,
         /*order_fail=*/iree_memory_order_acquire);
diff --git a/runtime/src/iree-amd-aie/driver/hsa/pending_queue_actions.h b/runtime/src/iree-amd-aie/driver/hsa/pending_queue_actions.h
index 54f274fce4..88b8089810 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/pending_queue_actions.h
+++ b/runtime/src/iree-amd-aie/driver/hsa/pending_queue_actions.h
@@ -40,7 +40,7 @@ typedef void(IREE_API_PTR* iree_hal_hsa_pending_action_cleanup_callback_t)(
 // Enqueues the given list of |command_buffers| that waits on
 // |wait_semaphore_list| and signals |signal_semaphore_lsit|.
 //
-// |cleanup_callback|, if not NULL, will run after the action completes but
+// |cleanup_callback|, if not nullptr, will run after the action completes but
 // before releasing all retained resources.
 iree_status_t iree_hal_hsa_pending_queue_actions_enqueue_execution(
     iree_hal_device_t* device, hsa_queue_t* dispatch_queue,
diff --git a/runtime/src/iree-amd-aie/driver/hsa/queue_command_buffer.c b/runtime/src/iree-amd-aie/driver/hsa/queue_command_buffer.cc
similarity index 98%
rename from runtime/src/iree-amd-aie/driver/hsa/queue_command_buffer.c
rename to runtime/src/iree-amd-aie/driver/hsa/queue_command_buffer.cc
index e2f0828629..98d2a94f5f 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/queue_command_buffer.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/queue_command_buffer.cc
@@ -47,8 +47,10 @@ typedef struct iree_hal_hsa_queue_command_buffer_t {
   } descriptor_sets[IREE_HAL_HSA_MAX_DESCRIPTOR_SET_COUNT];
 } iree_hal_hsa_queue_command_buffer_t;
 
-static const iree_hal_command_buffer_vtable_t
+namespace {
+extern const iree_hal_command_buffer_vtable_t
     iree_hal_hsa_queue_command_buffer_vtable;
+}
 
 static iree_hal_hsa_queue_command_buffer_t*
 iree_hal_hsa_queue_command_buffer_cast(iree_hal_command_buffer_t* base_value) {
@@ -68,7 +70,7 @@ iree_status_t iree_hal_hsa_queue_command_buffer_create(
   IREE_ASSERT_ARGUMENT(device);
   IREE_ASSERT_ARGUMENT(hsa_symbols);
   IREE_ASSERT_ARGUMENT(out_command_buffer);
-  *out_command_buffer = NULL;
+  *out_command_buffer = nullptr;
 
   if (binding_capacity > 0) {
     return iree_make_status(IREE_STATUS_UNIMPLEMENTED,
@@ -77,7 +79,7 @@ iree_status_t iree_hal_hsa_queue_command_buffer_create(
 
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  iree_hal_hsa_queue_command_buffer_t* command_buffer = NULL;
+  iree_hal_hsa_queue_command_buffer_t* command_buffer = nullptr;
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_allocator_malloc(host_allocator, sizeof(*command_buffer),
                                 (void**)&command_buffer));
@@ -280,7 +282,7 @@ static iree_status_t iree_hal_hsa_queue_command_buffer_update_buffer(
   // operation and get the wrong data.
   const uint8_t* src = (const uint8_t*)source_buffer + source_offset;
   if (command_buffer->arena.block_pool) {
-    uint8_t* storage = NULL;
+    uint8_t* storage = nullptr;
     IREE_RETURN_AND_END_ZONE_IF_ERROR(
         z0, iree_arena_allocate(&command_buffer->arena, target_ref.length,
                                 (void**)&storage));
@@ -360,7 +362,7 @@ static iree_status_t iree_hal_hsa_queue_command_buffer_push_descriptor_set(
       command_buffer->descriptor_sets[set].bindings;
   for (iree_host_size_t i = 0; i < binding_count; i++) {
     const iree_hal_buffer_ref_t* binding = &bindings[i];
-    hsa_device_pointer_t device_ptr = NULL;
+    hsa_device_pointer_t device_ptr = nullptr;
     if (binding->buffer) {
       IREE_RETURN_AND_END_ZONE_IF_ERROR(
           z0, iree_hal_resource_set_insert(command_buffer->resource_set, 1,
@@ -424,7 +426,7 @@ static iree_status_t iree_hal_hsa_queue_command_buffer_dispatch(
   };
 
   iree_device_size_t kern_arg_allocation_size = total_size;
-  iree_hal_buffer_t* kern_arg_allocation_buffer = NULL;
+  iree_hal_buffer_t* kern_arg_allocation_buffer = nullptr;
   iree_status_t result = iree_hal_allocator_allocate_buffer(
       command_buffer->device_allocator, buffer_params, kern_arg_allocation_size,
       &kern_arg_allocation_buffer);
@@ -483,7 +485,7 @@ static iree_status_t iree_hal_hsa_queue_command_buffer_dispatch(
 
   hsa_signal_value_t signal_value = 1;
   uint32_t num_consumers = 0;
-  const hsa_agent_t* consumers = NULL;
+  const hsa_agent_t* consumers = nullptr;
   iree_status_t status = IREE_HSA_RESULT_TO_STATUS(
       command_buffer->hsa_symbols,
       hsa_signal_create(signal_value, num_consumers, consumers,
diff --git a/runtime/src/iree-amd-aie/driver/hsa/queue_command_buffer.h b/runtime/src/iree-amd-aie/driver/hsa/queue_command_buffer.h
index e37e5c2a58..de8d02a90f 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/queue_command_buffer.h
+++ b/runtime/src/iree-amd-aie/driver/hsa/queue_command_buffer.h
@@ -10,7 +10,6 @@
 
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 #include "iree-amd-aie/driver/hsa/hsa_headers.h"
-// #include "iree-amd-aie/driver/hsa/tracing.h"
 #include "iree/base/internal/arena.h"
 #include "iree/hal/api.h"
 
@@ -21,14 +20,14 @@ extern "C" {
 // Creates command buffer that immediately issues commands against the given
 // HSA |stream|. Access to |stream| must be synchronized by the user.
 //
-// If |block_pool| is non-NULL then the stream command buffer will retain copies
-// of input data until reset. If NULL then the caller must ensure the lifetime
-// of input data outlives the command buffer.
+// If |block_pool| is non-nullptr then the stream command buffer will retain
+// copies of input data until reset. If nullptr then the caller must ensure the
+// lifetime of input data outlives the command buffer.
 //
 // This command buffer is used to replay deferred command buffers. When
 // replaying the scratch data required for things like buffer updates is
 // retained by the source deferred command buffer and as such the |block_pool|
-// and can be NULL to avoid a double copy.
+// and can be nullptr to avoid a double copy.
 iree_status_t iree_hal_hsa_queue_command_buffer_create(
     iree_hal_device_t* device,
     const iree_hal_hsa_dynamic_symbols_t* hsa_symbols,
diff --git a/runtime/src/iree-amd-aie/driver/hsa/status_util.c b/runtime/src/iree-amd-aie/driver/hsa/status_util.cc
similarity index 99%
rename from runtime/src/iree-amd-aie/driver/hsa/status_util.c
rename to runtime/src/iree-amd-aie/driver/hsa/status_util.cc
index 9676094409..bc2175fb26 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/status_util.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/status_util.cc
@@ -7,8 +7,6 @@
 
 #include "iree-amd-aie/driver/hsa/status_util.h"
 
-#include <stddef.h>
-
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 #include "iree/base/status.h"
 
@@ -178,7 +176,7 @@ iree_status_t iree_hal_hsa_result_to_status(
 
   const char* error_name = hsa_status_to_string(result);
 
-  const char* error_string = NULL;
+  const char* error_string = nullptr;
   hsa_status_t status_string_result =
       syms->hsa_status_string(result, &error_string);
   if (status_string_result != HSA_STATUS_SUCCESS) {
diff --git a/runtime/src/iree-amd-aie/driver/hsa/status_util.h b/runtime/src/iree-amd-aie/driver/hsa/status_util.h
index 4f20070788..84fa14a55f 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/status_util.h
+++ b/runtime/src/iree-amd-aie/driver/hsa/status_util.h
@@ -8,7 +8,7 @@
 #ifndef IREE_EXPERIMENTAL_HSA_STATUS_UTIL_H_
 #define IREE_EXPERIMENTAL_HSA_STATUS_UTIL_H_
 
-#include <stdint.h>
+#include <cstdint>
 
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 #include "iree/base/api.h"
diff --git a/runtime/src/iree-amd-aie/driver/hsa/timepoint_pool.c b/runtime/src/iree-amd-aie/driver/hsa/timepoint_pool.cc
similarity index 95%
rename from runtime/src/iree-amd-aie/driver/hsa/timepoint_pool.c
rename to runtime/src/iree-amd-aie/driver/hsa/timepoint_pool.cc
index 78d5c0a1a4..e489436b43 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/timepoint_pool.c
+++ b/runtime/src/iree-amd-aie/driver/hsa/timepoint_pool.cc
@@ -7,9 +7,8 @@
 
 #include "iree-amd-aie/driver/hsa/timepoint_pool.h"
 
-#include <stdbool.h>
-#include <stddef.h>
-#include <string.h>
+#include <cstddef>
+#include <cstring>
 
 #include "iree-amd-aie/driver/hsa/dynamic_symbols.h"
 #include "iree-amd-aie/driver/hsa/event_pool.h"
@@ -30,10 +29,10 @@ static iree_status_t iree_hal_hsa_timepoint_allocate(
     iree_hal_hsa_timepoint_t** out_timepoint) {
   IREE_ASSERT_ARGUMENT(pool);
   IREE_ASSERT_ARGUMENT(out_timepoint);
-  *out_timepoint = NULL;
+  *out_timepoint = nullptr;
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  iree_hal_hsa_timepoint_t* timepoint = NULL;
+  iree_hal_hsa_timepoint_t* timepoint = nullptr;
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_allocator_malloc(host_allocator, sizeof(*timepoint),
                                 (void**)&timepoint));
@@ -108,10 +107,10 @@ iree_status_t iree_hal_hsa_timepoint_pool_allocate(
   IREE_ASSERT_ARGUMENT(host_event_pool);
   IREE_ASSERT_ARGUMENT(device_event_pool);
   IREE_ASSERT_ARGUMENT(out_timepoint_pool);
-  *out_timepoint_pool = NULL;
+  *out_timepoint_pool = nullptr;
   IREE_TRACE_ZONE_BEGIN(z0);
 
-  iree_hal_hsa_timepoint_pool_t* timepoint_pool = NULL;
+  iree_hal_hsa_timepoint_pool_t* timepoint_pool = nullptr;
   iree_host_size_t total_size =
       sizeof(*timepoint_pool) +
       available_capacity * sizeof(*timepoint_pool->available_list);
@@ -219,8 +218,8 @@ iree_status_t iree_hal_hsa_timepoint_pool_acquire_host_wait(
 
   // Acquire host events to wrap up. This should happen before acquiring the
   // timepoints to avoid nested locks.
-  iree_event_t* host_events = iree_alloca(
-      timepoint_count * sizeof((*out_timepoints)->timepoint.host_wait));
+  iree_event_t* host_events = static_cast<iree_event_t*>(iree_alloca(
+      timepoint_count * sizeof((*out_timepoints)->timepoint.host_wait)));
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_event_pool_acquire(timepoint_pool->host_event_pool,
                                   timepoint_count, host_events));
@@ -245,8 +244,9 @@ iree_status_t iree_hal_hsa_timepoint_pool_acquire_device_signal(
 
   // Acquire device events to wrap up. This should happen before acquiring the
   // timepoints to avoid nested locks.
-  iree_hal_hsa_event_t** device_events = iree_alloca(
-      timepoint_count * sizeof((*out_timepoints)->timepoint.device_signal));
+  iree_hal_hsa_event_t** device_events = static_cast<iree_hal_hsa_event_t**>(
+      iree_alloca(timepoint_count *
+                  sizeof((*out_timepoints)->timepoint.device_signal)));
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_hal_hsa_event_pool_acquire(timepoint_pool->device_event_pool,
                                           timepoint_count, device_events));
@@ -271,8 +271,9 @@ iree_status_t iree_hal_hsa_timepoint_pool_acquire_device_wait(
 
   // Acquire device events to wrap up. This should happen before acquiring the
   // timepoints to avoid nested locks.
-  iree_hal_hsa_event_t** device_events = iree_alloca(
-      timepoint_count * sizeof((*out_timepoints)->timepoint.device_wait));
+  iree_hal_hsa_event_t** device_events =
+      static_cast<iree_hal_hsa_event_t**>(iree_alloca(
+          timepoint_count * sizeof((*out_timepoints)->timepoint.device_wait)));
   IREE_RETURN_AND_END_ZONE_IF_ERROR(
       z0, iree_hal_hsa_event_pool_acquire(timepoint_pool->device_event_pool,
                                           timepoint_count, device_events));
diff --git a/runtime/src/iree-amd-aie/driver/hsa/timepoint_pool.h b/runtime/src/iree-amd-aie/driver/hsa/timepoint_pool.h
index e1ad690121..08b5a97cc0 100644
--- a/runtime/src/iree-amd-aie/driver/hsa/timepoint_pool.h
+++ b/runtime/src/iree-amd-aie/driver/hsa/timepoint_pool.h
@@ -58,7 +58,7 @@ typedef struct iree_hal_hsa_timepoint_t {
   union {
     iree_event_t host_wait;
     iree_hal_hsa_event_t* device_signal;
-    // The device event to wait. NULL means no device event available to wait
+    // The device event to wait. nullptr means no device event available to wait
     // for this timepoint at the moment.
     iree_hal_hsa_event_t* device_wait;
   } timepoint;