Add: add scalar float type data dump support.

docs/dfx/tensor-dump.md

@@ -26,9 +26,11 @@ saw, without the timing distortion of inline printing.
 - **Manifest + binary payload.** `--dump-tensor` writes a JSON
   manifest plus one `.bin` payload per run; tensor entries carry
   `bin_offset` / `bin_size`, while scalar entries stay manifest-only.
-- **Unified scalar args.** Scalar callable slots are emitted as
+- **Unified scalar args.** Scalar values are emitted as
   `kind: scalar`, `stage: before_dispatch`, zero-dim records in
-  `tensor_dump.json`; there is no separate args-only manifest.
+  `tensor_dump.json`; there is no separate args-only manifest. Their
+  final `dtype` is registered at submit time in a dump-only per-task side table
+  and resolved by AICPU when writing each scalar dump record.
 - **Cross-architecture.** Same flags and on-disk layout family on
   `a2a3` and `a5`. Both runtimes are wired through.
 
@@ -274,8 +276,9 @@ What you can read out of `tensor_dump.json` + `tensor_dump.bin`:
 - **Per-task input snapshots** (`role: input`, `stage:
   before_dispatch`) — what each kernel was given.
 - **Per-dispatch scalar args** (`kind: scalar`, `stage:
-  before_dispatch`) — the raw callable-slot scalar values AICPU
-  handed to the kernel.
+  before_dispatch`) — the raw per-task scalar-slot values AICPU
+  handed to the kernel, tagged with the dump-only dtype captured at
+  submit time.
 - **Per-task output snapshots** (`role: output`, `stage:
   after_completion`) — what each kernel produced. The barrier
   ensures these reflect the kernel's final writes.
@@ -284,8 +287,8 @@ What you can read out of `tensor_dump.json` + `tensor_dump.bin`:
 - **Logical view reconstruction** — `shape` / `strides` /
   `start_offset` / `is_contiguous` plus the gathered
   logical-contiguous payload.
-- **Per-task identity** — `task_id` / `subtask_id` / `func_id`
-  correlates dump entries with swimlane and PMU rows.
+- **Per-task identity** — `task_id`, `stage`, `role`, and `arg_index`
+  identify each dumped argument within a task.
 - **Loss accounting** — `truncated` / `overwritten` per-record
   flags, plus aggregate `dropped_records` / `dropped_overwrite` in
   the summary.
@@ -361,10 +364,13 @@ Each runtime's scheduler dispatch code calls
 └──────────────────────────────────────┘
 ```
 
-`dump_tensors_for_task` walks the formal callable signature,
-matches each non-scalar slot to a `TensorDumpInfo` (dtype + shape +
-strides + start offset + device address), and calls
-`dump_tensor_record` for slots that match the current stage.
+`dump_tensors_for_task` walks the formal callable signature for
+non-scalar tensor slots, matches each tensor slot to a
+`TensorDumpInfo` (dtype + shape + strides + start offset + device
+address), and calls `dump_tensor_record` for slots that match the
+current stage. Scalar values are dumped separately from the flat
+payload `scalars[]`; their dtype table is registered at submit time
+and emitted as a dump-only metadata record at `BEFORE_DISPATCH`.
 
 When dump is enabled, AICore executors also issue
 `pipe_barrier(PIPE_ALL)` after kernel execution and before writing

simpler_setup/tools/dump_viewer.py

@@ -51,8 +51,11 @@
     "int64": ("q", 8),
     "uint64": ("Q", 8),
     "int16": ("h", 2),
+    "uint16": ("H", 2),
     "int8": ("b", 1),
     "uint8": ("B", 1),
+    "uint32": ("I", 4),
+    "bool": ("?", 1),
 }
 
 
@@ -103,6 +106,7 @@ def write_tensor(tensor: dict, bin_path: Path, out):
     out.write(f"# stage: {t['stage']}\n")
     out.write(f"# arg_index: {t['arg_index']}\n")
     out.write(f"# dtype: {t['dtype']}\n")
+    out.write(f"# kind: {t.get('kind', 'tensor')}\n")
     out.write(f"# is_contiguous: {t['is_contiguous']}\n")
     out.write(f"# shape: {t['shape']}\n")
     out.write(f"# strides: {t['strides']}\n")
