fix vck5000 tests (#390)

Fixing tests to make them build correctly after rocm integration (#367).

python/air/compiler/aircc/configure.py.in

@@ -5,5 +5,5 @@
 
 air_link_with_xchesscc = @CONFIG_LINK_WITH_XCHESSCC@
 air_compile_with_xchesscc = @CONFIG_COMPILE_WITH_XCHESSCC@
-libxaie_path = "@LibXAIE_ROOT@"
+libxaie_path = "@XILINX_XAIE_DIR@"
 rocm_path = "@hsa-runtime64_DIR@"

test/02_mb_dispatch/run.lit

@@ -6,5 +6,5 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: %CLANG %S/test.cpp -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -ltest_lib -I%aie_runtime_lib%/test_lib/include -L%aie_runtime_lib%/test_lib/lib %airhost_libs% -o %T/test.elf
+// RUN: %CLANG %S/test.cpp -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/06_air_link_shared/run.lit

@@ -7,5 +7,5 @@
 //===----------------------------------------------------------------------===//
 
 // RUN: aircc.py --shared -row-offset=2 -col-offset=7 %S/air.mlir -o aie_ctrl.so
-// RUN: %CLANG %S/test.cpp -ltest_lib -L%aie_runtime_lib%/test_lib/lib -g -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%aie_runtime_lib%/test_lib/include -I./ -I./air_project/segment_0 -I%air_runtime_lib%/airhost/include -rdynamic -lxaiengine %airhost_libs% -o %T/test.elf
+// RUN: %CLANG %S/test.cpp -Iair_project/segment_0 -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/07_mb_beef_maker/run.lit

@@ -6,5 +6,5 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aiecc.py %S/aie.mlir -I%air_runtime_lib%/airhost/include -I%aie_runtime_lib%/test_lib/include -ltest_lib -L%aie_runtime_lib%/test_lib/lib -L%air_runtime_lib%/airhost %S/test.cpp -Wl,--whole-archive -lairhost -Wl,--no-whole-archive -lstdc++ -ldl -o %T/test.elf
+// RUN: aiecc.py %S/aie.mlir %S/test.cpp -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/09_mb_hsa_herd_lock/run.lit

@@ -6,5 +6,5 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aiecc.py %S/aie.mlir -I%air_runtime_lib%/airhost/include -I%aie_runtime_lib%/test_lib/include -ltest_lib -L%aie_runtime_lib%/test_lib/lib -L%air_runtime_lib%/airhost %S/test.cpp -Wl,--whole-archive -lairhost -Wl,--no-whole-archive -lstdc++ -ldl -o %T/test.elf
+// RUN: aiecc.py %S/aie.mlir %S/test.cpp -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/10_mb_shim_dma_to_tile_dma/run.lit

@@ -6,5 +6,5 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aiecc.py %S/aie.mlir -I%air_runtime_lib%/airhost/include -I%aie_runtime_lib%/test_lib/include -ltest_lib -L%aie_runtime_lib%/test_lib/lib -L%air_runtime_lib%/airhost %S/test.cpp -Wl,--whole-archive -lairhost -Wl,--no-whole-archive -lstdc++ -ldl -o %T/test.elf
+// RUN: aiecc.py %S/aie.mlir -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% %S/test.cpp -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/11_mb_shim_dma_from_tile_dma/run.lit

@@ -6,5 +6,5 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aiecc.py %S/aie.mlir -I%air_runtime_lib%/airhost/include -I%aie_runtime_lib%/test_lib/include -ltest_lib -L%aie_runtime_lib%/test_lib/lib -L%air_runtime_lib%/airhost %S/test.cpp -Wl,--whole-archive -lairhost -Wl,--no-whole-archive -lstdc++ -ldl -o %T/test.elf
+// RUN: aiecc.py %S/aie.mlir -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% %S/test.cpp -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/12_dual_channel_shim_dma/run.lit

@@ -6,5 +6,5 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aiecc.py %S/aie.mlir -I%air_runtime_lib%/airhost/include -I%aie_runtime_lib%/test_lib/include -ltest_lib -L%aie_runtime_lib%/test_lib/lib -L%air_runtime_lib%/airhost %S/test.cpp -Wl,--whole-archive -lairhost -Wl,--no-whole-archive -lstdc++ -ldl -o %T/test.elf
+// RUN: aiecc.py %S/aie.mlir -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% %S/test.cpp -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/13_mb_add_one/run.lit

@@ -6,5 +6,5 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aiecc.py %S/aie.mlir -I%air_runtime_lib%/airhost/include -I%aie_runtime_lib%/test_lib/include -ltest_lib -L%aie_runtime_lib%/test_lib/lib -L%air_runtime_lib%/airhost %S/test.cpp -Wl,--whole-archive -lairhost -Wl,--no-whole-archive -lstdc++ -ldl -o %T/test.elf
+// RUN: aiecc.py %S/aie.mlir -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% %S/test.cpp -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/14_multi_shim_dma/run.lit

@@ -6,5 +6,5 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aiecc.py %S/aie.mlir -I%air_runtime_lib%/airhost/include -I%aie_runtime_lib%/test_lib/include -ltest_lib -L%aie_runtime_lib%/test_lib/lib -L%air_runtime_lib%/airhost %S/test.cpp -Wl,--whole-archive -lairhost -Wl,--no-whole-archive -lstdc++ -ldl -o %T/test.elf
+// RUN: aiecc.py %S/aie.mlir -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% %S/test.cpp -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/15_dual_mb_dual_herd_add_one/run.lit

@@ -6,5 +6,5 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aiecc.py %S/aie.mlir -I%air_runtime_lib%/airhost/include -I%aie_runtime_lib%/test_lib/include -ltest_lib -L%aie_runtime_lib%/test_lib/lib -L%air_runtime_lib%/airhost %S/test.cpp -Wl,--whole-archive -lairhost -Wl,--no-whole-archive -lstdc++ -ldl -o %T/test.elf
+// RUN: aiecc.py %S/aie.mlir -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% %S/test.cpp -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/15_dual_mb_dual_herd_add_one/test.cpp

@@ -23,6 +23,9 @@
 
 #include "aie_inc.cpp"
 
+#include "hsa/hsa.h"
+#include "hsa/hsa_ext_amd.h"
+
 #define XAIE_NUM_COLS 36
 
 int
@@ -38,7 +41,7 @@ main(int argc, char *argv[])
     return -1;
   }
 
-  std::vector<air_agent_t> agents;
+  std::vector<hsa_agent_t> agents;
   auto get_agents_ret = air_get_agents(agents);
   assert(get_agents_ret == HSA_STATUS_SUCCESS && "failed to get agents!");
 
@@ -49,14 +52,21 @@ main(int argc, char *argv[])
 
   std::cout << "Found " << agents.size() << " agents" << std::endl;
 
-  std::vector<queue_t *> queues;
+  uint32_t aie_max_queue_size(0);
+  hsa_agent_get_info(agents[0], HSA_AGENT_INFO_QUEUE_MAX_SIZE,
+                     &aie_max_queue_size);
+
+  std::cout << "Max AIE queue size: " << aie_max_queue_size << std::endl;
+
+  std::vector<hsa_queue_t *> queues;
   for (auto agent : agents) {
     // create the queue
-    queue_t *q = nullptr;
-    auto create_queue_ret =
-        air_queue_create(MB_QUEUE_SIZE, HSA_QUEUE_TYPE_SINGLE, &q, agent.handle,
-                         0 /* device_id (optional) */);
-    assert(create_queue_ret == 0 && "failed to create queue!");
+    hsa_queue_t *q = NULL;
+    auto queue_create_status =
+        hsa_queue_create(agents[0], aie_max_queue_size, HSA_QUEUE_TYPE_SINGLE,
+                         nullptr, nullptr, 0, 0, &q);
+    if (queue_create_status != HSA_STATUS_SUCCESS)
+      std::cout << "hsa_queue_create failed" << std::endl;
     queues.push_back(q);
   }
 
@@ -75,46 +85,38 @@ main(int argc, char *argv[])
   //
   // Set up a 1x3 herd starting 7,0
   //
-  uint64_t wr_idx = queue_add_write_index(queues[0], 1);
+  uint64_t wr_idx = hsa_queue_add_write_index_relaxed(queues[0], 1);
   uint64_t packet_id = wr_idx % queues[0]->size;
-  dispatch_packet_t *segment_pkt =
-      (dispatch_packet_t *)(queues[0]->base_address_vaddr) + packet_id;
-  air_packet_segment_init(segment_pkt, 0, col, 1, row, 3);
-  air_queue_dispatch_and_wait(queues[0], wr_idx, segment_pkt);
+  hsa_agent_dispatch_packet_t segment_pkt;
+  air_packet_segment_init(&segment_pkt, 0, col, 1, row, 3);
+  air_queue_dispatch_and_wait(&agents[0], queues[0], packet_id, wr_idx,
+                              &segment_pkt);
   //
   // Set up a 1x3 herd starting 34,0
   //
-  uint64_t wr_idx2 = queue_add_write_index(queues[1], 1);
+  uint64_t wr_idx2 = hsa_queue_add_write_index_relaxed(queues[1], 1);
   uint64_t packet_id2 = wr_idx2 % queues[1]->size;
-  dispatch_packet_t *segment_pkt2 =
-      (dispatch_packet_t *)(queues[1]->base_address_vaddr) + packet_id2;
-  air_packet_segment_init(segment_pkt2, 0, col2, 1, row, 3);
-  air_queue_dispatch_and_wait(queues[1], wr_idx2, segment_pkt2);
+  hsa_agent_dispatch_packet_t segment_pkt2;
+  air_packet_segment_init(&segment_pkt2, 0, col2, 1, row, 3);
+  air_queue_dispatch_and_wait(&agents[0], queues[1], packet_id, wr_idx2,
+                              &segment_pkt2);
 
-  wr_idx = queue_add_write_index(queues[0], 1);
+  wr_idx = hsa_queue_add_write_index_relaxed(queues[0], 1);
   packet_id = wr_idx % queues[0]->size;
-  dispatch_packet_t *shim_pkt =
-      (dispatch_packet_t *)(queues[0]->base_address_vaddr) + packet_id;
-  air_packet_device_init(shim_pkt, XAIE_NUM_COLS);
-  air_queue_dispatch_and_wait(queues[0], wr_idx, shim_pkt);
+  hsa_agent_dispatch_packet_t shim_pkt;
+  air_packet_device_init(&shim_pkt, XAIE_NUM_COLS);
+  air_queue_dispatch_and_wait(&agents[0], queues[0], packet_id, wr_idx,
+                              &shim_pkt);
 
   mlir_aie_configure_cores(xaie);
   mlir_aie_configure_switchboxes(xaie);
   mlir_aie_initialize_locks(xaie);
   mlir_aie_configure_dmas(xaie);
   mlir_aie_start_cores(xaie);
 
-  // Want to initializing the device memory allocator
-  if (air_init_dev_mem_allocator(0x8000 /* dev_mem_size */,
-                                 0 /* device_id (optional)*/)) {
-    std::cout << "Error creating device memory allocator" << std::endl;
-    return -1;
-  }
-
 #define DMA_COUNT 16
 
-  uint32_t *bram_ptr =
-      (uint32_t *)air_dev_mem_alloc(4 * DMA_COUNT * sizeof(uint32_t));
+  uint32_t *bram_ptr = (uint32_t *)air_malloc(4 * DMA_COUNT * sizeof(uint32_t));
   if (bram_ptr != NULL) {
     for (int i=0; i<DMA_COUNT; i++) {
       bram_ptr[i] = i+1;
@@ -135,56 +137,54 @@ main(int argc, char *argv[])
   }
 
   // create the queues
-  queue_t *q = queues[0];
-  queue_t *q2 = queues[1];
+  hsa_queue_t *q = queues[0];
+  hsa_queue_t *q2 = queues[1];
 
   //
   // send the data
   //
 
-  wr_idx = queue_add_write_index(q, 1);
+  wr_idx = hsa_queue_add_write_index_relaxed(q, 1);
   packet_id = wr_idx % q->size;
-  dispatch_packet_t *pkt1 = (dispatch_packet_t*)(q->base_address_vaddr) + packet_id;
-  air_packet_nd_memcpy(pkt1, 0, 7, 1, 0, 4, 2, air_dev_mem_get_pa(bram_ptr),
+  hsa_agent_dispatch_packet_t pkt1;
+  air_packet_nd_memcpy(&pkt1, 0, 7, 1, 0, 4, 2, (size_t)bram_ptr,
                        DMA_COUNT * sizeof(float), 1, 0, 1, 0, 1, 0);
 
   //
   // read the data
   //
 
-  wr_idx = queue_add_write_index(q, 1);
+  wr_idx = hsa_queue_add_write_index_relaxed(q, 1);
   packet_id = wr_idx % q->size;
-  dispatch_packet_t *pkt2 = (dispatch_packet_t*)(q->base_address_vaddr) + packet_id;
-  air_packet_nd_memcpy(pkt2, 0, 7, 0, 0, 4, 2,
-                       air_dev_mem_get_pa(bram_ptr) +
-                           (DMA_COUNT * sizeof(float)),
+  hsa_agent_dispatch_packet_t pkt2;
+  air_packet_nd_memcpy(&pkt2, 0, 7, 0, 0, 4, 2,
+                       (size_t)(bram_ptr + (DMA_COUNT * sizeof(float))),
                        DMA_COUNT * sizeof(float), 1, 0, 1, 0, 1, 0);
 
   //
   // send the data
   //
 
-  wr_idx2 = queue_add_write_index(q2, 1);
+  wr_idx2 = hsa_queue_add_write_index_relaxed(q2, 1);
   packet_id2 = wr_idx2 % q2->size;
-  dispatch_packet_t *pkt12 = (dispatch_packet_t*)(q2->base_address_vaddr) + packet_id2;
-  air_packet_nd_memcpy(pkt12, 0, 34, 1, 0, 4, 2, air_dev_mem_get_pa(bram_ptr),
+  hsa_agent_dispatch_packet_t pkt12;
+  air_packet_nd_memcpy(&pkt12, 0, 34, 1, 0, 4, 2, (size_t)bram_ptr,
                        DMA_COUNT * sizeof(float), 1, 0, 1, 0, 1, 0);
 
   //
   // read the data
   //
 
-  wr_idx2 = queue_add_write_index(q2, 1);
+  wr_idx2 = hsa_queue_add_write_index_relaxed(q2, 1);
   packet_id2 = wr_idx2 % q2->size;
-  dispatch_packet_t *pkt22 = (dispatch_packet_t*)(q2->base_address_vaddr) + packet_id2;
-  air_packet_nd_memcpy(pkt22, 0, 34, 0, 0, 4, 2,
-                       air_dev_mem_get_pa(bram_ptr) +
-                           (2 * DMA_COUNT * sizeof(float)),
+  hsa_agent_dispatch_packet_t pkt22;
+  air_packet_nd_memcpy(&pkt22, 0, 34, 0, 0, 4, 2,
+                       (size_t)(bram_ptr + (2 * DMA_COUNT * sizeof(float))),
                        DMA_COUNT * sizeof(float), 1, 0, 1, 0, 1, 0);
 
-  air_queue_dispatch(q, wr_idx, pkt2);
-  air_queue_dispatch_and_wait(q2, wr_idx2, pkt22);
-  air_queue_wait(q, pkt2);
+  air_queue_dispatch(q, packet_id, wr_idx, &pkt2);
+  air_queue_dispatch_and_wait(&agents[0], q2, packet_id, wr_idx2, &pkt22);
+  air_queue_wait(q, &pkt2);
 
   int errors = 0;
 
@@ -233,8 +233,6 @@ main(int argc, char *argv[])
     }
   }
 
-  air_dev_mem_allocator_free();
-
   if (!errors) {
     printf("PASS!\n");
     return 0;

test/16_multi_shim_dma_all_channels/run.lit

@@ -6,5 +6,5 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aiecc.py %S/aie.mlir -I%air_runtime_lib%/airhost/include -I%aie_runtime_lib%/test_lib/include -ltest_lib -L%aie_runtime_lib%/test_lib/lib -L%air_runtime_lib%/airhost %S/test.cpp -Wl,--whole-archive -lairhost -Wl,--no-whole-archive -lstdc++ -ldl -o %T/test.elf
+// RUN: aiecc.py %S/aie.mlir -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% %S/test.cpp -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/18_2d_tile_dma_transpose/run.lit

@@ -8,5 +8,5 @@
 
 // Expected to fail unless mlir-aie pull #55 is merged
 // XFAIL: *
-// RUN: aiecc.py %S/aie.mlir -I%air_runtime_lib%/airhost/include -I%aie_runtime_lib%/test_lib/include -ltest_lib -L%aie_runtime_lib%/test_lib/lib -L%air_runtime_lib%/airhost %S/test.cpp -Wl,--whole-archive -lairhost -Wl,--no-whole-archive -lstdc++ -ldl -o %T/test.elf
+// RUN: aiecc.py %S/aie.mlir -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% %S/test.cpp -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/19_air_nd_memcpy_to_tile_dma/run.lit

@@ -6,6 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aircc.py -row-offset=4 -col-offset=5 %S/air.mlir -o air.mlir.a
-// RUN: %CLANG %S/test.cpp -ltest_lib -L%aie_runtime_lib%/test_lib/lib -Wl,--whole-archive air.mlir.a -Wl,--no-whole-archive -g -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%aie_runtime_lib%/test_lib/include -I%air_runtime_lib%/airhost/include -rdynamic -lxaiengine %airhost_libs% -o %T/test.elf
+// RUN: aircc.py -row-offset=4 -col-offset=5 %S/air.mlir -o %T/air.a
+// RUN: %CLANG %S/test.cpp -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% -Wl,--whole-archive %T/air.a -Wl,--no-whole-archive -rdynamic -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/20_air_nd_memcpy_from_tile_dma/run.lit

@@ -6,6 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aircc.py -row-offset=4 -col-offset=5 %S/air.mlir -o air.mlir.a
-// RUN: %CLANG %S/test.cpp -ltest_lib -L%aie_runtime_lib%/test_lib/lib -Wl,--whole-archive air.mlir.a -Wl,--no-whole-archive -g -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%aie_runtime_lib%/test_lib/include -I%air_runtime_lib%/airhost/include -rdynamic -lxaiengine %airhost_libs% -o %T/test.elf
+// RUN: aircc.py -row-offset=4 -col-offset=5 %S/air.mlir -o %T/air.a
+// RUN: %CLANG %S/test.cpp -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% -Wl,--whole-archive %T/air.a -Wl,--no-whole-archive -rdynamic -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/21_air_nd_memcpy_2d/run.lit

@@ -6,6 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aircc.py -row-offset=4 -col-offset=5 %S/air.mlir -o air.mlir.a
-// RUN: %CLANG %S/test.cpp -ltest_lib -L%aie_runtime_lib%/test_lib/lib -Wl,--whole-archive air.mlir.a -Wl,--no-whole-archive -g -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%aie_runtime_lib%/test_lib/include -I%air_runtime_lib%/airhost/include -rdynamic -lxaiengine %airhost_libs% -o %T/test.elf
+// RUN: aircc.py -row-offset=4 -col-offset=5 %S/air.mlir -o %T/air.a
+// RUN: %CLANG %S/test.cpp -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% -Wl,--whole-archive %T/air.a -Wl,--no-whole-archive -rdynamic -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/23_air_shim_dma_to_tile_dma/run.lit

@@ -6,6 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aircc.py -row-offset=2 -col-offset=7 %S/air.mlir -o air.mlir.a
-// RUN: %CLANG %S/test.cpp -ltest_lib -L%aie_runtime_lib%/test_lib/lib -Wl,--whole-archive air.mlir.a -Wl,--no-whole-archive -g -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%aie_runtime_lib%/test_lib/include -I%air_runtime_lib%/airhost/include -rdynamic -lxaiengine %airhost_libs% -o %T/test.elf
+// RUN: aircc.py -row-offset=2 -col-offset=7 %S/air.mlir -o %T/air.a
+// RUN: %CLANG %S/test.cpp -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% -Wl,--whole-archive %T/air.a -Wl,--no-whole-archive -rdynamic -o %T/test.elf
 // RUN: %run_on_board %T/test.elf

test/24_air_shim_dma_from_tile_dma/run.lit

@@ -6,6 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-// RUN: aircc.py -row-offset=2 -col-offset=7 %S/air.mlir -o air.mlir.a
-// RUN: %CLANG %S/test.cpp -ltest_lib -L%aie_runtime_lib%/test_lib/lib -Wl,--whole-archive air.mlir.a -Wl,--no-whole-archive -g -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%aie_runtime_lib%/test_lib/include -I%air_runtime_lib%/airhost/include -rdynamic -lxaiengine %airhost_libs% -o %T/test.elf
+// RUN: aircc.py -row-offset=2 -col-offset=7 %S/air.mlir -o %T/air.a
+// RUN: %CLANG %S/test.cpp -I%HSA_DIR%/include -L%HSA_DIR%/lib -lhsa-runtime64 -I%LIBXAIE_DIR%/include -L%LIBXAIE_DIR%/lib -lxaiengine -I%AIE_RUNTIME_DIR%/test_lib/include -L%AIE_RUNTIME_DIR%/test_lib/lib -ltest_lib %airhost_libs% -Wl,--whole-archive %T/air.a -Wl,--no-whole-archive -rdynamic -o %T/test.elf
 // RUN: %run_on_board %T/test.elf