clang-format

Xilinx · Mar 20, 2024 · a0cecc4 · a0cecc4
1 parent 957b7d9
commit a0cecc4
Show file tree

Hide file tree

Showing 30 changed files with 240 additions and 206 deletions.
diff --git 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
@@ -56,7 +56,6 @@ getAllocOpForSymbol(AIE::DeviceOp dev, StringRef sym_name) {
   return std::nullopt;
 }
 
-
 mlir::LogicalResult AIETranslateToHSA(ModuleOp module, raw_ostream &output) {
 
   DenseMap<TileID, Operation *> 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<DeviceOp>().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<mlir::func::FuncOp>().empty()) {
+  if (targetOp.getOps<mlir::func::FuncOp>().empty()) {
     return success();
   }
 
   // Getting the sequence function op which contains the instructions
   mlir::func::FuncOp funcOp;
   for (auto op : targetOp.getOps<mlir::func::FuncOp>()) {
-    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<IpuDmaMemcpyNdOp>()) {
     // 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<IpuDmaMemcpyNdOp>()) {
     auto dev = funcOp->getParentOfType<AIE::DeviceOp>();
     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<int64_t, 3> 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<int64_t, 4> 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<int64_t, 4> 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<hsa_agent_dispatch_packet_t>(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<hsa_agent_dispatch_packet_t>(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
@@ -19,11 +19,11 @@
 #include <vector>
 #include <xaiengine.h>
 
-#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
@@ -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
@@ -19,11 +19,11 @@
 #include <vector>
 #include <xaiengine.h>
 
-#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
@@ -19,11 +19,11 @@
 #include <vector>
 #include <xaiengine.h>
 
-#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
@@ -19,11 +19,11 @@
 #include <vector>
 #include <xaiengine.h>
 
-#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
@@ -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;
     }
   }