From a0cecc49aa638022bcd6d7e8c1a97399c4b3436d Mon Sep 17 00:00:00 2001 From: Eddie Richter Date: Wed, 20 Mar 2024 11:21:09 -0600 Subject: [PATCH] clang-format --- include/aie/Targets/AIETargets.h | 2 +- lib/Targets/AIETargetHSA.cpp | 61 +++++++++++-------- .../test_vck5000.cpp | 11 ++-- .../IRON-examples/matrix_add_one/test.cpp | 11 ++-- .../matrix_add_one/test_vck5000.cpp | 22 ++++--- .../matrix_multiplication_scalar/test.cpp | 22 ++++--- .../passthrough_hardware/test_vck5000.cpp | 11 ++-- .../IRON-examples/vector_add/test.cpp | 9 +-- .../IRON-examples/vector_add/test_vck5000.cpp | 11 ++-- .../IRON-examples/vector_max/test.cpp | 9 +-- .../IRON-examples/vector_max/test_vck5000.cpp | 30 +++++---- .../IRON-examples/vector_min/test.cpp | 9 +-- .../IRON-examples/vector_min/test_vck5000.cpp | 21 ++++--- .../IRON-examples/vector_mult/test.cpp | 9 +-- .../vector_mult/test_vck5000.cpp | 11 ++-- .../vector_scalar/test_vck5000.cpp | 11 ++-- .../vector_scalar_kernel/scale.cc | 1 - .../vector_scalar_kernel/test.cpp | 11 ++-- .../IRON-examples/vector_sum/test.cpp | 9 +-- .../IRON-examples/vector_sum/test_vck5000.cpp | 20 +++--- .../dynamic_dma_config_add_one/test.cpp | 13 ++-- .../test.cpp | 13 ++-- .../test.cpp | 11 ++-- .../ipu-xrt/matrix_add_one/test.cpp | 11 ++-- reference_designs/ipu-xrt/vector_max/test.cpp | 9 +-- reference_designs/ipu-xrt/vector_min/test.cpp | 9 +-- reference_designs/ipu-xrt/vector_sum/test.cpp | 9 +-- runtime_lib/test_lib/test_library.cpp | 28 ++++----- runtime_lib/test_lib/test_library.h | 40 ++++++------ runtime_lib/xaiengine/aie-rt | 2 +- 30 files changed, 240 insertions(+), 206 deletions(-) diff --git a/include/aie/Targets/AIETargets.h b/include/aie/Targets/AIETargets.h index 21a3a54fbb..b9b960d798 100644 --- a/include/aie/Targets/AIETargets.h +++ b/include/aie/Targets/AIETargets.h @@ -20,7 +20,7 @@ namespace AIE { mlir::LogicalResult AIETranslateToXAIEV2(mlir::ModuleOp module, llvm::raw_ostream &output); mlir::LogicalResult AIETranslateToHSA(mlir::ModuleOp module, - llvm::raw_ostream &output); + llvm::raw_ostream &output); mlir::LogicalResult AIEFlowsToJSON(mlir::ModuleOp module, llvm::raw_ostream &output); mlir::LogicalResult ADFGenerateCPPGraph(mlir::ModuleOp module, diff --git a/lib/Targets/AIETargetHSA.cpp b/lib/Targets/AIETargetHSA.cpp index 83dafdc66b..a79c01cc63 100644 --- a/lib/Targets/AIETargetHSA.cpp +++ b/lib/Targets/AIETargetHSA.cpp @@ -56,7 +56,6 @@ getAllocOpForSymbol(AIE::DeviceOp dev, StringRef sym_name) { return std::nullopt; } - mlir::LogicalResult AIETranslateToHSA(ModuleOp module, raw_ostream &output) { DenseMap tiles; @@ -66,30 +65,30 @@ mlir::LogicalResult AIETranslateToHSA(ModuleOp module, raw_ostream &output) { return module.emitOpError("expected AIE.device operation at toplevel"); DeviceOp targetOp = *(module.getOps().begin()); - // Putting the standard header + // Putting the standard header output << hsa_cpp_file_header; // Getting the func op which has the data movement - if(targetOp.getOps().empty()) { + if (targetOp.getOps().empty()) { return success(); } // Getting the sequence function op which contains the instructions mlir::func::FuncOp funcOp; for (auto op : targetOp.getOps()) { - if(op.getName().str().compare("sequence") == 0) { + if (op.getName().str().compare("sequence") == 0) { funcOp = op; } } - collectTiles(targetOp, tiles); collectBuffers(targetOp, buffers); // Generate dynamic data movement output << "void invoke_data_movement(hsa_queue_t *q, hsa_agent_t *a"; - // Looping over every Memcpy operation so we take the correct number of buffers + // Looping over every Memcpy operation so we take the correct number of + // buffers int num_ops = 0; for (auto op : funcOp.getOps()) { // Getting the IDs of the buffers @@ -112,14 +111,14 @@ mlir::LogicalResult AIETranslateToHSA(ModuleOp module, raw_ostream &output) { output << "\tuint64_t wr_idx = 0;\n"; output << "\tuint64_t packet_id = 0;\n"; - int op_count = 0; + int op_count = 0; for (auto op : funcOp.getOps()) { auto dev = funcOp->getParentOfType(); if (!dev) { op.emitOpError("couldn't get DeviceOp"); return failure(); } - + auto infoOp = getAllocOpForSymbol(dev, op.getMetadata()); if (!infoOp) { op.emitOpError("couldn't find shim_dma_allocation op"); @@ -132,14 +131,14 @@ mlir::LogicalResult AIETranslateToHSA(ModuleOp module, raw_ostream &output) { int col = infoOp->getCol(); llvm::SmallVector strides = llvm::map_to_vector( - llvm::reverse(op.getMixedStrides()), - [](OpFoldResult s) { return getConstantIntValue(s).value(); }); + llvm::reverse(op.getMixedStrides()), + [](OpFoldResult s) { return getConstantIntValue(s).value(); }); ::SmallVector sizes = llvm::map_to_vector( - llvm::reverse(op.getMixedSizes()), - [](OpFoldResult s) { return getConstantIntValue(s).value(); }); + llvm::reverse(op.getMixedSizes()), + [](OpFoldResult s) { return getConstantIntValue(s).value(); }); ::SmallVector offsets = llvm::map_to_vector( - llvm::reverse(op.getMixedOffsets()), - [](OpFoldResult s) { return getConstantIntValue(s).value(); }); + llvm::reverse(op.getMixedOffsets()), + [](OpFoldResult s) { return getConstantIntValue(s).value(); }); // buffer_offset size_t stride = 1; @@ -171,24 +170,39 @@ mlir::LogicalResult AIETranslateToHSA(ModuleOp module, raw_ostream &output) { output << "\thsa_agent_dispatch_packet_t pkt" << op_count << " ;\n"; output << "\twr_idx = hsa_queue_add_write_index_relaxed(q, 1);\n"; output << "\tpacket_id = wr_idx % q->size;\n"; - output << "\tmlir_aie_packet_nd_memcpy(&pkt" << op_count << ", 0 /* herd_id */, " << col << " /* col */, " << isMM2S << " /* dir */, " << ChannelId << "/* channel */, 4 /* Burst length */, 2 /* Memory space */, (uint64_t)buf" << arg_idx << " + " << offset << " /* Address */, " << sizes[0] * 4 << " /* 1d_length */, " << (strides[0] ? sizes[1] : 1) << " /* 2d_length */, " << (strides[0] ? strides[0] * 4: 0) << " /* 2d_stride */, " << (strides[1] ? sizes[2] : 1) << " /* 3d_length */, " << (strides[1] ? strides[1] * 4: 0) << " /* 3d_stride */ , 1 /* 4d_length */, 0 /* 4d_stride */);\n"; + output << "\tmlir_aie_packet_nd_memcpy(&pkt" << op_count + << ", 0 /* herd_id */, " << col << " /* col */, " << isMM2S + << " /* dir */, " << ChannelId + << "/* channel */, 4 /* Burst length */, 2 /* Memory space */, " + "(uint64_t)buf" + << arg_idx << " + " << offset << " /* Address */, " << sizes[0] * 4 + << " /* 1d_length */, " << (strides[0] ? sizes[1] : 1) + << " /* 2d_length */, " << (strides[0] ? strides[0] * 4 : 0) + << " /* 2d_stride */, " << (strides[1] ? sizes[2] : 1) + << " /* 3d_length */, " << (strides[1] ? strides[1] * 4 : 0) + << " /* 3d_stride */ , 1 /* 4d_length */, 0 /* 4d_stride */);\n"; bool last_op = op_count == (num_ops - 1); // Only ring the doorbell on the last packet - if(last_op) { - output << "\tmlir_aie_queue_dispatch_and_wait(a, q, packet_id, wr_idx, &pkt" << op_count << ", false);\n\n"; + if (last_op) { + output + << "\tmlir_aie_queue_dispatch_and_wait(a, q, packet_id, wr_idx, &pkt" + << op_count << ", false);\n\n"; + } else { + output << "\thsa_amd_signal_create_on_agent(1, 0, nullptr, a, 0, &pkt" + << op_count << ".completion_signal);\n"; + output << "\tmlir_aie_write_pkt(q, " + "packet_id, &pkt" + << op_count << ");\n\n"; } - else { - output << "\thsa_amd_signal_create_on_agent(1, 0, nullptr, a, 0, &pkt" << op_count << ".completion_signal);\n"; - output << "\tmlir_aie_write_pkt(q, packet_id, &pkt" << op_count << ");\n\n"; - } - + op_count++; } // Waiting to make sure each DMA is complete for (int i = 0; i < op_count; i++) { - output << "\twhile (hsa_signal_wait_scacquire(pkt" << i << ".completion_signal,\n"; + output << "\twhile (hsa_signal_wait_scacquire(pkt" << i + << ".completion_signal,\n"; output << "\tHSA_SIGNAL_CONDITION_EQ, 0, 0x80000,\n"; output << "\tHSA_WAIT_STATE_ACTIVE) != 0);\n"; } @@ -200,7 +214,6 @@ mlir::LogicalResult AIETranslateToHSA(ModuleOp module, raw_ostream &output) { output << "}\n"; - return success(); } } // namespace xilinx::AIE diff --git a/reference_designs/IRON-examples/add_one_objFifo_no_mem_tile/test_vck5000.cpp b/reference_designs/IRON-examples/add_one_objFifo_no_mem_tile/test_vck5000.cpp index 438cd92a14..fbd2776740 100644 --- a/reference_designs/IRON-examples/add_one_objFifo_no_mem_tile/test_vck5000.cpp +++ b/reference_designs/IRON-examples/add_one_objFifo_no_mem_tile/test_vck5000.cpp @@ -19,11 +19,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -54,8 +54,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -121,7 +122,7 @@ int main(int argc, char *argv[]) { // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/IRON-examples/matrix_add_one/test.cpp b/reference_designs/IRON-examples/matrix_add_one/test.cpp index f8c6626922..c8804fecfc 100644 --- a/reference_designs/IRON-examples/matrix_add_one/test.cpp +++ b/reference_designs/IRON-examples/matrix_add_one/test.cpp @@ -179,15 +179,16 @@ int main(int argc, const char *argv[]) { uint32_t d = bufOut[i]; if (row < TILE_HEIGHT && col < TILE_WIDTH) { - if(d != s + 1) { + if (d != s + 1) { errors++; printf("[ERROR] row %d and col %d, %d != %d\n", row, col, s, d); } - } - else { - if(d == s + 1) { + } else { + if (d == s + 1) { errors++; - printf("[ERROR] row %d and col %d, %d == %d -- this was not supposed to be changed\n", row, col, s, d); + printf("[ERROR] row %d and col %d, %d == %d -- this was not supposed " + "to be changed\n", + row, col, s, d); } } diff --git a/reference_designs/IRON-examples/matrix_add_one/test_vck5000.cpp b/reference_designs/IRON-examples/matrix_add_one/test_vck5000.cpp index 81c78e6ce3..36f20ca322 100644 --- a/reference_designs/IRON-examples/matrix_add_one/test_vck5000.cpp +++ b/reference_designs/IRON-examples/matrix_add_one/test_vck5000.cpp @@ -19,11 +19,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -65,8 +65,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -124,15 +125,16 @@ int main(int argc, char *argv[]) { uint32_t d = out[i]; if (row < TILE_HEIGHT && col < TILE_WIDTH) { - if(d != s + 1) { + if (d != s + 1) { errors++; printf("[ERROR] row %d and col %d, %d != %d\n", row, col, s, d); } - } - else { - if(d == s + 1) { + } else { + if (d == s + 1) { errors++; - printf("[ERROR] row %d and col %d, %d == %d -- this was not supposed to be changed\n", row, col, s, d); + printf("[ERROR] row %d and col %d, %d == %d -- this was not supposed " + "to be changed\n", + row, col, s, d); } } @@ -142,7 +144,7 @@ int main(int argc, char *argv[]) { // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/IRON-examples/matrix_multiplication_scalar/test.cpp b/reference_designs/IRON-examples/matrix_multiplication_scalar/test.cpp index 81c78e6ce3..36f20ca322 100644 --- a/reference_designs/IRON-examples/matrix_multiplication_scalar/test.cpp +++ b/reference_designs/IRON-examples/matrix_multiplication_scalar/test.cpp @@ -19,11 +19,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -65,8 +65,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -124,15 +125,16 @@ int main(int argc, char *argv[]) { uint32_t d = out[i]; if (row < TILE_HEIGHT && col < TILE_WIDTH) { - if(d != s + 1) { + if (d != s + 1) { errors++; printf("[ERROR] row %d and col %d, %d != %d\n", row, col, s, d); } - } - else { - if(d == s + 1) { + } else { + if (d == s + 1) { errors++; - printf("[ERROR] row %d and col %d, %d == %d -- this was not supposed to be changed\n", row, col, s, d); + printf("[ERROR] row %d and col %d, %d == %d -- this was not supposed " + "to be changed\n", + row, col, s, d); } } @@ -142,7 +144,7 @@ int main(int argc, char *argv[]) { // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/IRON-examples/passthrough_hardware/test_vck5000.cpp b/reference_designs/IRON-examples/passthrough_hardware/test_vck5000.cpp index 967ef78115..53f310d6aa 100644 --- a/reference_designs/IRON-examples/passthrough_hardware/test_vck5000.cpp +++ b/reference_designs/IRON-examples/passthrough_hardware/test_vck5000.cpp @@ -19,11 +19,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -54,8 +54,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -119,7 +120,7 @@ int main(int argc, char *argv[]) { // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/IRON-examples/vector_add/test.cpp b/reference_designs/IRON-examples/vector_add/test.cpp index b460156b39..c784db343d 100644 --- a/reference_designs/IRON-examples/vector_add/test.cpp +++ b/reference_designs/IRON-examples/vector_add/test.cpp @@ -175,13 +175,14 @@ int main(int argc, const char *argv[]) { int errors = 0; for (uint32_t i = 0; i < 64; i++) { - if (*(bufOut + i) != *(bufInA + i) + *(bufInB + i) ) { - std::cout << "Error in output " << *(bufOut + i) << " != " << *(bufInA + i) << " + " << *(bufInB + i) + if (*(bufOut + i) != *(bufInA + i) + *(bufInB + i)) { + std::cout << "Error in output " << *(bufOut + i) + << " != " << *(bufInA + i) << " + " << *(bufInB + i) << std::endl; errors++; } else { - std::cout << "Correct output " << *(bufOut + i) << " == " << *(bufInA + i) + *(bufInB + i) - << std::endl; + std::cout << "Correct output " << *(bufOut + i) + << " == " << *(bufInA + i) + *(bufInB + i) << std::endl; } } diff --git a/reference_designs/IRON-examples/vector_add/test_vck5000.cpp b/reference_designs/IRON-examples/vector_add/test_vck5000.cpp index 390de4cc16..c51a0f4041 100644 --- a/reference_designs/IRON-examples/vector_add/test_vck5000.cpp +++ b/reference_designs/IRON-examples/vector_add/test_vck5000.cpp @@ -19,11 +19,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -54,8 +54,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -123,7 +124,7 @@ int main(int argc, char *argv[]) { // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/IRON-examples/vector_max/test.cpp b/reference_designs/IRON-examples/vector_max/test.cpp index 34d28e220f..ffe277e862 100644 --- a/reference_designs/IRON-examples/vector_max/test.cpp +++ b/reference_designs/IRON-examples/vector_max/test.cpp @@ -151,7 +151,7 @@ int main(int argc, const char *argv[]) { srcVecA.push_back(i + 1); memcpy(bufInA, srcVecA.data(), (srcVecA.size() * sizeof(uint32_t))); - bufInA[IN_SIZE / 2] = 654321; + bufInA[IN_SIZE / 2] = 654321; bufInA[IN_SIZE - 1] = 100; void *bufInstr = bo_instr.map(); @@ -170,7 +170,7 @@ int main(int argc, const char *argv[]) { uint32_t *bufOut = bo_out.map(); int errors = 0; - + uint32_t max_val = 0; for (uint32_t i = 0; i < IN_SIZE; i++) { if (*(bufInA + i) > max_val) { @@ -178,8 +178,9 @@ int main(int argc, const char *argv[]) { } } - if(*bufOut != max_val) { - std::cout << "[ERROR] Maximum value is " << max_val << " but kernel returned " << *bufOut << "\n"; + if (*bufOut != max_val) { + std::cout << "[ERROR] Maximum value is " << max_val + << " but kernel returned " << *bufOut << "\n"; errors++; } diff --git a/reference_designs/IRON-examples/vector_max/test_vck5000.cpp b/reference_designs/IRON-examples/vector_max/test_vck5000.cpp index fb323f0759..4c039f3ca2 100644 --- a/reference_designs/IRON-examples/vector_max/test_vck5000.cpp +++ b/reference_designs/IRON-examples/vector_max/test_vck5000.cpp @@ -19,11 +19,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -54,8 +54,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -87,7 +88,8 @@ int main(int argc, char *argv[]) { ext_mem_model_t buf0, buf1, buf2; uint32_t *in_a = (uint32_t *)mlir_aie_mem_alloc(xaie, buf0, DMA_COUNT); uint32_t *in_b = (uint32_t *)mlir_aie_mem_alloc(xaie, buf1, DMA_COUNT); - uint32_t *out = (uint32_t *)mlir_aie_mem_alloc(xaie, buf2, 4 /* For some reason can't do 1 */); + uint32_t *out = (uint32_t *)mlir_aie_mem_alloc( + xaie, buf2, 4 /* For some reason can't do 1 */); mlir_aie_sync_mem_dev(buf0); mlir_aie_sync_mem_dev(buf1); mlir_aie_sync_mem_dev(buf2); @@ -97,16 +99,16 @@ int main(int argc, char *argv[]) { return -1; } - out[0] = 0xdeface; for (int i = 0; i < DMA_COUNT; i++) { in_a[i] = i + 1; } - in_a[DMA_COUNT / 2] = 123456; + in_a[DMA_COUNT / 2] = 123456; in_a[DMA_COUNT - 1] = 100; - //printf("[EDDIE DEBUG] max_val before data movement is %d\n", mlir_aie_read_buffer_max_val(xaie, 0)); + // printf("[EDDIE DEBUG] max_val before data movement is %d\n", + // mlir_aie_read_buffer_max_val(xaie, 0)); // Pass arguments in the order of dma_memcpys in the mlir invoke_data_movement(queues[0], &agents[0], out, in_a); @@ -116,21 +118,23 @@ int main(int argc, char *argv[]) { uint32_t max_val = 0; for (int i = 0; i < DMA_COUNT; i++) { uint32_t s = in_a[i]; - if(max_val < s) { + if (max_val < s) { max_val = s; } } - //printf("[EDDIE DEBUG] max_val before data movement is %d\n", mlir_aie_read_buffer_max_val(xaie, 0)); + // printf("[EDDIE DEBUG] max_val before data movement is %d\n", + // mlir_aie_read_buffer_max_val(xaie, 0)); - if(*out != max_val) { + if (*out != max_val) { errors++; - printf("[ERROR] Maximum value is %d but kernel returned %d\n", max_val, *out); + printf("[ERROR] Maximum value is %d but kernel returned %d\n", max_val, + *out); } // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/IRON-examples/vector_min/test.cpp b/reference_designs/IRON-examples/vector_min/test.cpp index a70649c16e..61339fc01e 100644 --- a/reference_designs/IRON-examples/vector_min/test.cpp +++ b/reference_designs/IRON-examples/vector_min/test.cpp @@ -151,7 +151,7 @@ int main(int argc, const char *argv[]) { srcVecA.push_back(i + 1); memcpy(bufInA, srcVecA.data(), (srcVecA.size() * sizeof(uint32_t))); - bufInA[IN_SIZE / 2] = 4; + bufInA[IN_SIZE / 2] = 4; bufInA[IN_SIZE - 1] = 100; void *bufInstr = bo_instr.map(); @@ -170,7 +170,7 @@ int main(int argc, const char *argv[]) { uint32_t *bufOut = bo_out.map(); int errors = 0; - + uint32_t min_val = 1000000; for (uint32_t i = 0; i < IN_SIZE; i++) { if (*(bufInA + i) < min_val) { @@ -178,8 +178,9 @@ int main(int argc, const char *argv[]) { } } - if(*bufOut != min_val) { - std::cout << "[ERROR] Min value is " << min_val << " but kernel returned " << *bufOut << "\n"; + if (*bufOut != min_val) { + std::cout << "[ERROR] Min value is " << min_val << " but kernel returned " + << *bufOut << "\n"; errors++; } diff --git a/reference_designs/IRON-examples/vector_min/test_vck5000.cpp b/reference_designs/IRON-examples/vector_min/test_vck5000.cpp index f3881e7694..28db01e2ff 100644 --- a/reference_designs/IRON-examples/vector_min/test_vck5000.cpp +++ b/reference_designs/IRON-examples/vector_min/test_vck5000.cpp @@ -19,11 +19,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -54,8 +54,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -87,7 +88,8 @@ int main(int argc, char *argv[]) { ext_mem_model_t buf0, buf1, buf2; uint32_t *in_a = (uint32_t *)mlir_aie_mem_alloc(xaie, buf0, DMA_COUNT); uint32_t *in_b = (uint32_t *)mlir_aie_mem_alloc(xaie, buf1, DMA_COUNT); - uint32_t *out = (uint32_t *)mlir_aie_mem_alloc(xaie, buf2, 4 /* For some reason can't do 1 */); + uint32_t *out = (uint32_t *)mlir_aie_mem_alloc( + xaie, buf2, 4 /* For some reason can't do 1 */); mlir_aie_sync_mem_dev(buf0); mlir_aie_sync_mem_dev(buf1); mlir_aie_sync_mem_dev(buf2); @@ -97,13 +99,12 @@ int main(int argc, char *argv[]) { return -1; } - out[0] = 0xdeface; for (int i = 0; i < DMA_COUNT; i++) { in_a[i] = i + 100; } - in_a[DMA_COUNT / 2] = 3; + in_a[DMA_COUNT / 2] = 3; in_a[DMA_COUNT - 1] = 100; // Pass arguments in the order of dma_memcpys in the mlir @@ -114,19 +115,19 @@ int main(int argc, char *argv[]) { uint32_t min_val = 1000; for (int i = 0; i < DMA_COUNT; i++) { uint32_t s = in_a[i]; - if(min_val > s) { + if (min_val > s) { min_val = s; } } - if(*out != min_val) { + if (*out != min_val) { errors++; printf("[ERROR] Min value is %d but kernel returned %d\n", min_val, *out); } // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/IRON-examples/vector_mult/test.cpp b/reference_designs/IRON-examples/vector_mult/test.cpp index 39cf7a7e79..3e9990f0ee 100644 --- a/reference_designs/IRON-examples/vector_mult/test.cpp +++ b/reference_designs/IRON-examples/vector_mult/test.cpp @@ -176,13 +176,14 @@ int main(int argc, const char *argv[]) { for (uint32_t i = 0; i < 64; i++) { uint32_t ref = i + 2; - if (*(bufOut + i) != *(bufInA + i) * *(bufInB + i) ) { - std::cout << "Error in output " << *(bufOut + i) << " != " << *(bufInA + i) << " + " << *(bufInB + i) + if (*(bufOut + i) != *(bufInA + i) * *(bufInB + i)) { + std::cout << "Error in output " << *(bufOut + i) + << " != " << *(bufInA + i) << " + " << *(bufInB + i) << std::endl; errors++; } else { - std::cout << "Correct output " << *(bufOut + i) << " == " << *(bufInA + i) * *(bufInB + i) - << std::endl; + std::cout << "Correct output " << *(bufOut + i) + << " == " << *(bufInA + i) * *(bufInB + i) << std::endl; } } diff --git a/reference_designs/IRON-examples/vector_mult/test_vck5000.cpp b/reference_designs/IRON-examples/vector_mult/test_vck5000.cpp index d637dcdd0d..3b1e815694 100644 --- a/reference_designs/IRON-examples/vector_mult/test_vck5000.cpp +++ b/reference_designs/IRON-examples/vector_mult/test_vck5000.cpp @@ -19,11 +19,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -54,8 +54,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -123,7 +124,7 @@ int main(int argc, char *argv[]) { // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/IRON-examples/vector_scalar/test_vck5000.cpp b/reference_designs/IRON-examples/vector_scalar/test_vck5000.cpp index 8caf0a81c2..4e60c2875c 100644 --- a/reference_designs/IRON-examples/vector_scalar/test_vck5000.cpp +++ b/reference_designs/IRON-examples/vector_scalar/test_vck5000.cpp @@ -19,11 +19,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -54,8 +54,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -121,7 +122,7 @@ int main(int argc, char *argv[]) { // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/IRON-examples/vector_scalar_kernel/scale.cc b/reference_designs/IRON-examples/vector_scalar_kernel/scale.cc index 2d21db1801..c04f08031f 100755 --- a/reference_designs/IRON-examples/vector_scalar_kernel/scale.cc +++ b/reference_designs/IRON-examples/vector_scalar_kernel/scale.cc @@ -8,7 +8,6 @@ // //===----------------------------------------------------------------------===// - template void scale(T_in *a, T_out *c, T_in factor) { for (int i = 0; i < N; i++) { diff --git a/reference_designs/IRON-examples/vector_scalar_kernel/test.cpp b/reference_designs/IRON-examples/vector_scalar_kernel/test.cpp index 8caf0a81c2..4e60c2875c 100644 --- a/reference_designs/IRON-examples/vector_scalar_kernel/test.cpp +++ b/reference_designs/IRON-examples/vector_scalar_kernel/test.cpp @@ -19,11 +19,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -54,8 +54,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -121,7 +122,7 @@ int main(int argc, char *argv[]) { // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/IRON-examples/vector_sum/test.cpp b/reference_designs/IRON-examples/vector_sum/test.cpp index 64a8bf79e8..b87b46b888 100644 --- a/reference_designs/IRON-examples/vector_sum/test.cpp +++ b/reference_designs/IRON-examples/vector_sum/test.cpp @@ -151,7 +151,7 @@ int main(int argc, const char *argv[]) { srcVecA.push_back(i + 1); memcpy(bufInA, srcVecA.data(), (srcVecA.size() * sizeof(uint32_t))); - bufInA[IN_SIZE / 2] = 654321; + bufInA[IN_SIZE / 2] = 654321; bufInA[IN_SIZE - 1] = 100; void *bufInstr = bo_instr.map(); @@ -170,14 +170,15 @@ int main(int argc, const char *argv[]) { uint32_t *bufOut = bo_out.map(); int errors = 0; - + uint32_t sum_val = 0; for (uint32_t i = 0; i < IN_SIZE; i++) { sum_val += *(bufInA + i); } - if(*bufOut != sum_val) { - std::cout << "[ERROR] Maximum value is " << sum_val << " but kernel returned " << *bufOut << "\n"; + if (*bufOut != sum_val) { + std::cout << "[ERROR] Maximum value is " << sum_val + << " but kernel returned " << *bufOut << "\n"; errors++; } diff --git a/reference_designs/IRON-examples/vector_sum/test_vck5000.cpp b/reference_designs/IRON-examples/vector_sum/test_vck5000.cpp index fd405400ba..34abda033f 100644 --- a/reference_designs/IRON-examples/vector_sum/test_vck5000.cpp +++ b/reference_designs/IRON-examples/vector_sum/test_vck5000.cpp @@ -19,11 +19,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -54,8 +54,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -87,7 +88,8 @@ int main(int argc, char *argv[]) { ext_mem_model_t buf0, buf1, buf2; uint32_t *in_a = (uint32_t *)mlir_aie_mem_alloc(xaie, buf0, DMA_COUNT); uint32_t *in_b = (uint32_t *)mlir_aie_mem_alloc(xaie, buf1, DMA_COUNT); - uint32_t *out = (uint32_t *)mlir_aie_mem_alloc(xaie, buf2, 4 /* For some reason can't do 1 */); + uint32_t *out = (uint32_t *)mlir_aie_mem_alloc( + xaie, buf2, 4 /* For some reason can't do 1 */); mlir_aie_sync_mem_dev(buf0); mlir_aie_sync_mem_dev(buf1); mlir_aie_sync_mem_dev(buf2); @@ -97,7 +99,6 @@ int main(int argc, char *argv[]) { return -1; } - out[0] = 0xdeface; for (int i = 0; i < DMA_COUNT; i++) { in_a[i] = i + 1; @@ -113,14 +114,15 @@ int main(int argc, char *argv[]) { sum_val += in_a[i]; } - if(*out != sum_val) { + if (*out != sum_val) { errors++; - printf("[ERROR] Maximum value is %d but kernel returned %d\n", sum_val, *out); + printf("[ERROR] Maximum value is %d but kernel returned %d\n", sum_val, + *out); } // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/dynamic_dma_config_add_one/test.cpp b/reference_designs/dynamic_dma_config_add_one/test.cpp index 32494d6f99..c26cb76355 100644 --- a/reference_designs/dynamic_dma_config_add_one/test.cpp +++ b/reference_designs/dynamic_dma_config_add_one/test.cpp @@ -18,8 +18,8 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" #include "aie_inc.cpp" @@ -52,8 +52,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -151,8 +152,8 @@ int main(int argc, char *argv[]) { for (int i = 0; i < DMA_COUNT; i++) { uint32_t s = src[i]; uint32_t d = dst[i]; - //printf("src[%d] = 0x%lx\n", i, src[i]); - //printf("dst[%d] = 0x%lx\n", i, dst[i]); + // printf("src[%d] = 0x%lx\n", i, src[i]); + // printf("dst[%d] = 0x%lx\n", i, dst[i]); if (d != (s + 1)) { errors++; printf("mismatch %x != 1 + %x\n", d, s); @@ -161,7 +162,7 @@ int main(int argc, char *argv[]) { // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/dynamic_dma_config_add_one_kernel/test.cpp b/reference_designs/dynamic_dma_config_add_one_kernel/test.cpp index 32494d6f99..c26cb76355 100644 --- a/reference_designs/dynamic_dma_config_add_one_kernel/test.cpp +++ b/reference_designs/dynamic_dma_config_add_one_kernel/test.cpp @@ -18,8 +18,8 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" #include "aie_inc.cpp" @@ -52,8 +52,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -151,8 +152,8 @@ int main(int argc, char *argv[]) { for (int i = 0; i < DMA_COUNT; i++) { uint32_t s = src[i]; uint32_t d = dst[i]; - //printf("src[%d] = 0x%lx\n", i, src[i]); - //printf("dst[%d] = 0x%lx\n", i, dst[i]); + // printf("src[%d] = 0x%lx\n", i, src[i]); + // printf("dst[%d] = 0x%lx\n", i, dst[i]); if (d != (s + 1)) { errors++; printf("mismatch %x != 1 + %x\n", d, s); @@ -161,7 +162,7 @@ int main(int argc, char *argv[]) { // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/dynamic_dma_config_add_one_kernel_obj_fifo/test.cpp b/reference_designs/dynamic_dma_config_add_one_kernel_obj_fifo/test.cpp index 7d89631970..68b523f873 100644 --- a/reference_designs/dynamic_dma_config_add_one_kernel_obj_fifo/test.cpp +++ b/reference_designs/dynamic_dma_config_add_one_kernel_obj_fifo/test.cpp @@ -18,11 +18,11 @@ #include #include -#include "test_library.h" #include "memory_allocator.h" +#include "test_library.h" -#include "aie_inc.cpp" #include "aie_data_movement.cpp" +#include "aie_inc.cpp" #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" @@ -53,8 +53,9 @@ int main(int argc, char *argv[]) { // and get an agent int ret = mlir_aie_init_device(xaie); - if(ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" << std::endl; + if (ret) { + std::cout << "[ERROR] Error when calling mlir_aie_init_device)" + << std::endl; return -1; } @@ -114,7 +115,7 @@ int main(int argc, char *argv[]) { // destroying the queue hsa_queue_destroy(queues[0]); - + // Shutdown AIR and HSA mlir_aie_deinit_libxaie(xaie); diff --git a/reference_designs/ipu-xrt/matrix_add_one/test.cpp b/reference_designs/ipu-xrt/matrix_add_one/test.cpp index f8c6626922..c8804fecfc 100644 --- a/reference_designs/ipu-xrt/matrix_add_one/test.cpp +++ b/reference_designs/ipu-xrt/matrix_add_one/test.cpp @@ -179,15 +179,16 @@ int main(int argc, const char *argv[]) { uint32_t d = bufOut[i]; if (row < TILE_HEIGHT && col < TILE_WIDTH) { - if(d != s + 1) { + if (d != s + 1) { errors++; printf("[ERROR] row %d and col %d, %d != %d\n", row, col, s, d); } - } - else { - if(d == s + 1) { + } else { + if (d == s + 1) { errors++; - printf("[ERROR] row %d and col %d, %d == %d -- this was not supposed to be changed\n", row, col, s, d); + printf("[ERROR] row %d and col %d, %d == %d -- this was not supposed " + "to be changed\n", + row, col, s, d); } } diff --git a/reference_designs/ipu-xrt/vector_max/test.cpp b/reference_designs/ipu-xrt/vector_max/test.cpp index 34d28e220f..ffe277e862 100644 --- a/reference_designs/ipu-xrt/vector_max/test.cpp +++ b/reference_designs/ipu-xrt/vector_max/test.cpp @@ -151,7 +151,7 @@ int main(int argc, const char *argv[]) { srcVecA.push_back(i + 1); memcpy(bufInA, srcVecA.data(), (srcVecA.size() * sizeof(uint32_t))); - bufInA[IN_SIZE / 2] = 654321; + bufInA[IN_SIZE / 2] = 654321; bufInA[IN_SIZE - 1] = 100; void *bufInstr = bo_instr.map(); @@ -170,7 +170,7 @@ int main(int argc, const char *argv[]) { uint32_t *bufOut = bo_out.map(); int errors = 0; - + uint32_t max_val = 0; for (uint32_t i = 0; i < IN_SIZE; i++) { if (*(bufInA + i) > max_val) { @@ -178,8 +178,9 @@ int main(int argc, const char *argv[]) { } } - if(*bufOut != max_val) { - std::cout << "[ERROR] Maximum value is " << max_val << " but kernel returned " << *bufOut << "\n"; + if (*bufOut != max_val) { + std::cout << "[ERROR] Maximum value is " << max_val + << " but kernel returned " << *bufOut << "\n"; errors++; } diff --git a/reference_designs/ipu-xrt/vector_min/test.cpp b/reference_designs/ipu-xrt/vector_min/test.cpp index a70649c16e..61339fc01e 100644 --- a/reference_designs/ipu-xrt/vector_min/test.cpp +++ b/reference_designs/ipu-xrt/vector_min/test.cpp @@ -151,7 +151,7 @@ int main(int argc, const char *argv[]) { srcVecA.push_back(i + 1); memcpy(bufInA, srcVecA.data(), (srcVecA.size() * sizeof(uint32_t))); - bufInA[IN_SIZE / 2] = 4; + bufInA[IN_SIZE / 2] = 4; bufInA[IN_SIZE - 1] = 100; void *bufInstr = bo_instr.map(); @@ -170,7 +170,7 @@ int main(int argc, const char *argv[]) { uint32_t *bufOut = bo_out.map(); int errors = 0; - + uint32_t min_val = 1000000; for (uint32_t i = 0; i < IN_SIZE; i++) { if (*(bufInA + i) < min_val) { @@ -178,8 +178,9 @@ int main(int argc, const char *argv[]) { } } - if(*bufOut != min_val) { - std::cout << "[ERROR] Min value is " << min_val << " but kernel returned " << *bufOut << "\n"; + if (*bufOut != min_val) { + std::cout << "[ERROR] Min value is " << min_val << " but kernel returned " + << *bufOut << "\n"; errors++; } diff --git a/reference_designs/ipu-xrt/vector_sum/test.cpp b/reference_designs/ipu-xrt/vector_sum/test.cpp index 64a8bf79e8..b87b46b888 100644 --- a/reference_designs/ipu-xrt/vector_sum/test.cpp +++ b/reference_designs/ipu-xrt/vector_sum/test.cpp @@ -151,7 +151,7 @@ int main(int argc, const char *argv[]) { srcVecA.push_back(i + 1); memcpy(bufInA, srcVecA.data(), (srcVecA.size() * sizeof(uint32_t))); - bufInA[IN_SIZE / 2] = 654321; + bufInA[IN_SIZE / 2] = 654321; bufInA[IN_SIZE - 1] = 100; void *bufInstr = bo_instr.map(); @@ -170,14 +170,15 @@ int main(int argc, const char *argv[]) { uint32_t *bufOut = bo_out.map(); int errors = 0; - + uint32_t sum_val = 0; for (uint32_t i = 0; i < IN_SIZE; i++) { sum_val += *(bufInA + i); } - if(*bufOut != sum_val) { - std::cout << "[ERROR] Maximum value is " << sum_val << " but kernel returned " << *bufOut << "\n"; + if (*bufOut != sum_val) { + std::cout << "[ERROR] Maximum value is " << sum_val + << " but kernel returned " << *bufOut << "\n"; errors++; } diff --git a/runtime_lib/test_lib/test_library.cpp b/runtime_lib/test_lib/test_library.cpp index eaca569d32..8d214eb5c0 100644 --- a/runtime_lib/test_lib/test_library.cpp +++ b/runtime_lib/test_lib/test_library.cpp @@ -26,7 +26,7 @@ #ifdef HSA_RUNTIME hsa_status_t mlir_aie_packet_req_translation(hsa_agent_dispatch_packet_t *pkt, - uint64_t va) { + uint64_t va) { pkt->arg[0] = 0; pkt->arg[0] = va; @@ -37,14 +37,13 @@ hsa_status_t mlir_aie_packet_req_translation(hsa_agent_dispatch_packet_t *pkt, return HSA_STATUS_SUCCESS; } -hsa_status_t -mlir_aie_packet_nd_memcpy(hsa_agent_dispatch_packet_t *pkt, uint16_t herd_id, - uint8_t col, uint8_t direction, uint8_t channel, - uint8_t burst_len, uint8_t memory_space, - uint64_t phys_addr, uint32_t transfer_length1d, - uint32_t transfer_length2d, uint32_t transfer_stride2d, - uint32_t transfer_length3d, uint32_t transfer_stride3d, - uint32_t transfer_length4d, uint32_t transfer_stride4d) { +hsa_status_t mlir_aie_packet_nd_memcpy( + hsa_agent_dispatch_packet_t *pkt, uint16_t herd_id, uint8_t col, + uint8_t direction, uint8_t channel, uint8_t burst_len, uint8_t memory_space, + uint64_t phys_addr, uint32_t transfer_length1d, uint32_t transfer_length2d, + uint32_t transfer_stride2d, uint32_t transfer_length3d, + uint32_t transfer_stride3d, uint32_t transfer_length4d, + uint32_t transfer_stride4d) { pkt->arg[0] = 0; pkt->arg[0] |= ((uint64_t)memory_space) << 16; @@ -106,10 +105,9 @@ hsa_status_t get_global_mem_pool(hsa_amd_memory_pool_t pool, void *data) { return status; } -hsa_status_t mlir_aie_queue_dispatch_and_wait(hsa_agent_t *agent, hsa_queue_t *q, - uint64_t packet_id, uint64_t doorbell, - hsa_agent_dispatch_packet_t *pkt, - bool destroy_signal) { +hsa_status_t mlir_aie_queue_dispatch_and_wait( + hsa_agent_t *agent, hsa_queue_t *q, uint64_t packet_id, uint64_t doorbell, + hsa_agent_dispatch_packet_t *pkt, bool destroy_signal) { // dispatch and wait has blocking semantics so we can internally create the // signal @@ -137,7 +135,7 @@ hsa_status_t mlir_aie_queue_dispatch_and_wait(hsa_agent_t *agent, hsa_queue_t *q } hsa_status_t mlir_aie_packet_device_init(hsa_agent_dispatch_packet_t *pkt, - uint32_t num_cols) { + uint32_t num_cols) { pkt->arg[0] = 0; pkt->arg[0] |= (AIR_ADDRESS_ABSOLUTE_RANGE << 48); @@ -226,7 +224,7 @@ int mlir_aie_init_device(aie_libxaie_ctx_t *ctx, uint32_t device_id) { hsa_agent_dispatch_packet_t shim_pkt; mlir_aie_packet_device_init(&shim_pkt, 50); mlir_aie_queue_dispatch_and_wait(&(ctx->agents[0]), q, packet_id, wr_idx, - &shim_pkt, true); + &shim_pkt, true); // Attaching the queue to the context so we can send more packets if needed ctx->cmd_queue = q; diff --git a/runtime_lib/test_lib/test_library.h b/runtime_lib/test_lib/test_library.h index fc7cd5751c..115a31cd49 100644 --- a/runtime_lib/test_lib/test_library.h +++ b/runtime_lib/test_lib/test_library.h @@ -20,12 +20,11 @@ #include "hsa_ext_air.h" #endif - #ifdef HSA_RUNTIME - template - inline void mlir_aie_write_pkt(hsa_queue_t *q, uint32_t packet_id, T *pkt) { - reinterpret_cast(q->base_address)[packet_id] = *pkt; - } +template +inline void mlir_aie_write_pkt(hsa_queue_t *q, uint32_t packet_id, T *pkt) { + reinterpret_cast(q->base_address)[packet_id] = *pkt; +} #endif extern "C" { @@ -126,24 +125,19 @@ class EventMonitor { // This is a more elegant solution #ifdef HSA_RUNTIME hsa_status_t mlir_aie_packet_req_translation(hsa_agent_dispatch_packet_t *pkt, - uint64_t va); - -hsa_status_t -mlir_aie_packet_nd_memcpy(hsa_agent_dispatch_packet_t *pkt, uint16_t herd_id, - uint8_t col, uint8_t direction, uint8_t channel, - uint8_t burst_len, uint8_t memory_space, - uint64_t phys_addr, uint32_t transfer_length1d, - uint32_t transfer_length2d, uint32_t transfer_stride2d, - uint32_t transfer_length3d, uint32_t transfer_stride3d, - uint32_t transfer_length4d, uint32_t transfer_stride4d); - -hsa_status_t mlir_aie_queue_dispatch_and_wait(hsa_agent_t *agent, hsa_queue_t *q, - uint64_t packet_id, uint64_t doorbell, - hsa_agent_dispatch_packet_t *pkt, - bool destroy_signal = true); - - - + uint64_t va); + +hsa_status_t mlir_aie_packet_nd_memcpy( + hsa_agent_dispatch_packet_t *pkt, uint16_t herd_id, uint8_t col, + uint8_t direction, uint8_t channel, uint8_t burst_len, uint8_t memory_space, + uint64_t phys_addr, uint32_t transfer_length1d, uint32_t transfer_length2d, + uint32_t transfer_stride2d, uint32_t transfer_length3d, + uint32_t transfer_stride3d, uint32_t transfer_length4d, + uint32_t transfer_stride4d); + +hsa_status_t mlir_aie_queue_dispatch_and_wait( + hsa_agent_t *agent, hsa_queue_t *q, uint64_t packet_id, uint64_t doorbell, + hsa_agent_dispatch_packet_t *pkt, bool destroy_signal = true); #endif diff --git a/runtime_lib/xaiengine/aie-rt b/runtime_lib/xaiengine/aie-rt index 4f8a992c5b..caee13c40d 160000 --- a/runtime_lib/xaiengine/aie-rt +++ b/runtime_lib/xaiengine/aie-rt @@ -1 +1 @@ -Subproject commit 4f8a992c5b46542592a7a213664daef4c259c09a +Subproject commit caee13c40d0c4c8d022bddd311dea0d9f75503fd