rdcycle_sync_diff & simx changes

Author: EricSongXinLe

kernel/include/vx_intrinsics.h

@@ -284,6 +284,94 @@ static inline uint64_t vx_rdcycle_sync() {
     return cycles;
 }
 
+typedef struct {
+    uint32_t hi;
+    uint32_t lo;
+} __rdcycle_time;
+
+static inline __attribute__((always_inline)) __rdcycle_time vx_rdcycle_sync_begin(void) {
+    __rdcycle_time t;
+#if __riscv_xlen == 32
+    __asm__ volatile (
+        ".insn r %2, 7, 0, x0, x0, x0\n\t"
+        "csrr %0, %3\n\t"
+        "csrr %1, %4\n\t"
+        : "=r" (t.hi), "=r" (t.lo)
+        : "i" (RISCV_CUSTOM0), "i" (VX_CSR_MCYCLE_H), "i" (VX_CSR_MCYCLE)
+        : "memory"
+    );
+#elif __riscv_xlen == 64
+    uint64_t cycles;
+    __asm__ volatile (
+        ".insn r %1, 7, 0, x0, x0, x0\n\t"
+        "csrr %0, %2\n\t"
+        : "=r" (cycles)
+        : "i" (RISCV_CUSTOM0), "i" (VX_CSR_MCYCLE)
+        : "memory"
+    );
+    t.hi = (uint32_t)(cycles >> 32);
+    t.lo = (uint32_t)cycles;
+#else
+#error "Unsupported RISC-V XLEN"
+#endif
+    return t;
+}
+
+static inline __attribute__((always_inline)) __rdcycle_time vx_rdcycle_sync_end(void) {
+    __rdcycle_time t;
+#if __riscv_xlen == 32
+    __asm__ volatile (
+        ".insn r %2, 7, 0, x0, x0, x0\n\t"
+        "csrr %0, %3\n\t"
+        "csrr %1, %4\n\t"
+        : "=r" (t.lo), "=r" (t.hi)
+        : "i" (RISCV_CUSTOM0), "i" (VX_CSR_MCYCLE), "i" (VX_CSR_MCYCLE_H)
+        : "memory"
+    );
+#elif __riscv_xlen == 64
+    uint64_t cycles;
+    __asm__ volatile (
+        ".insn r %1, 7, 0, x0, x0, x0\n\t"
+        "csrr %0, %2\n\t"
+        : "=r" (cycles)
+        : "i" (RISCV_CUSTOM0), "i" (VX_CSR_MCYCLE)
+        : "memory"
+    );
+    t.hi = (uint32_t)(cycles >> 32);
+    t.lo = (uint32_t)cycles;
+#else
+#error "Unsupported RISC-V XLEN"
+#endif
+    return t;
+}
+
+static inline __attribute__((always_inline)) uint64_t vx_rdcycle_sync_diff(__rdcycle_time start, __rdcycle_time end) {
+#if __riscv_xlen == 32
+    uint32_t diff_hi = end.hi;
+    uint32_t diff_lo = end.lo;
+    uint32_t tmp = start.hi;
+    uint32_t start_lo = start.lo;
+
+    __asm__ volatile (
+        "sub  %0, %0, %2\n\t"
+        "sltu %2, %1, %3\n\t"
+        "sub  %1, %1, %3\n\t"
+        "sub  %0, %0, %2\n\t"
+        : "+r" (diff_hi), "+r" (diff_lo), "+r" (tmp)
+        : "r" (start_lo)
+        : "memory"
+    );
+
+    return ((uint64_t)diff_hi << 32) | diff_lo;
+#elif __riscv_xlen == 64
+    uint64_t s = ((uint64_t)start.hi << 32) | start.lo;
+    uint64_t e = ((uint64_t)end.hi << 32) | end.lo;
+    return e - s;
+#else
+#error "Unsupported RISC-V XLEN"
+#endif
+}
+
 // Memory fence
 inline void vx_fence() {
     __asm__ volatile ("fence iorw, iorw");

sim/simx/core.cpp

@@ -476,11 +476,6 @@ void Core::commit() {
 
     // update scoreboard
     if (trace->eop) {
-      bool release_wsync = false;
-      if (auto wctl_type = std::get_if<WctlType>(&trace->op_type)) {
-        release_wsync = (*wctl_type == WctlType::WSYNC);
-      } //if the instruction is a wsync, set release_wsync to true.  warp will continue to run after commits
-
       if (trace->wb) {
         operands_.at(iw)->writeback(trace);
         scoreboard_.release(trace);
@@ -511,10 +506,6 @@ void Core::commit() {
     #endif
       // instruction completed
       pending_instrs_.remove(trace);
-
-      if (release_wsync) {
-        this->resume(trace->wid); //allows the scheduler to select this warp to fetch new instructions
-      }
     }
 
     // delete the trace
@@ -536,11 +527,10 @@ bool Core::running() const {
 }
 
 bool Core::warp_sync_ready(uint32_t wid, uint64_t uuid) const {
+  (void)uuid;
   for (auto trace : pending_instrs_) {
     if (trace->wid != wid)
       continue;
-    if (trace->uuid < uuid)
-      return false;
   }
   return true;
 }

sim/simx/func_unit.cpp

@@ -459,7 +459,7 @@ void SfuUnit::tick() {
 				if (trace->eop && !core_->warp_sync_ready(trace->wid, trace->uuid))
 					continue; //skips the rest of the loop. Does not pop the input.
 				output.send(trace, 2+delay);
-				release_warp = false;
+				release_warp = true;
 				break;
 			case WctlType::BAR: {
 				output.send(trace, 2+delay);

tests/regression/sgemm_tcu/kernel.cpp

@@ -1,6 +1,7 @@
 #include "common.h"
 #include <vx_spawn.h>
 #include <vx_tensor.h>
+#include <vx_intrinsics.h>
 
 namespace vt = vortex::tensor;
 using ctx = vt::wmma_context<NUM_THREADS, vt::ITYPE, vt::OTYPE>;
@@ -9,6 +10,7 @@ void kernel_body(kernel_arg_t *__UNIFORM__ arg) {
   auto pA = reinterpret_cast<ctx::input_t *>(arg->A_addr);
   auto pB = reinterpret_cast<ctx::input_t *>(arg->B_addr);
   auto pC = reinterpret_cast<ctx::output_t *>(arg->C_addr);
+  auto pCycles = reinterpret_cast<uint64_t *>(arg->cycles_addr);
 
   uint32_t M = arg->M;
   uint32_t N = arg->N;
@@ -25,6 +27,7 @@ void kernel_body(kernel_arg_t *__UNIFORM__ arg) {
   // Initialize accumulator tile to zero
   ctx::fill_fragment(fragC, 0);
 
+  uint64_t cycles = 0;
   for (int i = 0; i < K; i += ctx::tileK) {
     auto pTileA = pA + tile_row * K + i;
 
@@ -42,7 +45,15 @@ void kernel_body(kernel_arg_t *__UNIFORM__ arg) {
     }
 
     // Matrix multiply-accumulate: c += a * b
+    __rdcycle_time t0 = vx_rdcycle_sync_begin();
     ctx::mma_sync(fragC, fragA, fragB, fragC);
+    __rdcycle_time t1 = vx_rdcycle_sync_end();
+    cycles += vx_rdcycle_sync_diff(t0, t1);
+  }
+
+  if (0 == vx_thread_id()) {
+    uint32_t block_id = blockIdx.y * arg->grid_dim[0] + blockIdx.x;
+    pCycles[block_id] = cycles;
   }
 
   // Store the computed C tile