@@ -114,6 +118,13 @@ def write_tensor(tensor: dict, bin_path: Path, out):
     if t.get("truncated"):
         out.write("# DATA TRUNCATED (tensor too large for arena)\n")
 
+    if t.get("kind") == "scalar":
+        val = t.get("value")
+        if t.get("dtype", "").upper() == "BOOL":
+            val = "true" if val else "false"
+        out.write(f"# value: {val}\n")
+        return
+
     bin_size = t.get("bin_size", 0)
     if bin_size == 0:
         out.write("# (no data)\n")

src/a2a3/runtime/tensormap_and_ringbuffer/runtime/pto_orchestrator.cpp

@@ -36,10 +36,9 @@
 #include "pto_types.h"
 #include "tensor.h"
 
-extern "C" void set_dump_tensor_task_mask(uint64_t task_id, uint64_t mask);
-
 #if PTO2_PROFILING
 #include "aicpu/scope_stats_collector_aicpu.h"
+#include "aicpu/tensor_dump_aicpu.h"
 #endif
 
 // Verify the captured Tensor blob size in DepGenRecord matches the runtime
@@ -648,12 +647,19 @@ static TaskOutputTensors submit_task_common(
 
     payload.init(args, result, prepared.alloc_result, layout);
 #if PTO2_PROFILING
-    // Selective vs full dump is latched at dump_tensor_init from DumpDataHeader
-    // (host-decided before any dispatch), so it is race-free regardless of
-    // submission order. Here we only record each marked task's arg mask, which
-    // selective collection consults.
-    if (args.tensor_dump_arg_mask() != 0) {
-        set_dump_tensor_task_mask(task_id.raw, args.tensor_dump_arg_mask());
+    if (is_dump_tensor_enabled()) {
+        if (args.scalar_count() > 0) {
+            set_dump_tensor_task_scalar_dtypes(
+                task_id.raw, static_cast<uint32_t>(args.scalar_count()), args.scalar_dtypes()
+            );
+        }
+        // Selective vs full dump is latched at dump_tensor_init from DumpDataHeader
+        // (host-decided before any dispatch), so it is race-free regardless of
+        // submission order. Here we only record each marked task's arg mask, which
+        // selective collection consults.
+        if (args.tensor_dump_arg_mask() != 0) {
+            set_dump_tensor_task_mask(task_id.raw, args.tensor_dump_arg_mask());
+        }
     }
 #endif
 

src/a2a3/runtime/tensormap_and_ringbuffer/runtime/pto_types.h

@@ -32,6 +32,7 @@
 #include <arm_neon.h>
 #endif
 
+#include "data_type.h"
 #include "pto_submit_types.h"
 #include "task_args.h"
 #include "tensor.h"
@@ -183,6 +184,7 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MAX_TENSOR_ARGS, MAX_SCALAR_ARGS,
         tensor_dump_arg_mask_ = 0;
         explicit_deps_ = nullptr;
         explicit_dep_count_ = 0;
+        memset(scalar_dtypes_, 0, sizeof(scalar_dtypes_));
     }
 
     void set_error(const char *msg) {
@@ -312,7 +314,9 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MAX_TENSOR_ARGS, MAX_SCALAR_ARGS,
             set_error("Too many scalar args (exceeds MAX_SCALAR_ARGS=32)");
             return;
         }
-        ((scalars_[scalar_count_++] = to_u64(args)), ...);
+        ((scalars_[scalar_count_] = to_u64(args),
+          scalar_dtypes_[scalar_count_] = dtype_of<std::remove_cv_t<std::remove_reference_t<Args>>>(), ++scalar_count_),
+         ...);
     }
 
     void add_scalars(const uint64_t *values, int count) {
@@ -370,15 +374,19 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MAX_TENSOR_ARGS, MAX_SCALAR_ARGS,
             return;
         }
         memcpy(&scalars_[scalar_count_], &src.scalars_[src_offset], count * sizeof(uint64_t));
+        memcpy(&scalar_dtypes_[scalar_count_], &src.scalar_dtypes_[src_offset], count * sizeof(uint8_t));
         scalar_count_ += count;
     }
 
