Niko1221/Strata
1
0 Fork 0
mirror of https://github.com/Niko1221/Strata.git synced 2026-11-07 14:21:06 -10:00
Code Issues Packages Projects Releases 1 Wiki Activity

SYCL: STRATA_UC_PROBE also prints the live device allocations, GPU-stamped kernel times and the host launch cost, at nine points of start-up (test only, does nothing when unset)

Browse Source
This commit is contained in:
Niko1221 committed 2026-11-02 06:06:07 +01:00
1 parent f9e14bdf8e
commit 5123338339
2 files changed
+56 -4
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

sycl/include/strata/uc_probe.hpp
+47 -4
View File
@@ -13,11 +13,13 @@
#include <dpct/dpct.hpp>
#include "strata/sycl_queue.hpp"
#include "strata/kernels/native_mmvq.hpp"
#include <algorithm>
#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <string>
#include <vector>
namespace strata {
namespace ucprobe {
@@ -79,16 +81,18 @@ inline void uc_probe(const char* tag) {
{
ucprobe::Bufs b(q);
namespace k = strata::kernels;
const double head = ucprobe::time_us(q, [&] {
double head = 1e30, q6 = 1e30, fma = 1e30, tiny = 1e30;
for (int rep = 0; rep < 3; ++rep) { // the first pass warms the GPU clock; the line shows the best of three
const double head_ = ucprobe::time_us(q, [&] {
k::native_quantize_q8_1(b.x, b.xq, 2048, 7, &q);
k::native_mmvq(8, b.w8, b.xq, b.y, 2048, 16384, 7, &q);
}, 40);
const double q6 = ucprobe::time_us(q, [&] {
const double q6_ = ucprobe::time_us(q, [&] {
k::native_quantize_q8_1(b.x, b.xq, 2560, 4, &q);
k::native_mmvq(14, b.w6, b.xq, b.y, 2560, 10240, 4, &q);
}, 80);
float* fo = b.fo;
const double fma = ucprobe::time_us(q, [&] {
const double fma_ = ucprobe::time_us(q, [&] {
q.parallel_for(sycl::nd_range<1>(256 * 128, 128), [=](sycl::nd_item<1> it) {
float a = 1.0f + 1e-3f * (float) it.get_global_id(0);
for (int i = 0; i < 4096; ++i) a = a * 1.0000001f + 1e-7f;
@@ -96,13 +100,52 @@ inline void uc_probe(const char* tag) {
});
}, 40);
int* ti = b.ti;
const double tiny = ucprobe::time_us(q, [&] {
const double tiny_ = ucprobe::time_us(q, [&] {
q.parallel_for(sycl::nd_range<1>(64, 64), [=](sycl::nd_item<1> it) {
int a = (int) it.get_global_id(0);
for (int i = 0; i < 400; ++i) a = a * 3 + i;
ti[it.get_global_id(0)] = a;
});
}, 200);
head = std::min(head, head_); q6 = std::min(q6, q6_); fma = std::min(fma, fma_); tiny = std::min(tiny, tiny_);
}
double fma_g = -1, tiny_g = -1, empty_w = -1;
float* fo = b.fo; int* ti = b.ti;
try { // GPU stamps of the same two kernels (needs a profiling queue: STRATA_KPROF=1), and the host cost of an empty launch
auto gpu_us = [&](auto&& launch, int n) {
std::vector<sycl::event> ev; ev.reserve(n);
for (int i = 0; i < 5; ++i) launch();
q.wait();
for (int i = 0; i < n; ++i) ev.push_back(launch());
q.wait();
double g = 0;
for (auto& e : ev) g += (double) (e.get_profiling_info<sycl::info::event_profiling::command_end>() - e.get_profiling_info<sycl::info::event_profiling::command_start>());
return g / 1e3 / n;
};
fma_g = gpu_us([&] { return q.parallel_for(sycl::nd_range<1>(256 * 128, 128), [=](sycl::nd_item<1> it) {
float a = 1.0f + 1e-3f * (float) it.get_global_id(0);
for (int i = 0; i < 4096; ++i) a = a * 1.0000001f + 1e-7f;
fo[it.get_global_id(0)] = a; }); }, 40);
tiny_g = gpu_us([&] { return q.parallel_for(sycl::nd_range<1>(64, 64), [=](sycl::nd_item<1> it) {
int a = (int) it.get_global_id(0);
for (int i = 0; i < 400; ++i) a = a * 3 + i;
ti[it.get_global_id(0)] = a; }); }, 200);
} catch (...) {}
{
for (int i = 0; i < 5; ++i) q.single_task([] {});
q.wait();
const auto t0 = std::chrono::steady_clock::now();
for (int i = 0; i < 200; ++i) q.single_task([] {});
const double submit = std::chrono::duration<double, std::micro>(std::chrono::steady_clock::now() - t0).count() / 200;
q.wait();
empty_w = submit;
}
{
std::lock_guard<std::mutex> lk(strata::device_alloc_mutex());
size_t tot = 0; for (const auto& kv : strata::device_alloc_map()) tot += kv.second;
std::fprintf(stderr, "UCPROBE %-14s registered device allocations: %zu (%.2f GiB)\n", tag, strata::device_alloc_map().size(), (double) tot / 1073741824.0);
}
std::fprintf(stderr, "UCPROBE %-14s gpu-stamps: fma %7.1f us tiny %7.1f us host cost of an empty launch %6.1f us\n", tag, fma_g, tiny_g, empty_w);
std::fprintf(stderr, "UCPROBE %-14s head %8.1f us q6k %8.1f us fma %8.1f us tiny %7.1f us\n", tag, head, q6, fma, tiny);
}
if (calls++ == 0) {
sycl/src/program/generate.cpp
+9
View File
@@ -3043,6 +3043,7 @@ int main(int argc, char **argv) try {
std::fprintf(stderr, "strata generate: %llu MiB of weights loaded from %s in %.1f s (%zu canonical tensors "
"skipped: served natively)\n",
(unsigned long long) (pool_bytes >> 20), o.pack.c_str(), load_s(), skip.size());
strata::uc_probe("weights");
strata::core::NativeDense native_dense;
if (!o.native_dense_gguf.empty()) {
@@ -3054,6 +3055,7 @@ int main(int argc, char **argv) try {
native_dense.tensor_count(), (double) native_dense.weight_bytes() / (1024.0 * 1024.0), load_s());
}
strata::uc_probe("native-dense");
strata::kernels::gr_set_fp32_activations(o.gr_fp32_activations);
strata::kernels::gr_set_native_mmvf(o.gr_native_mmvf);
// Plan v0.3 P3: the fused hyper-connection read rides the native (FP32-activation) contract; the per-stage
@@ -3649,6 +3651,7 @@ int main(int argc, char **argv) try {
"behind rank %zu (the peer's)\n", f, n_moved, at, fill);
}
}
strata::uc_probe("after-profile");
std::fprintf(stderr, "strata generate: profile %s: %zu ranked pairs, built for %lld slots\n",
o.expert_profile.c_str(), profile.size(), (long long) pslots);
}
@@ -4338,6 +4341,7 @@ int main(int argc, char **argv) try {
(unsigned long long) native_head.weight_bytes(),
std::chrono::duration<double>(std::chrono::steady_clock::now() - head_t0).count());
}
strata::uc_probe("after-head");
std::vector<float> logits((size_t) n_vocab);
float* d_logits = nullptr;
if (DPCT_CHECK_ERROR(
@@ -5080,6 +5084,7 @@ int main(int argc, char **argv) try {
std::fprintf(stderr, "strata generate: %s\n", err.c_str());
return 1;
}
strata::uc_probe("after-fill");
mem_mark("the profile fill");
std::fprintf(stderr, "strata generate: pre-filled %lld of %lld slots from the profile in %.1f s (%.0f MB/s); "
"slot 0 verified\n",
@@ -5619,6 +5624,7 @@ int main(int argc, char **argv) try {
return 2;
}
strata::uc_probe("cache-done");
mem_mark("the expert cache and the graphs");
std::fprintf(stderr, "strata generate: session is up (engine %s)\n", STRATA_VERSION);
if (verify_all_slots) verify_all_slots("session up");
@@ -12345,6 +12351,7 @@ int main(int argc, char **argv) try {
// then starts at the last prompt position, whose prediction is the first generated token.
int64_t pos_start = 0;
int64_t spec_pos = 0; // plan v0.3 P6: where the speculative loop starts (0 = not used)
strata::uc_probe("before-prefill");
strata::prefill::Prefill prefill;
bool kvg_started = false; // the elastic K/V took this run's cells
double prefill_batched_ms = 0;
@@ -12459,6 +12466,7 @@ int main(int argc, char **argv) try {
}
prefill_batched_ms = std::chrono::duration<double, std::milli>(Clock::now() - tp0).count();
prefill_ms += prefill_batched_ms;
strata::uc_probe("after-prefill");
pos_start = n_batched;
tok = o.tokens[(size_t) pos_start];
// the PLE window of the token path: the two tokens before `pos_start`
@@ -12993,6 +13001,7 @@ int main(int argc, char **argv) try {
}
std::fprintf(stderr, "strata generate: window and draft graphs captured in %.0f ms\n",
std::chrono::duration<double, std::milli>(Clock::now() - tw).count());
strata::uc_probe("captured");
}
}
if (use_mtp && !first_window &&