|
43 | 43 | #include <cstring> |
44 | 44 | #include <ctime> |
45 | 45 | #include <functional> |
| 46 | +#include <vector> |
| 47 | +#include <algorithm> |
| 48 | +#include <random> |
46 | 49 | #include "timing_common.cuh" |
47 | 50 | #include <type_traits> |
48 | 51 |
|
@@ -316,6 +319,14 @@ static double g_row_spread = 0.0; |
316 | 319 | // trajectory in every trial of every contender. |
317 | 320 | static std::function<void()> g_pre_trial; |
318 | 321 |
|
| 322 | +// Opt-in randomized measurement order (GLASS_SHUFFLE_ORDER=<seed>, nonzero): |
| 323 | +// each row's (contender, launch-shape) cells execute in a shuffled order while |
| 324 | +// the printed row keeps the canonical BLOCK|WARP|THREAD|AUTO|NV layout, so |
| 325 | +// parsers and the default protocol are byte-identical. Guards against |
| 326 | +// systematic order effects (thermal/clock ramp) biasing later contenders. |
| 327 | +static uint64_t g_shuffle_seed = 0; |
| 328 | +static std::mt19937_64 g_shuffle_rng; |
| 329 | + |
319 | 330 | template<typename F> |
320 | 331 | static double time_ns_per_prob(F launch, int reps) { |
321 | 332 | double ns = tc_time_ns_per_prob_pre( |
@@ -423,41 +434,73 @@ static void bench_size(Op op, int reps) { |
423 | 434 | cudaDeviceSynchronize(); |
424 | 435 | }; |
425 | 436 |
|
| 437 | + g_row_spread = 0.0; |
| 438 | + // MEASURE phase: every (contender, launch-shape) cell is one item; with |
| 439 | + // GLASS_SHUFFLE_ORDER set the items execute in shuffled order. The PRINT |
| 440 | + // phase below always emits the canonical row layout from the stored |
| 441 | + // results, so output is byte-identical either way. |
| 442 | + constexpr int TBS[4] = {32, 64, 128, 256}; |
| 443 | + constexpr int WPBS[6] = {1, 2, 4, 8, 16, 32}; |
| 444 | + double r_block[4], r_warp[6], r_thread[4], r_auto[4], r_nv = -1.0; |
| 445 | + for (double* a : {r_block, r_auto, r_thread}) for (int i=0;i<4;i++) a[i] = 1e30; |
| 446 | + for (int i=0;i<6;i++) r_warp[i] = 1e30; |
| 447 | + std::vector<std::function<void()>> items; |
| 448 | + for (int i=0;i<4;i++) items.push_back([&,i]{ |
| 449 | + int TB = TBS[i]; |
| 450 | + r_block[i] = time_ns_per_prob([&]{ launch_block<T,N>(op, TB, A, B, C, x, y); }, reps); }); |
| 451 | + for (int i=0;i<6;i++) { |
| 452 | + if (WPBS[i] > NPROB) break; |
| 453 | + items.push_back([&,i]{ |
| 454 | + int WPB = WPBS[i]; |
| 455 | + r_warp[i] = time_ns_per_prob([&]{ launch_warp<T,N>(op, WPB, A, B, C, x, y); }, reps); }); |
| 456 | + } |
| 457 | + if constexpr (thread_ok<T,N>()) { |
| 458 | + for (int i=0;i<4;i++) items.push_back([&,i]{ |
| 459 | + int TPB = TBS[i]; |
| 460 | + r_thread[i] = time_ns_per_prob([&]{ launch_thread<T,N>(op, TPB, A, B, C, x, y); }, reps); }); |
| 461 | + } |
| 462 | + // AUTO: the bare shipped-default face at the block launch shapes. |
| 463 | + // AUDIT-ONLY (see header) — excluded from the winner verdict and ignored by |
| 464 | + // tune.py's parsers (its summary token trails the ones _ROW_RE captures). |
| 465 | + for (int i=0;i<4;i++) items.push_back([&,i]{ |
| 466 | + int TB = TBS[i]; |
| 467 | + r_auto[i] = time_ns_per_prob([&]{ launch_auto<T,N>(op, TB, A, B, C, x, y); }, reps); }); |
| 468 | + items.push_back([&]{ r_nv = nv_dispatch<T,N>(op, A, B, C, x, y, reps); }); |
| 469 | + if (g_shuffle_seed) std::shuffle(items.begin(), items.end(), g_shuffle_rng); |
| 470 | + for (auto& it : items) it(); |
| 471 | + |
| 472 | + // PRINT phase: canonical layout, identical to the historical in-order path. |
426 | 473 | double best_block=1e30, best_warp=1e30, best_thread=1e30; |
427 | 474 | int best_tb=0, best_wpb=0, best_tpb=0; |
428 | | - g_row_spread = 0.0; |
429 | 475 | printf("%-5s N=%-3d | BLOCK", op_name(op), N); |
430 | | - for (int TB : {32, 64, 128, 256}) { |
431 | | - double ns = time_ns_per_prob([&]{ launch_block<T,N>(op, TB, A, B, C, x, y); }, reps); |
432 | | - if (ns < 1e29) printf(" tb%d=%.4f", TB, ns); else printf(" tb%d=FAIL", TB); |
433 | | - if (ns < best_block) { best_block = ns; best_tb = TB; } |
| 476 | + for (int i=0;i<4;i++) { |
| 477 | + double ns = r_block[i]; |
| 478 | + if (ns < 1e29) printf(" tb%d=%.4f", TBS[i], ns); else printf(" tb%d=FAIL", TBS[i]); |
| 479 | + if (ns < best_block) { best_block = ns; best_tb = TBS[i]; } |
434 | 480 | } |
435 | 481 | printf(" | WARP"); |
436 | | - for (int WPB : {1, 2, 4, 8, 16, 32}) { |
437 | | - if (WPB > NPROB) break; |
438 | | - double ns = time_ns_per_prob([&]{ launch_warp<T,N>(op, WPB, A, B, C, x, y); }, reps); |
439 | | - if (ns < 1e29) printf(" w%d=%.4f", WPB, ns); else printf(" w%d=FAIL", WPB); |
440 | | - if (ns < best_warp) { best_warp = ns; best_wpb = WPB; } |
| 482 | + for (int i=0;i<6;i++) { |
| 483 | + if (WPBS[i] > NPROB) break; |
| 484 | + double ns = r_warp[i]; |
| 485 | + if (ns < 1e29) printf(" w%d=%.4f", WPBS[i], ns); else printf(" w%d=FAIL", WPBS[i]); |
| 486 | + if (ns < best_warp) { best_warp = ns; best_wpb = WPBS[i]; } |
441 | 487 | } |
442 | 488 | if constexpr (thread_ok<T,N>()) { |
443 | 489 | printf(" | THREAD"); |
444 | | - for (int TPB : {32, 64, 128, 256}) { |
445 | | - double ns = time_ns_per_prob([&]{ launch_thread<T,N>(op, TPB, A, B, C, x, y); }, reps); |
446 | | - if (ns < 1e29) printf(" t%d=%.4f", TPB, ns); else printf(" t%d=FAIL", TPB); |
447 | | - if (ns < best_thread) { best_thread = ns; best_tpb = TPB; } |
| 490 | + for (int i=0;i<4;i++) { |
| 491 | + double ns = r_thread[i]; |
| 492 | + if (ns < 1e29) printf(" t%d=%.4f", TBS[i], ns); else printf(" t%d=FAIL", TBS[i]); |
| 493 | + if (ns < best_thread) { best_thread = ns; best_tpb = TBS[i]; } |
448 | 494 | } |
449 | 495 | } |
450 | | - // AUTO: the bare shipped-default face at the block launch shapes. |
451 | | - // AUDIT-ONLY (see header) — excluded from the winner verdict and ignored by |
452 | | - // tune.py's parsers (its summary token trails the ones _ROW_RE captures). |
453 | 496 | double best_auto = 1e30; int best_atb = 0; |
454 | 497 | printf(" | AUTO"); |
455 | | - for (int TB : {32, 64, 128, 256}) { |
456 | | - double ns = time_ns_per_prob([&]{ launch_auto<T,N>(op, TB, A, B, C, x, y); }, reps); |
457 | | - if (ns < 1e29) printf(" a%d=%.4f", TB, ns); else printf(" a%d=FAIL", TB); |
458 | | - if (ns < best_auto) { best_auto = ns; best_atb = TB; } |
| 498 | + for (int i=0;i<4;i++) { |
| 499 | + double ns = r_auto[i]; |
| 500 | + if (ns < 1e29) printf(" a%d=%.4f", TBS[i], ns); else printf(" a%d=FAIL", TBS[i]); |
| 501 | + if (ns < best_auto) { best_auto = ns; best_atb = TBS[i]; } |
459 | 502 | } |
460 | | - double nv = nv_dispatch<T,N>(op, A, B, C, x, y, reps); |
| 503 | + double nv = r_nv; |
461 | 504 |
|
462 | 505 | // 4-way winner. thread/warp/block are all dependency-free pure SIMT, so they |
463 | 506 | // compete on raw time; only nvidia (MathDx) must clear tune_pick's margin — |
@@ -508,6 +551,12 @@ int main(int argc, char** argv) { |
508 | 551 | printf("# mega sweep | NPROB=%d reps=%d dtype=%s | ns/problem (lower=better) | optin_smem=%zuKB\n", NPROB, reps, f64 ? "f64" : "f32", g_optin_smem/1024); |
509 | 552 | printf("# contenders: BLOCK(SIMT, TB swept) | WARP(WPB swept) | THREAD(SIMT, TPB swept) | AUTO(bare glass::, shipped dispatch, TB swept; figure-only) | NV(cuBLASDx/cuSOLVERDx, forced; f32<=128, f64<=64)\n"); |
510 | 553 | printf("# THREAD stages operands global->registers->global in the per-problem-contiguous layout (uncoalesced; the layout tax is IN the timing).\n"); |
| 554 | + if (const char* s = getenv("GLASS_SHUFFLE_ORDER"); s && strtoull(s, nullptr, 10) != 0) { |
| 555 | + g_shuffle_seed = strtoull(s, nullptr, 10); |
| 556 | + g_shuffle_rng.seed(g_shuffle_seed); |
| 557 | + printf("# measurement_order=shuffled seed=%llu (per-cell execution order randomized; printed layout canonical)\n", |
| 558 | + (unsigned long long)g_shuffle_seed); |
| 559 | + } |
511 | 560 | tc_warm_gpu(); // steady boost clocks before the first timed cell |
512 | 561 | if (f64) run_all<double>(reps); |
513 | 562 | else run_all<float>(reps); |
|
0 commit comments