+    const uint8_t *scalar_dtypes() const { return scalar_dtypes_; }
+
 private:
     // Caller-owned dependency array; lifetime must extend through submit.
     static_assert(MAX_TENSOR_ARGS <= 64, "tensor dump arg mask assumes at most 64 tensor arguments");
     uint64_t tensor_dump_arg_mask_{0};
     const PTO2TaskId *explicit_deps_{nullptr};
     uint32_t explicit_dep_count_{0};
+    uint8_t scalar_dtypes_[MAX_SCALAR_ARGS] = {};
 
     // No-arg dump(): mark every tensor arg already added to this Arg.
     void mark_all_tensor_dump_arg() {

src/a5/runtime/tensormap_and_ringbuffer/runtime/pto_orchestrator.cpp

@@ -36,7 +36,9 @@
 #include "pto_types.h"
 #include "tensor.h"
 
-extern "C" void set_dump_tensor_task_mask(uint64_t task_id, uint64_t mask);
+#if PTO2_PROFILING
+#include "aicpu/tensor_dump_aicpu.h"
+#endif
 
 // Verify the captured Tensor blob size in DepGenRecord matches the runtime
 // Tensor layout. The platform header defines DEP_GEN_TENSOR_SIZE without
@@ -647,12 +649,19 @@ static TaskOutputTensors submit_task_common(
 
     payload.init(args, result, prepared.alloc_result, layout);
 #if PTO2_PROFILING
-    // Selective vs full dump is latched at dump_tensor_init from DumpDataHeader
-    // (host-decided before any dispatch), so it is race-free regardless of
-    // submission order. Here we only record each marked task's arg mask, which
-    // selective collection consults.
-    if (args.tensor_dump_arg_mask() != 0) {
-        set_dump_tensor_task_mask(task_id.raw, args.tensor_dump_arg_mask());
+    if (is_dump_tensor_enabled()) {
+        if (args.scalar_count() > 0) {
+            set_dump_tensor_task_scalar_dtypes(
+                task_id.raw, static_cast<uint32_t>(args.scalar_count()), args.scalar_dtypes()
+            );
+        }
+        // Selective vs full dump is latched at dump_tensor_init from DumpDataHeader
+        // (host-decided before any dispatch), so it is race-free regardless of
+        // submission order. Here we only record each marked task's arg mask, which
+        // selective collection consults.
+        if (args.tensor_dump_arg_mask() != 0) {
+            set_dump_tensor_task_mask(task_id.raw, args.tensor_dump_arg_mask());
+        }
     }
 #endif
 

src/a5/runtime/tensormap_and_ringbuffer/runtime/pto_types.h

@@ -28,12 +28,13 @@
 #include <stdint.h>
 #include <string.h>
 
-#include <utility>
-
 #if defined(__aarch64__)
 #include <arm_neon.h>
 #endif
 
+#include <utility>
+
+#include "data_type.h"
 #include "pto_submit_types.h"
 #include "task_args.h"
 #include "tensor.h"
@@ -181,6 +182,7 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MAX_TENSOR_ARGS, MAX_SCALAR_ARGS,
         tensor_dump_arg_mask_ = 0;
         explicit_deps_ = nullptr;
         explicit_dep_count_ = 0;
+        memset(scalar_dtypes_, 0, sizeof(scalar_dtypes_));
     }
 
     void reset() {
@@ -316,7 +318,9 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MAX_TENSOR_ARGS, MAX_SCALAR_ARGS,
             set_error("Too many scalar args (exceeds MAX_SCALAR_ARGS=32)");
             return;
         }
-        ((scalars_[scalar_count_++] = to_u64(args)), ...);
+        ((scalars_[scalar_count_] = to_u64(args),
+          scalar_dtypes_[scalar_count_] = dtype_of<std::remove_cv_t<std::remove_reference_t<Args>>>(), ++scalar_count_),
+         ...);
     }
 
     void add_scalars(const uint64_t *values, int count) {
@@ -374,15 +378,19 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MAX_TENSOR_ARGS, MAX_SCALAR_ARGS,
             return;
         }
         memcpy(&scalars_[scalar_count_], &src.scalars_[src_offset], count * sizeof(uint64_t));
+        memcpy(&scalar_dtypes_[scalar_count_], &src.scalar_dtypes_[src_offset], count * sizeof(uint8_t));
         scalar_count_ += count;
     }
 
+    const uint8_t *scalar_dtypes() const { return scalar_dtypes_; }
+
 private:
     // Caller-owned dependency array; lifetime must extend through submit.
     static_assert(MAX_TENSOR_ARGS <= 64, "tensor dump arg mask assumes at most 64 tensor arguments");
     uint64_t tensor_dump_arg_mask_{0};
     const PTO2TaskId *explicit_deps_{nullptr};
     uint32_t explicit_dep_count_{0};
+    uint8_t scalar_dtypes_[MAX_SCALAR_ARGS] = {};
 
     // No-arg dump(): mark every tensor arg already added to this Arg.
     void mark_all_tensor_dump_arg() {

src/common/platform/include/aicpu/tensor_dump_aicpu.h

@@ -76,6 +76,8 @@ void set_dump_tensor_selective_mode(bool enable);
 bool is_dump_tensor_selective_mode();
 void set_dump_tensor_task_mask(uint64_t task_id, TensorDumpArgMask mask);
 TensorDumpArgMask get_dump_tensor_task_mask(uint64_t task_id);
+void set_dump_tensor_task_scalar_dtypes(uint64_t task_id, uint32_t scalar_count, const uint8_t *scalar_dtypes);
+bool get_dump_tensor_task_scalar_dtypes(uint64_t task_id, uint32_t *scalar_count, uint8_t *scalar_dtypes);
 
 #ifdef __cplusplus
 }
@@ -123,8 +125,8 @@ inline void dump_tensors_for_task(
         if (try_log_tensor_dump_layout_mismatch()) {
             LOG_WARN(
                 "Thread %d: tensor dump skipped for task 0x%" PRIx64
-                ": active callable tensor count (%d) does not match payload tensor count (%d). "
-                "Task-level dump assumes payload tensors are concatenated by active subtask order.",
+                ": active callable tensor args (%d) do not match payload tensor args (%d). "
+                "Task-level dump assumes payload tensor args are concatenated by active subtask order.",
                 thread_idx, static_cast<uint64_t>(slot_state.task->task_id.raw), total_tensor_args, pl.tensor_count
             );
         }
@@ -172,19 +174,26 @@ inline void dump_tensors_for_task(
             arg_index++;
             tensor_index++;
         }
-        if (stage == TensorDumpStage::BEFORE_DISPATCH && pl.scalar_count > 0) {
-            for (int32_t i = 0; i < pl.scalar_count; i++) {
-                TensorDumpInfo info = {};
-                info.task_id = slot_state.task->task_id.raw;
-                info.role = TensorDumpRole::INPUT;
-                info.stage = stage;
-                info.dtype = static_cast<uint8_t>(DataType::UINT64);
-                info.ndims = 0;
-                info.arg_index = static_cast<uint32_t>(pl.tensor_count + i);
-                info.kind = static_cast<uint8_t>(TensorDumpKind::SCALAR);
-                info.scalar_value = pl.scalars[i];
-                dump_tensor_record(thread_idx, info);
-            }
+    }
+
+    if (stage == TensorDumpStage::BEFORE_DISPATCH && pl.scalar_count > 0) {
+        uint8_t scalar_dtypes[CORE_MAX_SCALAR_ARGS] = {};
+        uint32_t dtype_scalar_count = 0;
+        bool has_scalar_dtypes =
+            get_dump_tensor_task_scalar_dtypes(slot_state.task->task_id.raw, &dtype_scalar_count, scalar_dtypes);
+        for (int32_t scalar_index = 0; scalar_index < pl.scalar_count; scalar_index++) {
+            TensorDumpInfo info = {};
+            info.task_id = slot_state.task->task_id.raw;
+            info.role = TensorDumpRole::INPUT;
+            info.stage = stage;
+            info.dtype = (has_scalar_dtypes && scalar_index < static_cast<int32_t>(dtype_scalar_count)) ?
+                             scalar_dtypes[scalar_index] :
+                             static_cast<uint8_t>(DataType::UINT64);
+            info.ndims = 0;
+            info.arg_index = static_cast<uint32_t>(pl.tensor_count + scalar_index);
+            info.kind = static_cast<uint8_t>(TensorDumpKind::SCALAR);
+            info.scalar_value = pl.scalars[scalar_index];
+            dump_tensor_record(thread_idx, info);
         }
     }
 }