From c6aa2f12b0892f701a5c0528317d1c8202521f15 Mon Sep 17 00:00:00 2001 From: ZaneHam Date: Tue, 11 Aug 2026 00:13:55 +1200 Subject: [PATCH] tests: name tests by family and number, after z390 Family is the file stem, numbers are two digits, and the old descriptive name became a description the runner prints. th_check refuses to start if a name, number or family does not line up. --- CHANGELOG.md | 29 ++++ CONTRIBUTING.md | 40 +++++- Makefile | 4 +- src/fe/lexer.c | 9 ++ src/fe/lexer.h | 8 ++ tests/reprocheck.sh | 51 +++++++- tests/tabend.c | 88 ++++++------- tests/tbackend.c | 36 ++--- tests/tcbsync.c | 24 ++-- tests/tcfold.c | 24 ++-- tests/tcomp.c | 51 ++++---- tests/tdce.c | 28 ++-- tests/tenc.c | 36 ++--- tests/terrs.c | 28 ++-- tests/test_diag.opt | 7 + tests/test_errors.opt | 2 + tests/test_mfma.opt | 5 + tests/test_umulhi.opt | 2 + tests/tguard.c | 4 +- tests/tharns.h | 26 +++- tests/tinline.c | 16 +-- tests/tinsert.c | 8 +- tests/tmain.c | 298 ++++++++++++++++++++++++++++++++++-------- tests/tordr.c | 124 ++++++++++++++++++ tests/tphase.c | 32 ++--- tests/tra_ssa.c | 30 ++--- tests/tregalloc.c | 21 +-- tests/trpi.c | 53 ++++++++ tests/trv_buf.c | 28 ++-- tests/trv_elf.c | 132 +++++++++---------- tests/trv_enc.c | 100 +++++++------- tests/trv_isel.c | 260 ++++++++++++++++++------------------ tests/tsched.c | 16 +-- tests/tsmoke.c | 36 ++--- tests/tsoft_fp.c | 112 ++++++++-------- tests/tsroa.c | 8 +- tests/tstruct.c | 16 +-- tests/tsysprint.c | 36 ++--- tests/ttabs.c | 63 +++++++-- tests/ttdf.c | 156 +++++++++++----------- tests/ttmc.c | 16 +-- tests/ttriton.c | 108 +++++++-------- tests/ttypes.c | 20 +-- tests/twarpsize.c | 4 +- 44 files changed, 1385 insertions(+), 810 deletions(-) create mode 100644 tests/test_diag.opt create mode 100644 tests/test_errors.opt create mode 100644 tests/test_mfma.opt create mode 100644 tests/test_umulhi.opt create mode 100644 tests/tordr.c create mode 100644 tests/trpi.c diff --git a/CHANGELOG.md b/CHANGELOG.md index 0997dcb..9fe00ef 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -23,6 +23,35 @@ Booth — Changelog layout `bir.h` claims is checked rather than assumed (Zane Hambly, 2026-08-09) +- tests are named for their family and position, `rvi01` and `tdf39` rather + than `rv_isel_max_frame_slots_in_range`, after z390's `TESTDCB1`. The family + is the file stem, the old descriptive name became a description the runner + prints, and `fam_order` in `tests/tmain.c` is the one place a family is + declared (Zane Hambly, 2026-08-10) + +- the runner refuses to start if a test registers an unknown family, a name + that disagrees with its number, or a number already taken. 278 of the 380 + tests were registering under families `cat_order` did not list, so they ran + unheaded in link order, `--list` showed 102 of them and `--cat rv_enc` ran + none of them while exiting 0 (Zane Hambly, 2026-08-10) + +- `--fam` replaces `--cat`, which still works, and `--families` lists the + families with their files and counts (Zane Hambly, 2026-08-10) + +- the keyword table `lookup_keyword` binary searches is checked for ordering, + and every keyword is checked to still lex as a keyword. A misfiled entry lexed + as an identifier and surfaced as a parse error somewhere else entirely + (Zane Hambly, 2026-08-10) + +- `make repro` reads a sidecar `tests/NAME.opt` per fixture instead of counting + every refusal as a silent skip. Eight refusals were being hidden, one of them + a live `v_mfma` verifier failure, and an `xfail` that starts passing is now + reported too (Zane Hambly, 2026-08-10) + +- `tests/trpi.c` collects regressions for bugs that shipped, seeded with #160's + line-number corruption, which had four sites fixed and no test + (Zane Hambly, 2026-08-10) + ## 2026-08-07 Version 0.5.2. diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 2e82fe1..bcc08e6 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -159,14 +159,52 @@ Tenstorrent additionally sits above BIR through TDF (Tile DataFlow), a small IR # Build make -# Run the test suite (currently 274 tests across the frontends, +# Run the test suite (currently 390 tests across the frontends, # IR, backends, runtime, and SYSPRINT) make test +# List the families, or run just one of them +./trunner --families +./trunner --fam rvi +./trunner --list + # Run the emulator test suite (RDNA3, requires tinygrad mockgpu in WSL) python tests/emu/run_emu.py ``` +### Test naming + +Tests are named for their family and their position in it, `rvi01`, `tdf39`, +`smk04`, which is lifted straight off z390's `rt\test` directory where the +members run `TESTDCB1` through `TESTDCB9`. It is a bit cargo-culted from the +mainframe world and the eight character member limit it comes from stopped +mattering decades ago, but an alphabetical listing groups and orders itself, +grep agrees with the runner about what comes first, and nothing drifts. + +The family is the file stem and there is one family per file. `trv_isel.c` +carries `rvi`, `ttdf.c` carries `tdf`. Add a test with: + +```c +static void rvi66(void) { ...; PASS(); } +TH_REG("rvi", 66, "what it checks, in a phrase", rvi66) +``` + +The name says nothing on its own, so the description carries it and shows up in +`--list` and in the run. Keep it inside `TH_DESCW`. Every family has to appear +in `fam_order` in `tests/tmain.c`, which is also where the running order lives, +and `th_check` refuses to run if a name, number or family does not line up. +That check exists because 278 of the tests used to register under families the +runner had never heard of, where `--list` could not see them and `--cat` could +not reach them. + +Regression tests for bugs that actually shipped go in `tests/trpi.c` with the +issue number leading the description, z390's `RPI1540` habit with our numbering. + +`.cu` fixtures may carry a sidecar `tests/NAME.opt` declaring modes they are +known not to survive, one `xfail MODE reason` per line. `make repro` requires +every other mode to compile twice to identical bytes, and it reports an `xfail` +that starts passing as well, so a stale exception does not hide a fix. + Verify your changes don't introduce encoding regressions: ```bash llvm-objdump -d --mcpu=gfx1100 output.hsaco diff --git a/Makefile b/Makefile index 1ae4475..988adff 100644 --- a/Makefile +++ b/Makefile @@ -110,7 +110,9 @@ TSRC = tests/tmain.c tests/tsmoke.c tests/tcomp.c tests/tenc.c \ tests/tcbsync.c \ tests/tsoft_fp.c \ tests/tsysprint.c \ - tests/tbackend.c + tests/tbackend.c \ + tests/tordr.c \ + tests/trpi.c TOBJS = $(TSRC:%.c=$(OBJDIR)/%.o) COBJS = $(OBJDIR)/src/kauri_impl.o $(OBJDIR)/src/ir/bir.o $(OBJDIR)/src/ir/bir_print.o $(OBJDIR)/src/ir/bir_lower.o $(OBJDIR)/src/ir/bir_mem2reg.o $(OBJDIR)/src/ir/bir_cfold.o $(OBJDIR)/src/ir/bir_dce.o $(OBJDIR)/src/ir/bir_struct.o $(OBJDIR)/src/ir/bir_insert.o $(OBJDIR)/src/ir/bir_sroa.o $(OBJDIR)/src/ir/bir_inline.o \ diff --git a/src/fe/lexer.c b/src/fe/lexer.c index e95f1c8..c932845 100644 --- a/src/fe/lexer.c +++ b/src/fe/lexer.c @@ -115,6 +115,15 @@ static int lookup_keyword(const char *src, uint32_t len) return TOK_IDENT; } +/* The table stays static; the suite only needs to read the order. */ +int lexer_kw_count(void) { return NUM_KEYWORDS; } + +const char *lexer_kw_at(int i) +{ + if (i < 0 || i >= NUM_KEYWORDS) return NULL; + return keywords[i].name; +} + static const char *tok_names[] = { [TOK_INT_LIT] = "INT_LIT", [TOK_FLOAT_LIT] = "FLOAT_LIT", diff --git a/src/fe/lexer.h b/src/fe/lexer.h index d30c3f4..7e02437 100644 --- a/src/fe/lexer.h +++ b/src/fe/lexer.h @@ -26,4 +26,12 @@ int lexer_tokenize(lexer_t *L); int lexer_token_text(const lexer_t *L, const token_t *tok, char *buf, int bufsize); +/* lookup_keyword binary searches the keyword table, so keeping that table in + * order is VERY IMPORTANT. Put an entry in the wrong place and the keyword + * quietly lexes as an identifier, then blows up as a parse error somewhere + * else entirely. Handed out so the suite can check the order instead of + * trusting the joke above the table. */ +int lexer_kw_count(void); +const char *lexer_kw_at(int i); + #endif /* BARRACUDA_LEXER_H */ diff --git a/tests/reprocheck.sh b/tests/reprocheck.sh index e17c60c..b544aca 100644 --- a/tests/reprocheck.sh +++ b/tests/reprocheck.sh @@ -1,6 +1,13 @@ #!/bin/sh # Same input twice, same bytes out. Separate processes, so ASLR and any # pointer-ordered container shows up as a mismatch. +# +# Each fixture may carry a sidecar tests/NAME.opt listing the modes it is known +# not to survive, one "xfail MODE reason" per line. That is z390 keeping +# TESTDCB1.OPT next to TESTDCB1.MLC rather than parking a table of special +# cases off in the runner. This script used to read any refusal as a skip and +# say nothing about it, so eight of them sat quiet, and a backend that started +# refusing every fixture would still have come out green. set -u root=$(git rev-parse --show-toplevel) @@ -13,13 +20,47 @@ kath=./kath.exe modes="--amdgpu --nvidia-ptx --ir" fail=0 checked=0 -skipped=0 +declared=0 + +EMPTY=$(printf '' | sha256sum | cut -d' ' -f1) + +# Prints the declared reason when fixture $1 is xfail for mode $2, nothing +# otherwise. Always prints something on a match, so callers can test for it. +xfail_text() { + opt="${1%.cu}.opt" + [ -f "$opt" ] || return 0 + awk -v m="$2" ' + { sub(/#.*/, "") } + $1 == "xfail" && ($2 == "all" || $2 == m) { + $1 = ""; $2 = ""; sub(/^ +/, "") + print ($0 == "" ? "(no reason given)" : $0) + exit + }' "$opt" +} for f in tests/*.cu; do for m in $modes; do + why=$(xfail_text "$f" "$m") + + if [ -n "$why" ]; then + # Declared broken. Confirm it still is, so a fix does not slip past. + if "$kath" $m "$f" >/dev/null 2>&1; then + fail=$((fail + 1)) + echo "XPASS: $m $f succeeds now, drop its xfail from ${f%.cu}.opt" >&2 + else + declared=$((declared + 1)) + fi + continue + fi + a=$("$kath" $m "$f" 2>/dev/null | sha256sum 2>/dev/null | cut -d' ' -f1) - # A file the backend refuses is not a reproducibility failure. - if [ -z "$a" ]; then skipped=$((skipped + 1)); continue; fi + if [ -z "$a" ] || [ "$a" = "$EMPTY" ]; then + fail=$((fail + 1)) + echo "NO OUTPUT: $m $f emitted nothing and declares no xfail" >&2 + echo " add one to ${f%.cu}.opt if that is deliberate" >&2 + continue + fi + b=$("$kath" $m "$f" 2>/dev/null | sha256sum 2>/dev/null | cut -d' ' -f1) checked=$((checked + 1)) if [ "$a" != "$b" ]; then @@ -33,8 +74,8 @@ done if [ "$fail" -ne 0 ]; then echo >&2 - echo "reprocheck: $fail of $checked runs differed" >&2 + echo "reprocheck: $fail problem(s) alongside $checked reproducible runs" >&2 exit 1 fi -echo "reprocheck: $checked runs reproducible ($skipped skipped)" +echo "reprocheck: $checked runs reproducible, $declared declared xfail" diff --git a/tests/tabend.c b/tests/tabend.c index bc1d214..193c9e7 100644 --- a/tests/tabend.c +++ b/tests/tabend.c @@ -33,7 +33,7 @@ static int ab_snap(const ab_ctx_t *A, char *buf, int bsz) /* ---- ABEND Context Tests ---- */ -static void ab_init_def(void) +static void abd01(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -46,12 +46,12 @@ static void ab_init_def(void) free(A); PASS(); } -TH_REG("abend", ab_init_def) +TH_REG("abd", 1, "init def", abd01) /* Arms and confirms the flag; can't actually raise a fault in-process * without terminating the test. Real behaviour is exercised by the * --cpu launcher when a kernel goes sideways. */ -static void ab_arm_cpu_ok(void) +static void abd02(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -64,9 +64,9 @@ static void ab_arm_cpu_ok(void) free(A); PASS(); } -TH_REG("abend", ab_arm_cpu_ok) +TH_REG("abd", 2, "arm CPU ok", abd02) -static void ab_trak_one(void) +static void abd03(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -81,9 +81,9 @@ static void ab_trak_one(void) free(A); PASS(); } -TH_REG("abend", ab_trak_one) +TH_REG("abd", 3, "trak one", abd03) -static void ab_trak_full(void) +static void abd04(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -99,9 +99,9 @@ static void ab_trak_full(void) free(A); PASS(); } -TH_REG("abend", ab_trak_full) +TH_REG("abd", 4, "trak full", abd04) -static void ab_snag_snap(void) +static void abd05(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -132,9 +132,9 @@ static void ab_snag_snap(void) free(A); PASS(); } -TH_REG("abend", ab_snag_snap) +TH_REG("abd", 5, "snag snapshot", abd05) -static void ab_mstr_all(void) +static void abd06(void) { /* Every defined code should return a non-empty string */ uint16_t codes[] = { AB_G0C1, AB_G0C4, AB_G0C5, AB_G0C7, @@ -149,11 +149,11 @@ static void ab_mstr_all(void) CHECK(u != NULL); PASS(); } -TH_REG("abend", ab_mstr_all) +TH_REG("abd", 6, "mstr all", abd06) /* ---- Fault Correlation Tests ---- */ -static void ab_near_inside(void) +static void abd07(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -170,9 +170,9 @@ static void ab_near_inside(void) free(A); PASS(); } -TH_REG("abend", ab_near_inside) +TH_REG("abd", 7, "near inside", abd07) -static void ab_near_past(void) +static void abd08(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -190,9 +190,9 @@ static void ab_near_past(void) free(A); PASS(); } -TH_REG("abend", ab_near_past) +TH_REG("abd", 8, "near past", abd08) -static void ab_near_before(void) +static void abd09(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -209,9 +209,9 @@ static void ab_near_before(void) free(A); PASS(); } -TH_REG("abend", ab_near_before) +TH_REG("abd", 9, "near before", abd09) -static void ab_near_gap(void) +static void abd10(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -230,9 +230,9 @@ static void ab_near_gap(void) free(A); PASS(); } -TH_REG("abend", ab_near_gap) +TH_REG("abd", 10, "near gap", abd10) -static void ab_near_empty(void) +static void abd11(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -249,11 +249,11 @@ static void ab_near_empty(void) free(A); PASS(); } -TH_REG("abend", ab_near_empty) +TH_REG("abd", 11, "near empty", abd11) /* ---- Dump Formatting Tests ---- */ -static void ab_dump_hdr(void) +static void abd12(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -275,9 +275,9 @@ static void ab_dump_hdr(void) free(A); PASS(); } -TH_REG("abend", ab_dump_hdr) +TH_REG("abd", 12, "dump header", abd12) -static void ab_dump_mem(void) +static void abd13(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -298,9 +298,9 @@ static void ab_dump_mem(void) free(A); PASS(); } -TH_REG("abend", ab_dump_mem) +TH_REG("abd", 13, "dump memory", abd13) -static void ab_dump_disp(void) +static void abd14(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -325,9 +325,9 @@ static void ab_dump_disp(void) free(A); PASS(); } -TH_REG("abend", ab_dump_disp) +TH_REG("abd", 14, "dump displacement", abd14) -static void ab_dump_snap(void) +static void abd15(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -360,9 +360,9 @@ static void ab_dump_snap(void) free(A); PASS(); } -TH_REG("abend", ab_dump_snap) +TH_REG("abd", 15, "dump snapshot", abd15) -static void ab_dump_smap(void) +static void abd16(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -385,9 +385,9 @@ static void ab_dump_smap(void) free(A); PASS(); } -TH_REG("abend", ab_dump_smap) +TH_REG("abd", 16, "dump source map", abd16) -static void ab_dump_null(void) +static void abd17(void) { /* Null context should return error, not crash */ int rc = ab_dump(NULL, stderr); @@ -401,7 +401,7 @@ static void ab_dump_null(void) free(A); PASS(); } -TH_REG("abend", ab_dump_null) +TH_REG("abd", 17, "dump null", abd17) /* ---- Source Map Tests ---- */ @@ -485,7 +485,7 @@ static int build_test_elf(uint8_t *buf, int bsz, int n_entries) return (int)total; } -static void ab_slod_ok(void) +static void abd18(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -509,9 +509,9 @@ static void ab_slod_ok(void) free(A); PASS(); } -TH_REG("abend", ab_slod_ok) +TH_REG("abd", 18, "slod ok", abd18) -static void ab_slod_nosc(void) +static void abd19(void) { /* ELF with no .debug_bc section */ ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); @@ -530,9 +530,9 @@ static void ab_slod_nosc(void) free(A); PASS(); } -TH_REG("abend", ab_slod_nosc) +TH_REG("abd", 19, "slod no scratch", abd19) -static void ab_slod_bad(void) +static void abd20(void) { ab_ctx_t *A = calloc(1, sizeof(ab_ctx_t)); CHECK(A != NULL); @@ -554,11 +554,11 @@ static void ab_slod_bad(void) free(A); PASS(); } -TH_REG("abend", ab_slod_bad) +TH_REG("abd", 20, "slod bad", abd20) /* ---- Lang Integration Tests ---- */ -static void ab_lang_en(void) +static void abd21(void) { /* Compiled-in defaults should be present */ const char *s = ab_afmt(AB_G0C4); @@ -574,9 +574,9 @@ static void ab_lang_en(void) CHECK(s != NULL); PASS(); } -TH_REG("abend", ab_lang_en) +TH_REG("abd", 21, "lang en", abd21) -static void ab_lang_mi(void) +static void abd22(void) { /* Load Maori translations and check ABEND messages */ int rc = bc_eload("lang/mi.txt"); @@ -592,4 +592,4 @@ static void ab_lang_mi(void) bc_eload("lang/en.txt"); PASS(); } -TH_REG("abend", ab_lang_mi) +TH_REG("abd", 22, "lang mi", abd22) diff --git a/tests/tbackend.c b/tests/tbackend.c index e60ae18..b3ee4e1 100644 --- a/tests/tbackend.c +++ b/tests/tbackend.c @@ -5,7 +5,7 @@ #include "backend.h" #include -static void be_list_ok(void) +static void bkd01(void) { uint32_t n = 0; for (uint32_t i = 0; be_list[i] != NULL && i < 64; i++) { @@ -19,9 +19,9 @@ static void be_list_ok(void) CHECK(n >= 8); PASS(); } -TH_REG("backend", be_list_ok) +TH_REG("bkd", 1, "list ok", bkd01) -static void be_find_hits(void) +static void bkd02(void) { static const char * const known[] = { "amdgpu", "nvptx", "tensix", "tensix-rv32", @@ -34,18 +34,18 @@ static void be_find_hits(void) } PASS(); } -TH_REG("backend", be_find_hits) +TH_REG("bkd", 2, "find hits", bkd02) -static void be_find_miss(void) +static void bkd03(void) { CHECK(be_find(NULL) == NULL); CHECK(be_find("") == NULL); CHECK(be_find("this-is-not-a-real-backend") == NULL); PASS(); } -TH_REG("backend", be_find_miss) +TH_REG("bkd", 3, "find missing", bkd03) -static void be_feats_sane(void) +static void bkd04(void) { const uint32_t all = BE_F_SIMT | BE_F_SCALAR | BE_F_ATOMIC | BE_F_SHARED | @@ -57,7 +57,7 @@ static void be_feats_sane(void) } PASS(); } -TH_REG("backend", be_feats_sane) +TH_REG("bkd", 4, "feats sane", bkd04) /* ---- one target per run ---- * Every backend emits to cfg->output_file, so several at once used to @@ -65,7 +65,7 @@ TH_REG("backend", be_feats_sane) static char be_obuf[TH_BUFSZ]; -static void be_one_target(void) +static void bkd05(void) { int rc = th_run(BC_BIN " --amdgpu-bin --nvidia-ptx --metal " "tests/vector_add.cu -o be_multi.out", @@ -79,9 +79,9 @@ static void be_one_target(void) remove("be_multi.out"); PASS(); } -TH_REG("backend", be_one_target) +TH_REG("bkd", 5, "one target", bkd05) -static void be_one_target_ok(void) +static void bkd06(void) { int rc = th_run(BC_BIN " --nvidia-ptx tests/vector_add.cu -o be_single.ptx", be_obuf, TH_BUFSZ); @@ -89,14 +89,14 @@ static void be_one_target_ok(void) remove("be_single.ptx"); PASS(); } -TH_REG("backend", be_one_target_ok) +TH_REG("bkd", 6, "one target ok", bkd06) /* ---- no two backends claim the same flag ---- * Flags are declared in the descriptor rather than discovered, so a * collision is a static error we can catch here instead of a * first-past-the-post race decided by be_list order. */ -static void be_flags_unique(void) +static void bkd07(void) { for (uint32_t i = 0; be_list[i] != NULL && i < 64; i++) { if (be_list[i]->flags == NULL) continue; @@ -112,14 +112,14 @@ static void be_flags_unique(void) } PASS(); } -TH_REG("backend", be_flags_unique) +TH_REG("bkd", 7, "flags unique", bkd07) /* ---- a declared flag is a parsed flag ---- * The list and parse() could drift apart, which would leave a flag * routed to a backend that then ignores it. Offering each one back * through the registry proves the pairing still holds. */ -static void be_flags_parse(void) +static void bkd08(void) { for (uint32_t i = 0; be_list[i] != NULL && i < 64; i++) { if (be_list[i]->flags == NULL) continue; @@ -133,7 +133,7 @@ static void be_flags_parse(void) } PASS(); } -TH_REG("backend", be_flags_parse) +TH_REG("bkd", 8, "flags parse", bkd08) /* ---- every backend flag is documented ---- * Backends own their flags now, so --help can fall behind without @@ -141,7 +141,7 @@ TH_REG("backend", be_flags_parse) static char be_hbuf[TH_BUFSZ]; -static void be_flags_documented(void) +static void bkd09(void) { th_run(BC_BIN " --help", be_hbuf, TH_BUFSZ); for (uint32_t i = 0; be_list[i] != NULL && i < 64; i++) { @@ -155,4 +155,4 @@ static void be_flags_documented(void) } PASS(); } -TH_REG("backend", be_flags_documented) +TH_REG("bkd", 9, "flags documented", bkd09) diff --git a/tests/tcbsync.c b/tests/tcbsync.c index b4ad524..4ffc512 100644 --- a/tests/tcbsync.c +++ b/tests/tcbsync.c @@ -27,7 +27,7 @@ static int find_word(uint32_t w) /* ---- the signal programs the NIU for a full-width atomic increment ---- */ -static void cb_sem_inc_encodes_atomic(void) +static void cbs01(void) { rv_buf_init(&B); /* increment counter at remote 0x00009000, coords 0x00001234, by 1. */ @@ -54,11 +54,11 @@ static void cb_sem_inc_encodes_atomic(void) CHEQ(d[n - 1], rv_sw (RV_T1, RV_T0, 0x28)); PASS(); } -TH_REG("tensix", cb_sem_inc_encodes_atomic); +TH_REG("cbs", 1, "semaphore inc encodes atomic", cbs01); /* ---- the wait is a two-instruction spin that branches back on itself ---- */ -static void cb_wait_is_spin_loop(void) +static void cbs02(void) { rv_buf_init(&B); CHEQ(tt_sem_wait_ge(&B, 0x00009000u, 4u), BC_OK); @@ -72,11 +72,11 @@ static void cb_wait_is_spin_loop(void) CHEQ(d[n - 1], rv_bltu(RV_T1, RV_T2, -4)); PASS(); } -TH_REG("tensix", cb_wait_is_spin_loop); +TH_REG("cbs", 2, "wait is spin loop", cbs02); /* ---- the CB ops are the wait/signal pair under pipeline names ---- */ -static void cb_push_is_sem_inc(void) +static void cbs03(void) { static rv_buf_t A; rv_buf_init(&A); @@ -88,9 +88,9 @@ static void cb_push_is_sem_inc(void) CHECK(memcmp(rv_buf_data(&A), rv_buf_data(&B), n * 4u) == 0); PASS(); } -TH_REG("tensix", cb_push_is_sem_inc); +TH_REG("cbs", 3, "push is semaphore inc", cbs03); -static void cb_wait_front_is_acquire(void) +static void cbs04(void) { static rv_buf_t A; rv_buf_init(&A); @@ -103,11 +103,11 @@ static void cb_wait_front_is_acquire(void) CHECK(memcmp(rv_buf_data(&A), rv_buf_data(&B), n * 4u) == 0); PASS(); } -TH_REG("tensix", cb_wait_front_is_acquire); +TH_REG("cbs", 4, "wait front is acquire", cbs04); /* ---- acquire = wait_ge followed by an atomic decrement ---- */ -static void acquire_waits_then_decrements(void) +static void cbs05(void) { static rv_buf_t A; rv_buf_init(&A); @@ -120,11 +120,11 @@ static void acquire_waits_then_decrements(void) CHECK(memcmp(rv_buf_data(&A), rv_buf_data(&B), n * 4u) == 0); PASS(); } -TH_REG("tensix", acquire_waits_then_decrements); +TH_REG("cbs", 5, "acquire waits then decrements", cbs05); /* ---- compute weave brackets the issue stream with the CB handshake ---- */ -static void compute_weave_brackets_issue(void) +static void cbs06(void) { tt_module_t *m = (tt_module_t *)malloc(sizeof *m); tt_compute_sync_t s; @@ -175,4 +175,4 @@ static void compute_weave_brackets_issue(void) free(m); PASS(); } -TH_REG("tensix", compute_weave_brackets_issue); +TH_REG("cbs", 6, "compute weave brackets issue", cbs06); diff --git a/tests/tcfold.c b/tests/tcfold.c index 2904070..422c882 100644 --- a/tests/tcfold.c +++ b/tests/tcfold.c @@ -19,7 +19,7 @@ static const char *strnstr_range(const char *start, const char *end, /* ---- cf: integer arithmetic folded ---- */ -static void cf_int_arith(void) +static void cfd01(void) { int rc = th_run(BC_BIN " --ir tests/test_cf.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -35,11 +35,11 @@ static void cf_int_arith(void) CHECK(strnstr_range(body, fn_end, " 7") != NULL); PASS(); } -TH_REG("cf", cf_int_arith) +TH_REG("cfd", 1, "integer arithmetic folds", cfd01) /* ---- cf: chained constants fold ---- */ -static void cf_chain(void) +static void cfd02(void) { int rc = th_run(BC_BIN " --ir tests/test_cf.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -55,11 +55,11 @@ static void cf_chain(void) CHECK(strnstr_range(body, fn_end, " 20") != NULL); PASS(); } -TH_REG("cf", cf_chain) +TH_REG("cfd", 2, "folding follows a chain", cfd02) /* ---- cf: icmp + select with constant condition ---- */ -static void cf_icmp_select(void) +static void cfd03(void) { int rc = th_run(BC_BIN " --ir tests/test_cf.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -74,11 +74,11 @@ static void cf_icmp_select(void) CHECK(strnstr_range(body, fn_end, "= select") == NULL); PASS(); } -TH_REG("cf", cf_icmp_select) +TH_REG("cfd", 3, "icmp feeding select folds", cfd03) /* ---- cf: integer division by zero not folded ---- */ -static void cf_divzero(void) +static void cfd04(void) { int rc = th_run(BC_BIN " --ir tests/test_cf.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -92,11 +92,11 @@ static void cf_divzero(void) CHECK(strnstr_range(body, fn_end, "= sdiv") != NULL); PASS(); } -TH_REG("cf", cf_divzero) +TH_REG("cfd", 4, "sdiv by zero survives, folding it is UB", cfd04) /* ---- cf: integer/float conversions folded ---- */ -static void cf_conv(void) +static void cfd05(void) { int rc = th_run(BC_BIN " --ir tests/test_cf.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -111,11 +111,11 @@ static void cf_conv(void) CHECK(strnstr_range(body, fn_end, "= fptosi") == NULL); PASS(); } -TH_REG("cf", cf_conv) +TH_REG("cfd", 5, "conversions fold", cfd05) /* ---- cf: float arithmetic folded ---- */ -static void cf_float(void) +static void cfd06(void) { int rc = th_run(BC_BIN " --ir tests/test_cf.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -131,4 +131,4 @@ static void cf_float(void) CHECK(strnstr_range(body, fn_end, "f32 4") != NULL); PASS(); } -TH_REG("cf", cf_float) +TH_REG("cfd", 6, "float arithmetic folds", cfd06) diff --git a/tests/tcomp.c b/tests/tcomp.c index d7c47b3..79565f9 100644 --- a/tests/tcomp.c +++ b/tests/tcomp.c @@ -35,31 +35,34 @@ static int compile_cu(const char *cu, const char *extra) /* ---- compile: individual .cu files ---- */ -#define COMP_TEST(name, cu, extra) \ - static void name(void) { \ +/* fn and num are both spelled out rather than pasted together, because ## on a + * padded number hands TH_REG an octal constant and 09 does not even compile. + * th_check confirms the two agree at startup. */ +#define COMP_TEST(fn, num, desc, cu, extra) \ + static void fn(void) { \ CHECK(compile_cu(cu, extra) == 0); \ PASS(); \ } \ - TH_REG("compile", name) + TH_REG("cmp", num, desc, fn) -COMP_TEST(cmp_vecadd, "tests/vector_add.cu", "") -COMP_TEST(cmp_canon, "tests/canonical.cu", "") -COMP_TEST(cmp_feat, "tests/cuda_features.cu", "") -COMP_TEST(cmp_tier12, "tests/test_tier12.cu", "") -COMP_TEST(cmp_notgpt, "tests/notgpt.cu", "") -COMP_TEST(cmp_stress, "tests/stress.cu", "") -COMP_TEST(cmp_math, "tests/mymathhomework.cu", "") -COMP_TEST(cmp_launch, "tests/test_launch_bounds.cu", "") -COMP_TEST(cmp_coop, "tests/test_coop_groups.cu", "") -COMP_TEST(cmp_preproc, "tests/test_preproc.cu", "") -COMP_TEST(cmp_includ, "tests/test_include.cu", "-I tests") -COMP_TEST(cmp_tmpl, "tests/templates.cu", "") -COMP_TEST(cmp_unsign, "tests/test_unsigned.cu", "") -COMP_TEST(cmp_shr2d, "tests/test_shared2d.cu", "") -COMP_TEST(cmp_tgadd, "tests/tinygrad_add.cu", "") -COMP_TEST(cmp_tgmul, "tests/tinygrad_matmul.cu", "") -COMP_TEST(cmp_tgelm, "tests/tg_elem.cu", "") -COMP_TEST(cmp_tgvc4, "tests/tg_vec4.cu", "") -COMP_TEST(cmp_tgfil, "tests/tg_fill.cu", "") -COMP_TEST(cmp_tgred, "tests/tg_reduce.cu", "") -COMP_TEST(cmp_shfl, "tests/test_shfl.cu", "") +COMP_TEST(cmp01, 1, "vector add", "tests/vector_add.cu", "") +COMP_TEST(cmp02, 2, "canonical kernel", "tests/canonical.cu", "") +COMP_TEST(cmp03, 3, "CUDA feature sweep", "tests/cuda_features.cu", "") +COMP_TEST(cmp04, 4, "tier 1 and 2 ops", "tests/test_tier12.cu", "") +COMP_TEST(cmp05, 5, "LLM-written kernel", "tests/notgpt.cu", "") +COMP_TEST(cmp06, 6, "register stress", "tests/stress.cu", "") +COMP_TEST(cmp07, 7, "maths intrinsics", "tests/mymathhomework.cu", "") +COMP_TEST(cmp08, 8, "launch bounds", "tests/test_launch_bounds.cu", "") +COMP_TEST(cmp09, 9, "cooperative groups", "tests/test_coop_groups.cu", "") +COMP_TEST(cmp10, 10, "preprocessor", "tests/test_preproc.cu", "") +COMP_TEST(cmp11, 11, "include path", "tests/test_include.cu", "-I tests") +COMP_TEST(cmp12, 12, "templates", "tests/templates.cu", "") +COMP_TEST(cmp13, 13, "unsigned arithmetic", "tests/test_unsigned.cu", "") +COMP_TEST(cmp14, 14, "2D shared memory", "tests/test_shared2d.cu", "") +COMP_TEST(cmp15, 15, "tinygrad add", "tests/tinygrad_add.cu", "") +COMP_TEST(cmp16, 16, "tinygrad matmul", "tests/tinygrad_matmul.cu", "") +COMP_TEST(cmp17, 17, "tinygrad elementwise", "tests/tg_elem.cu", "") +COMP_TEST(cmp18, 18, "tinygrad vec4", "tests/tg_vec4.cu", "") +COMP_TEST(cmp19, 19, "tinygrad fill", "tests/tg_fill.cu", "") +COMP_TEST(cmp20, 20, "tinygrad reduce", "tests/tg_reduce.cu", "") +COMP_TEST(cmp21, 21, "warp shuffle", "tests/test_shfl.cu", "") diff --git a/tests/tdce.c b/tests/tdce.c index 67c9e30..b908d82 100644 --- a/tests/tdce.c +++ b/tests/tdce.c @@ -28,7 +28,7 @@ static int count_lines(const char *start, const char *end) /* ---- dce: dead chain eliminated ---- */ -static void dce_chain(void) +static void dce01(void) { int rc = th_run(BC_BIN " --ir tests/test_dce.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -45,11 +45,11 @@ static void dce_chain(void) CHECK(strnstr_range(fn, fn_end, "store ") != NULL); PASS(); } -TH_REG("dce", dce_chain) +TH_REG("dce", 1, "a dead chain goes", dce01) /* ---- dce: unused non-volatile load eliminated ---- */ -static void dce_load(void) +static void dce02(void) { int rc = th_run(BC_BIN " --ir tests/test_dce.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -65,11 +65,11 @@ static void dce_load(void) CHECK(strnstr_range(body, fn_end, "= gep") == NULL); PASS(); } -TH_REG("dce", dce_load) +TH_REG("dce", 2, "an unused non-volatile load goes", dce02) /* ---- dce: params survive even if unused ---- */ -static void dce_param(void) +static void dce03(void) { int rc1 = th_run(BC_BIN " --ir tests/test_dce.cu", obuf, TH_BUFSZ); CHEQ(rc1, 0); @@ -94,11 +94,11 @@ static void dce_param(void) CHECK(strstr(fn1, "i32 %3") != NULL); PASS(); } -TH_REG("dce", dce_param) +TH_REG("dce", 3, "params survive even when unused", dce03) /* ---- dce: side effects survive ---- */ -static void dce_side(void) +static void dce04(void) { int rc = th_run(BC_BIN " --ir tests/test_dce.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -110,11 +110,11 @@ static void dce_side(void) CHECK(strnstr_range(fn, fn_end, "barrier") != NULL); /* unique opcode */ PASS(); } -TH_REG("dce", dce_side) +TH_REG("dce", 4, "side effects survive", dce04) /* ---- dce: empty function unchanged ---- */ -static void dce_empty(void) +static void dce05(void) { int rc = th_run(BC_BIN " --ir tests/test_dce.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -125,11 +125,11 @@ static void dce_empty(void) CHECK(strnstr_range(fn, fn_end, "ret") != NULL); PASS(); } -TH_REG("dce", dce_empty) +TH_REG("dce", 5, "an empty function is unchanged", dce05) /* ---- dce: no dead code — output identical with and without DCE ---- */ -static void dce_nop(void) +static void dce06(void) { int rc1 = th_run(BC_BIN " --ir tests/test_dce.cu", obuf, TH_BUFSZ); CHEQ(rc1, 0); @@ -156,11 +156,11 @@ static void dce_nop(void) CHECK(strnstr_range(fn1, end1, "store ") != NULL); PASS(); } -TH_REG("dce", dce_nop) +TH_REG("dce", 6, "no dead code means identical output", dce06) /* ---- dce: instruction count drops ---- */ -static void dce_count(void) +static void dce07(void) { int rc1 = th_run(BC_BIN " --ir tests/test_dce.cu", obuf, TH_BUFSZ); CHEQ(rc1, 0); @@ -180,4 +180,4 @@ static void dce_count(void) CHECK(n_noopt > n_opt); PASS(); } -TH_REG("dce", dce_count) +TH_REG("dce", 7, "instruction count actually drops", dce07) diff --git a/tests/tenc.c b/tests/tenc.c index 0ce41d3..fb7ff11 100644 --- a/tests/tenc.c +++ b/tests/tenc.c @@ -69,7 +69,7 @@ static moperand_t imm(int32_t val) * [31:23]=0xBE8 prefix [22:16]=SDST(7) [15:8]=OP(0x00) [7:0]=SSRC0(2) * Expected: 0xBE870002 */ -static void enc_sop1(void) +static void enc01(void) { enc_setup(AMD_TARGET_GFX1100); minst_t *mi = &A->minsts[0]; @@ -83,14 +83,14 @@ static void enc_sop1(void) CHEQX(dw(0), 0xBE870002u); PASS(); } -TH_REG("encode", enc_sop1) +TH_REG("enc", 1, "SOP1 encodes s_mov_b32", enc01) /* ---- encode: SOP2 ---- */ /* s_add_u32 s7, s2, s3 -- GFX11 * [31:30]=10 [29:23]=OP(0x00) [22:16]=SDST(7) [15:8]=SSRC1(3) [7:0]=SSRC0(2) * Expected: 0x80070302 */ -static void enc_sop2(void) +static void enc02(void) { enc_setup(AMD_TARGET_GFX1100); minst_t *mi = &A->minsts[0]; @@ -105,14 +105,14 @@ static void enc_sop2(void) CHEQX(dw(0), 0x80070302u); PASS(); } -TH_REG("encode", enc_sop2) +TH_REG("enc", 2, "SOP2 encodes s_add_u32", enc02) /* ---- encode: VOP2 ---- */ /* v_add_f32 v1, v2, v3 -- GFX11 * [31]=0 [30:25]=OP(0x03) [24:17]=VDST(1) [16:9]=VSRC1(3) [8:0]=SRC0(v2=258) * Expected: 0x06020702 */ -static void enc_vop2(void) +static void enc03(void) { enc_setup(AMD_TARGET_GFX1100); minst_t *mi = &A->minsts[0]; @@ -127,14 +127,14 @@ static void enc_vop2(void) CHEQX(dw(0), 0x06020702u); PASS(); } -TH_REG("encode", enc_vop2) +TH_REG("enc", 3, "VOP2 encodes v_add_f32", enc03) /* ---- encode: VOP1 + literal ---- */ /* v_mov_b32 v5, 0xDEAD -- GFX11 * [31:25]=0x3F [24:17]=VDST(5) [16:9]=OP(0x01) [8:0]=SRC0(255=literal) * DW0: 0x7E0A02FF DW1: 0x0000DEAD */ -static void enc_vop1(void) +static void enc04(void) { enc_setup(AMD_TARGET_GFX1100); minst_t *mi = &A->minsts[0]; @@ -149,7 +149,7 @@ static void enc_vop1(void) CHEQX(dw(1), 0x0000DEADu); PASS(); } -TH_REG("encode", enc_vop1) +TH_REG("enc", 4, "VOP1 takes a literal operand", enc04) /* ---- encode: SMEM GFX11 vs GFX10 ---- */ /* s_load_dword s7, s[2:3], 0 @@ -157,7 +157,7 @@ TH_REG("encode", enc_vop1) * GFX11: DW1 SOFFSET null=0x7C -> 0xF8000000 * GFX10: DW1 SOFFSET null=0x7D -> 0xFA000000 */ -static void enc_smem(void) +static void enc05(void) { /* GFX11 */ enc_setup(AMD_TARGET_GFX1100); @@ -188,7 +188,7 @@ static void enc_smem(void) CHEQX(dw(1), 0xFA000000u); PASS(); } -TH_REG("encode", enc_smem) +TH_REG("enc", 5, "SMEM differs between GFX11 and GFX10", enc05) /* ---- encode: global_load_dword ---- */ /* global_load_dword v5, v2 -- GFX11 @@ -197,7 +197,7 @@ TH_REG("encode", enc_smem) * DW1: [31:24]=VDST(5) [23:16]=SADDR(0x7C=null) [15:8]=DATA(0) [7:0]=ADDR(2) * = 0x057C0002 */ -static void enc_glob(void) +static void enc06(void) { enc_setup(AMD_TARGET_GFX1100); minst_t *mi = &A->minsts[0]; @@ -212,7 +212,7 @@ static void enc_glob(void) CHEQX(dw(1), 0x057C0002u); PASS(); } -TH_REG("encode", enc_glob) +TH_REG("enc", 6, "global_load_dword encodes", enc06) /* ---- encode: DS ---- */ /* ds_load_b32 v5, v2 -- GFX11 @@ -221,7 +221,7 @@ TH_REG("encode", enc_glob) * DW1: [31:24]=VDST(5) [15:8]=DATA0(0) [7:0]=ADDR(2) * = 0x05000002 */ -static void enc_ds(void) +static void enc07(void) { enc_setup(AMD_TARGET_GFX1100); minst_t *mi = &A->minsts[0]; @@ -236,13 +236,13 @@ static void enc_ds(void) CHEQX(dw(1), 0x05000002u); PASS(); } -TH_REG("encode", enc_ds) +TH_REG("enc", 7, "DS encodes ds_read_b32", enc07) /* ---- encode: s_endpgm ---- */ /* GFX10: SOPP hw_op=0x01, SIMM16=0 * Expected: 0xBF810000 */ -static void enc_endp(void) +static void enc08(void) { enc_setup(AMD_TARGET_GFX1030); minst_t *mi = &A->minsts[0]; @@ -254,7 +254,7 @@ static void enc_endp(void) CHEQX(dw(0), 0xBF810000u); PASS(); } -TH_REG("encode", enc_endp) +TH_REG("enc", 8, "s_endpgm encodes", enc08) /* ---- encode: s_waitcnt vmcnt(0) ---- */ /* GFX10: hw_op=0x0C, vmcnt=0, lgkmcnt=63, expcnt=7 @@ -262,7 +262,7 @@ TH_REG("encode", enc_endp) * = 0x3F70 * DW = 0xBF8C3F70 */ -static void enc_wait(void) +static void enc09(void) { enc_setup(AMD_TARGET_GFX1030); minst_t *mi = &A->minsts[0]; @@ -275,4 +275,4 @@ static void enc_wait(void) CHEQX(dw(0), 0xBF8C3F70u); PASS(); } -TH_REG("encode", enc_wait) +TH_REG("enc", 9, "s_waitcnt vmcnt(0) encodes", enc09) diff --git a/tests/terrs.c b/tests/terrs.c index 11026af..c997567 100644 --- a/tests/terrs.c +++ b/tests/terrs.c @@ -7,7 +7,7 @@ static char obuf[TH_BUFSZ]; /* ---- errors: syntax ---- */ -static void err_synt(void) +static void err01(void) { int rc = th_run(BC_BIN " --amdgpu-bin tests/test_errors.cu -o err_test.hsaco", obuf, TH_BUFSZ); @@ -16,11 +16,11 @@ static void err_synt(void) remove("err_test.hsaco"); PASS(); } -TH_REG("errors", err_synt) +TH_REG("err", 1, "a syntax error is reported", err01) /* ---- errors: missing file ---- */ -static void err_miss(void) +static void err02(void) { int rc = th_run(BC_BIN " --amdgpu-bin nonexistent_file_42.cu -o err_test.hsaco", obuf, TH_BUFSZ); @@ -28,13 +28,13 @@ static void err_miss(void) remove("err_test.hsaco"); PASS(); } -TH_REG("errors", err_miss) +TH_REG("err", 2, "a missing file is reported", err02) /* ---- errors: bad output directory ---- */ /* /dev/null is a file, not a directory. You can't mkdir inside it. * Works on every Unix. On Windows we use NUL, same idea. */ -static void err_odir(void) +static void err03(void) { char cmd[TH_BUFSZ]; #ifdef _WIN32 @@ -52,13 +52,13 @@ static void err_odir(void) strstr(obuf, "error") != NULL); PASS(); } -TH_REG("errors", err_odir) +TH_REG("err", 3, "a bad output directory is reported", err03) /* ---- errors: diagnostic rendering + real token text ---- * A missing semicolon should render Clang-style (id in brackets, location * line, caret) and name the actual token it choked on, not the kind. */ -static void err_diag_render(void) +static void err04(void) { int rc = th_run(BC_BIN " --parse tests/test_diag.cu", obuf, TH_BUFSZ); (void)rc; @@ -69,13 +69,13 @@ static void err_diag_render(void) CHECK(strstr(obuf, "got 'IDENT'") == NULL); /* not the kind name */ PASS(); } -TH_REG("errors", err_diag_render) +TH_REG("err", 4, "diagnostic renders real token text", err04) /* ---- calls: past sixteen arguments ---- * Jorge Galvez's ocean kernels pass 23. Sema stopped counting at 16 and then * reported an arity mismatch against a count it had made up. */ -static void arg_many(void) +static void err05(void) { int rc = th_run(BC_BIN " --nvidia-ptx tests/many_args.cu -o many_args.ptx", obuf, TH_BUFSZ); @@ -84,13 +84,13 @@ static void arg_many(void) remove("many_args.ptx"); PASS(); } -TH_REG("errors", arg_many) +TH_REG("err", 5, "calls past sixteen arguments", err05) /* ---- calls: past the cap ---- * Overflowing has to say so. Silently dropping the tail would emit a call * with the wrong operands and nothing to show for it. */ -static void arg_cap(void) +static void err06(void) { FILE *f = fopen("argcap_test.cu", "w"); CHECK(f != NULL); @@ -108,14 +108,14 @@ static void arg_cap(void) remove("argcap_test.ptx"); PASS(); } -TH_REG("errors", arg_cap) +TH_REG("err", 6, "calls past the cap", err06) /* ---- errors: sema errors are fatal ---- * They used to be printed and then ignored by every mode but --sema, so the * backend ran on source we had already rejected and the exit status said the * compile went fine. */ -static void err_sema_fatal(void) +static void err07(void) { FILE *f = fopen("semafail_test.cu", "w"); CHECK(f != NULL); @@ -134,4 +134,4 @@ static void err_sema_fatal(void) remove("semafail_test.ptx"); PASS(); } -TH_REG("errors", err_sema_fatal) +TH_REG("err", 7, "sema errors are fatal", err07) diff --git a/tests/test_diag.opt b/tests/test_diag.opt new file mode 100644 index 0000000..adcf6b7 --- /dev/null +++ b/tests/test_diag.opt @@ -0,0 +1,7 @@ +# Options for test_diag.cu, read by reprocheck.sh. +# +# "xfail MODE reason" declares a mode this fixture is known not to survive. +# "xfail all reason" covers every mode. Anything not declared here has to +# compile, twice, to identical bytes. An xfail that starts passing is also +# reported, because a stale exception is how a fixed bug goes unnoticed. +xfail all a diagnostics fixture, it is meant to be rejected diff --git a/tests/test_errors.opt b/tests/test_errors.opt new file mode 100644 index 0000000..ae01a07 --- /dev/null +++ b/tests/test_errors.opt @@ -0,0 +1,2 @@ +# Options for test_errors.cu. See test_diag.opt for the format. +xfail all deliberately malformed source, rejection is the point diff --git a/tests/test_mfma.opt b/tests/test_mfma.opt new file mode 100644 index 0000000..46163bf --- /dev/null +++ b/tests/test_mfma.opt @@ -0,0 +1,5 @@ +# Options for test_mfma.cu. See test_diag.opt for the format. +# +# This one is a real bug rather than a missing feature. reprocheck used to +# count it as a skip and say nothing, which is how it stayed quiet. +xfail --amdgpu verify rejects VGPR in scalar source of v_mfma diff --git a/tests/test_umulhi.opt b/tests/test_umulhi.opt new file mode 100644 index 0000000..7380cd7 --- /dev/null +++ b/tests/test_umulhi.opt @@ -0,0 +1,2 @@ +# Options for test_umulhi.cu. See test_diag.opt for the format. +xfail --amdgpu __umul64hi is not implemented on the AMD backend diff --git a/tests/tguard.c b/tests/tguard.c index 8c831b4..45c4246 100644 --- a/tests/tguard.c +++ b/tests/tguard.c @@ -6,7 +6,7 @@ static char buf[1 << 16]; -static void guard_masks_not_ends(void) +static void grd01(void) { char cmd[TH_BUFSZ]; int rc; @@ -23,4 +23,4 @@ static void guard_masks_not_ends(void) PASS(); } -TH_REG("guard", guard_masks_not_ends) +TH_REG("grd", 1, "guarded lanes are masked, not ended", grd01) diff --git a/tests/tharns.h b/tests/tharns.h index 95b3b92..31fcb7e 100644 --- a/tests/tharns.h +++ b/tests/tharns.h @@ -9,27 +9,47 @@ typedef void (*tfunc_t)(void); +/* Names are family plus number, PDS member style, so an alphabetical listing + * groups and orders itself. I got this from living in z390 and mainframes for + * a while, so it is a bit cargo-culty, but the listing really does sort itself + * and nothing drifts. The name says nothing, so tdesc says it instead. */ typedef struct { const char *tname; - const char *tcats; + const char *tfam; + int tnum; + const char *tdesc; tfunc_t func; } tcase_t; #define TH_MAXTS 512 #define TH_BUFSZ 4096 +/* Column widths for the result line. TH_DESCW is a budget, not a hint. A + * description that overruns it is a description trying to become a comment. */ +#define TH_NAMEW 8 +#define TH_DESCW 46 + extern tcase_t th_list[]; extern int th_cnt; +extern int th_over; extern int npass, nfail, nskip; /* ---- Self-Registration ---- */ /* Works on gcc and clang, which is everyone who matters * and several who don't. MSVC users: you know what you did. */ -#define TH_REG(cat, fn) \ + +/* fam has to be in fam_order over in tmain.c and fn has to be spelled fam then + * num, padded. th_check enforces both, which is VERY IMPORTANT, because a test + * filed under a family nobody lists is a test that quietly stops running and + * still reports green. Overflow gets counted rather than swallowed for the + * same reason. */ +#define TH_REG(fam, num, desc, fn) \ __attribute__((constructor)) static void reg_##fn(void) { \ if (th_cnt < TH_MAXTS) \ - th_list[th_cnt++] = (tcase_t){#fn, cat, fn}; \ + th_list[th_cnt++] = (tcase_t){#fn, fam, num, desc, fn}; \ + else \ + th_over++; \ } /* ---- Assertions ---- */ diff --git a/tests/tinline.c b/tests/tinline.c index 03ded4c..00f1316 100644 --- a/tests/tinline.c +++ b/tests/tinline.c @@ -9,7 +9,7 @@ static char obuf[TH_BUFSZ]; /* On a GPU target every device call is inlined, so no call survives in the IR * (the standalone device bodies are inlined into their callers too). */ -static void inl_gpu_no_calls(void) +static void inl01(void) { int rc = th_run(BC_BIN " --amdgpu --ir tests/device_calls.cu", obuf, TH_BUFSZ); @@ -19,7 +19,7 @@ static void inl_gpu_no_calls(void) } /* And the kernel compiles the whole way to a .hsaco with the calls gone. */ -static void inl_gpu_binary(void) +static void inl02(void) { const char *out = "test_inline.hsaco"; int rc = th_run(BC_BIN " --amdgpu-bin tests/device_calls.cu " @@ -31,7 +31,7 @@ static void inl_gpu_binary(void) } /* NVIDIA and Tensix isel cannot emit a call either, so they inline too. */ -static void inl_other_gpu_no_calls(void) +static void inl03(void) { int rc = th_run(BC_BIN " --nvidia-ptx --ir tests/device_calls.cu", obuf, TH_BUFSZ); @@ -47,7 +47,7 @@ static void inl_other_gpu_no_calls(void) /* The CPU backend has a real SysV call ABI, so the inliner must not touch it: * the device calls stay as calls. */ -static void inl_cpu_keeps_calls(void) +static void inl04(void) { int rc = th_run(BC_BIN " --cpu --ir tests/device_calls.cu", obuf, TH_BUFSZ); @@ -56,7 +56,7 @@ static void inl_cpu_keeps_calls(void) PASS(); } -TH_REG("inline", inl_gpu_no_calls); -TH_REG("inline", inl_gpu_binary); -TH_REG("inline", inl_other_gpu_no_calls); -TH_REG("inline", inl_cpu_keeps_calls); +TH_REG("inl", 1, "a GPU kernel ends up with no calls", inl01); +TH_REG("inl", 2, "a GPU binary inlines", inl02); +TH_REG("inl", 3, "other GPU targets inline too", inl03); +TH_REG("inl", 4, "CPU keeps its calls", inl04); diff --git a/tests/tinsert.c b/tests/tinsert.c index 527cab7..9ff6f37 100644 --- a/tests/tinsert.c +++ b/tests/tinsert.c @@ -14,7 +14,7 @@ /* ---- insert: value refs slide, BR block ref stays ---- */ -static void insert_basic(void) +static void ins01(void) { bir_module_t *M = malloc(sizeof(*M)); uint32_t i32, vt, at; @@ -88,11 +88,11 @@ static void insert_basic(void) free(M); PASS(); } -TH_REG("insert", insert_basic) +TH_REG("ins", 1, "value refs slide, a BR block ref stays", ins01) /* ---- insert: PHI + BR_COND block halves survive, value halves slide ---- */ -static void insert_phi_blockrefs(void) +static void ins02(void) { bir_module_t *M = malloc(sizeof(*M)); uint32_t i32, vt, at; @@ -170,4 +170,4 @@ static void insert_phi_blockrefs(void) free(M); PASS(); } -TH_REG("insert", insert_phi_blockrefs) +TH_REG("ins", 2, "PHI block halves survive, values slide", ins02) diff --git a/tests/tmain.c b/tests/tmain.c index c6a9d3f..f155106 100644 --- a/tests/tmain.c +++ b/tests/tmain.c @@ -6,10 +6,11 @@ /* ---- Storage ---- */ tcase_t th_list[TH_MAXTS]; -int th_cnt = 0; -int npass = 0; -int nfail = 0; -int nskip = 0; +int th_cnt = 0; +int th_over = 0; +int npass = 0; +int nfail = 0; +int nskip = 0; /* ---- Utilities ---- */ @@ -48,28 +49,187 @@ int th_exist(const char *path) return 1; } -/* ---- Category Ordering ---- */ +/* ---- Families ---- */ -static const char *cat_order[] = { - "smoke", "compile", "encode", "tables", - "types", "errors", "phase", "sched", "abend", "regalloc", "ra_ssa", NULL +/* The one place a family is declared. This whole scheme is lifted from z390's + * rt\test directory, where the members are TESTDCB1 through TESTDCB9 and an + * alphabetical listing does the grouping for you. I got it from living in z390 + * and mainframes for a while so it is a bit cargo-culty, and the PDS eight + * character limit it comes from has not applied to anything here since about + * 1974. It still works better than what was here before. + * + * Order is pipeline order, not alphabetical, because you read a failing run top + * to bottom and want to see how far the compiler got before it fell over. + * + * Names, filters and listings are all checked against this table, so the file + * stem, the family and the test name cannot wander off from each other again. + * Two digits everywhere so grep and ls agree with the runner about what comes + * first. A family that outgrows 99 needs a wider width, and th_check will + * nag. */ +typedef struct { + const char *fam; + const char *file; + const char *what; + int width; +} tfam_t; + +static const tfam_t fam_order[] = { + { "smk", "tsmoke.c", "driver smoke", 2 }, + + { "pha", "tphase.c", "frontend phases", 2 }, + { "err", "terrs.c", "diagnostics", 2 }, + { "typ", "ttypes.c", "type table", 2 }, + { "tab", "ttabs.c", "static tables", 2 }, + + { "dce", "tdce.c", "dead code elimination", 2 }, + { "cfd", "tcfold.c", "constant folding", 2 }, + { "str", "tstruct.c", "structurisation", 2 }, + { "ins", "tinsert.c", "insertion points", 2 }, + { "sro", "tsroa.c", "scalar replacement", 2 }, + { "inl", "tinline.c", "inlining", 2 }, + + { "enc", "tenc.c", "AMD encoding", 2 }, + { "sch", "tsched.c", "AMD scheduling", 2 }, + { "ral", "tregalloc.c", "AMD register allocation", 2 }, + { "rss", "tra_ssa.c", "AMD SSA allocation", 2 }, + { "grd", "tguard.c", "AMD exec masking", 2 }, + + { "rve", "trv_enc.c", "RISC-V encoding", 2 }, + { "rvb", "trv_buf.c", "RISC-V code buffer", 2 }, + { "rvl", "trv_elf.c", "RISC-V ELF writer", 2 }, + { "rvi", "trv_isel.c", "RISC-V selection", 2 }, + { "cbs", "tcbsync.c", "Tensix CB sync", 2 }, + { "tmc", "ttmc.c", "Tensix machine code", 2 }, + { "tdf", "ttdf.c", "Tensix dataflow", 2 }, + + { "tri", "ttriton.c", "Triton frontend", 2 }, + + { "sfp", "tsoft_fp.c", "soft float", 2 }, + { "spr", "tsysprint.c", "SYSPRINT", 2 }, + { "abd", "tabend.c", "abend handling", 2 }, + + { "bkd", "tbackend.c", "backend registry", 2 }, + { "cmp", "tcomp.c", "compile matrix", 2 }, + { "wsz", "twarpsize.c", "warp size", 2 }, + + { "ord", "tordr.c", "harness ordering", 2 }, + { "rpi", "trpi.c", "shipped-bug regressions", 2 }, }; -static int cat_idx(const char *cat) +#define NFAM ((int)(sizeof(fam_order) / sizeof(fam_order[0]))) + +static int fam_idx(const char *fam) +{ + for (int i = 0; i < NFAM; i++) + if (strcmp(fam, fam_order[i].fam) == 0) return i; + return -1; +} + +/* ---- Startup Checks ---- */ + +/* Enforcing the naming is the whole trick, so this runs before anything else. + * It used to be possible to register under a family the runner had never heard + * of, and 278 of the 380 tests were doing exactly that, running unheaded at the + * end where --list could not see them and --cat could not reach them. Now a + * bad family, a misspelled name or a reused number stops the run and names the + * culprit. */ +static int th_check(void) +{ + int bad = 0; + + if (th_over) { + printf("harness: %d tests did not fit in TH_MAXTS (%d), raise it\n", + th_over, TH_MAXTS); + bad++; + } + + for (int i = 0; i < th_cnt; i++) { + const tcase_t *t = &th_list[i]; + int fi = fam_idx(t->tfam); + + if (fi < 0) { + printf("harness: %s registers family \"%s\", " + "which is not in fam_order\n", t->tname, t->tfam); + bad++; + continue; + } + + /* The function name has to spell out the family and number it + * registered with, or the three namespaces start drifting again. */ + char want[TH_NAMEW * 2]; + snprintf(want, sizeof want, "%s%0*d", + t->tfam, fam_order[fi].width, t->tnum); + if (strcmp(want, t->tname) != 0) { + printf("harness: %s registers as %s, name it %s\n", + t->tname, want, want); + bad++; + } + if ((int)strlen(want) > TH_NAMEW - 1) { + printf("harness: %s is wider than TH_NAMEW\n", want); + bad++; + } + if (t->tdesc == NULL || t->tdesc[0] == '\0') { + printf("harness: %s has no description\n", t->tname); + bad++; + } else if ((int)strlen(t->tdesc) > TH_DESCW) { + printf("harness: %s description is %d chars, budget is %d\n", + t->tname, (int)strlen(t->tdesc), TH_DESCW); + bad++; + } + + for (int j = i + 1; j < th_cnt; j++) { + if (th_list[j].tnum == t->tnum && + strcmp(th_list[j].tfam, t->tfam) == 0) { + printf("harness: %s and %s are both %s number %d\n", + t->tname, th_list[j].tname, t->tfam, t->tnum); + bad++; + } + } + } + + /* A family in the table with nothing registered means its file dropped out + * of TSRC, which reads as a smaller suite rather than a broken build. */ + for (int f = 0; f < NFAM; f++) { + int seen = 0; + for (int i = 0; i < th_cnt && !seen; i++) + if (strcmp(th_list[i].tfam, fam_order[f].fam) == 0) seen = 1; + if (!seen) { + printf("harness: family %s (%s) registered nothing, " + "is it still in TSRC?\n", + fam_order[f].fam, fam_order[f].file); + bad++; + } + } + + return bad; +} + +/* ---- Ordering ---- */ + +/* Family in pipeline order, then number. th_check has already proved the keys + * unique, so this does not care whether qsort is stable, and more to the point + * it does not care what order the constructors fired in. The old orphan pass + * did, which is why the tail of a run came out in link order. */ +static int case_cmp(const void *a, const void *b) { - for (int i = 0; cat_order[i]; i++) - if (strcmp(cat, cat_order[i]) == 0) return i; - return 99; + const tcase_t *x = a, *y = b; + int fx = fam_idx(x->tfam), fy = fam_idx(y->tfam); + if (fx != fy) return fx < fy ? -1 : 1; + if (x->tnum != y->tnum) return x->tnum < y->tnum ? -1 : 1; + return strcmp(x->tname, y->tname); } /* ---- Display ---- */ -static void print_result(const char *tname, const char *tag) +static void print_result(const tcase_t *tc, const char *tag) { - int nlen = (int)strlen(tname); - int dots = 30 - nlen; - if (dots < 3) dots = 3; - printf(" %s ", tname); + int dlen = (int)strlen(tc->tdesc); + int dots = TH_DESCW - dlen; + /* A description that fills the budget gets no dots at all, which is the + * point of ELDERBERRY over in tordr.c. th_check has already refused + * anything longer. */ + if (dots < 0) dots = 0; + printf(" %-*s %s ", TH_NAMEW, tc->tname, tc->tdesc); for (int i = 0; i < dots; i++) putchar('.'); printf(" %s\n", tag); } @@ -83,48 +243,85 @@ static void run_test(tcase_t *tc) tc->func(); if (nfail > was_fail) - print_result(tc->tname, "FAIL"); + print_result(tc, "FAIL"); else if (nskip > was_skip) - print_result(tc->tname, "SKIP"); + print_result(tc, "SKIP"); else if (npass > was_pass) - print_result(tc->tname, "PASS"); + print_result(tc, "PASS"); else { /* Test forgot to call PASS(). Benefit of the doubt, * like a lenient customs officer. */ npass++; - print_result(tc->tname, "PASS"); + print_result(tc, "PASS"); } } /* ---- Main ---- */ +static void usage(void) +{ + printf("usage: trunner [--all] [--fam FAM] [--test NAME] [--list] " + "[--families]\n"); +} + int main(int argc, char *argv[]) { - const char *filter_cat = NULL; + const char *filter_fam = NULL; const char *filter_test = NULL; int list_mode = 0; + int fam_mode = 0; for (int i = 1; i < argc; i++) { - if (strcmp(argv[i], "--cat") == 0 && i + 1 < argc) - filter_cat = argv[++i]; + /* --cat is what this used to be called. */ + if ((strcmp(argv[i], "--fam") == 0 || strcmp(argv[i], "--cat") == 0) + && i + 1 < argc) + filter_fam = argv[++i]; else if (strcmp(argv[i], "--test") == 0 && i + 1 < argc) filter_test = argv[++i]; else if (strcmp(argv[i], "--list") == 0) list_mode = 1; + else if (strcmp(argv[i], "--families") == 0) + fam_mode = 1; + else if (strcmp(argv[i], "--help") == 0) { + usage(); + return 0; + } /* --all is default, silently accepted */ } + if (th_check() != 0) { + printf("harness: refusing to run with the above unresolved\n"); + return 2; + } + + qsort(th_list, (size_t)th_cnt, sizeof th_list[0], case_cmp); + + /* A named family that does not exist used to run zero tests and exit 0, + * which reads exactly like a family that passed. */ + if (filter_fam && fam_idx(filter_fam) < 0) { + printf("trunner: no family \"%s\". --families lists them.\n", + filter_fam); + return 2; + } + + if (fam_mode) { + for (int f = 0; f < NFAM; f++) { + int n = 0; + for (int i = 0; i < th_cnt; i++) + if (strcmp(th_list[i].tfam, fam_order[f].fam) == 0) n++; + printf(" %-4s %-14s %3d %s\n", fam_order[f].fam, + fam_order[f].file, n, fam_order[f].what); + } + return 0; + } + if (list_mode) { - for (int ci = 0; cat_order[ci]; ci++) { - int printed_cat = 0; - for (int i = 0; i < th_cnt; i++) { - if (strcmp(th_list[i].tcats, cat_order[ci]) != 0) continue; - if (!printed_cat) { - printf("[%s]\n", cat_order[ci]); - printed_cat = 1; - } - printf(" %s\n", th_list[i].tname); - } + for (int i = 0; i < th_cnt; i++) { + const tcase_t *t = &th_list[i]; + if (filter_fam && strcmp(t->tfam, filter_fam) != 0) continue; + if (i == 0 || strcmp(t->tfam, th_list[i - 1].tfam) != 0) + printf("[%s] %s\n", t->tfam, fam_order[fam_idx(t->tfam)].what); + printf(" %-*s %s\n", TH_NAMEW, t->tname, t->tdesc); } return 0; } @@ -132,31 +329,17 @@ int main(int argc, char *argv[]) printf("Booth Test Suite\n"); printf("====================\n"); - /* Run tests in category order */ - for (int ci = 0; cat_order[ci]; ci++) { - if (filter_cat && strcmp(filter_cat, cat_order[ci]) != 0) - continue; - - int printed_cat = 0; - for (int i = 0; i < th_cnt; i++) { - if (strcmp(th_list[i].tcats, cat_order[ci]) != 0) continue; - if (filter_test && strcmp(filter_test, th_list[i].tname) != 0) - continue; + const char *last_fam = NULL; + for (int i = 0; i < th_cnt; i++) { + tcase_t *t = &th_list[i]; + if (filter_fam && strcmp(t->tfam, filter_fam) != 0) continue; + if (filter_test && strcmp(filter_test, t->tname) != 0) continue; - if (!printed_cat) { - printf("[%s]\n", cat_order[ci]); - printed_cat = 1; - } - run_test(&th_list[i]); + if (!last_fam || strcmp(t->tfam, last_fam) != 0) { + printf("[%s] %s\n", t->tfam, fam_order[fam_idx(t->tfam)].what); + last_fam = t->tfam; } - } - - /* Stray tests in unknown categories. Orphans, basically. */ - for (int i = 0; i < th_cnt; i++) { - if (cat_idx(th_list[i].tcats) < 99) continue; - if (filter_cat) continue; - if (filter_test && strcmp(filter_test, th_list[i].tname) != 0) continue; - run_test(&th_list[i]); + run_test(t); } int total = npass + nfail + nskip; @@ -164,6 +347,5 @@ int main(int argc, char *argv[]) printf("%d tests: %d passed, %d failed, %d skipped\n", total, npass, nfail, nskip); - (void)argc; return nfail > 0 ? 1 : 0; } diff --git a/tests/tordr.c b/tests/tordr.c new file mode 100644 index 0000000..6c03d9b --- /dev/null +++ b/tests/tordr.c @@ -0,0 +1,124 @@ +/* tordr.c -- the harness checks its own filing + * + * z390 tests a sort with data that narrates itself. APPLE says it is the first + * record, BANANA says it is the second, and ELDERBERRY is exactly ten + * characters so it fills the key field with no gap and walks the boundary on + * its way past. Get the sort wrong and you are not squinting at two columns of + * numbers, you are reading THIS IS THE THIRD RECORD sitting in second place. + * The data tells you off itself. + * + * Yes, this is cargo-culted from spending too long in z390 and mainframes. I + * make no apology for it. Every family now leans on the ordering in tmain.c so + * the ordering gets the same treatment. */ + +#include "tharns.h" + +/* Where a test sits among its own family, once tmain has sorted. */ +static int ord_place(const char *me) +{ + int place = 0; + for (int i = 0; i < th_cnt; i++) { + if (strcmp(th_list[i].tfam, "ord") != 0) continue; + place++; + if (strcmp(th_list[i].tname, me) == 0) return place; + } + return -1; +} + +/* claim is the position the description spells out in words. */ +static int ord_claims(const char *me, int claim, const char *word) +{ + int place = ord_place(me); + if (place < 0) { + printf(" %s is not in the listing at all\n", me); + return 0; + } + if (place != claim) { + printf(" %s SAYS IT IS THE %s RECORD, IT IS SITTING AT %d\n", + me, word, place); + return 0; + } + return 1; +} + +static void ord01(void) +{ + CHECK(ord_claims("ord01", 1, "FIRST")); + PASS(); +} +TH_REG("ord", 1, "THIS IS THE FIRST RECORD BY FAMILY ORDER", ord01) + +static void ord02(void) +{ + CHECK(ord_claims("ord02", 2, "SECOND")); + PASS(); +} +TH_REG("ord", 2, "THIS IS THE SECOND RECORD BY FAMILY ORDER", ord02) + +static void ord03(void) +{ + CHECK(ord_claims("ord03", 3, "THIRD")); + PASS(); +} +TH_REG("ord", 3, "THIS IS THE THIRD RECORD BY FAMILY ORDER", ord03) + +static void ord04(void) +{ + CHECK(ord_claims("ord04", 4, "FOURTH")); + PASS(); +} +TH_REG("ord", 4, "THIS IS THE FOURTH RECORD BY FAMILY ORDER", ord04) + +/* ELDERBERRY's job: fill the field exactly, so the boundary gets walked rather + * than assumed. This description is TH_DESCW characters, which means the dot + * run between it and the verdict is empty. Retyping it shorter is not a tidy, + * it is deleting the only test of the wide end of the column. */ +static void ord05(void) +{ + const char *me = NULL; + for (int i = 0; i < th_cnt; i++) + if (strcmp(th_list[i].tname, "ord05") == 0) me = th_list[i].tdesc; + + CHECK(me != NULL); + CHEQ((int)strlen(me), TH_DESCW); + CHECK(ord_claims("ord05", 5, "FIFTH")); + PASS(); +} +TH_REG("ord", 5, "THIS IS THE FIFTH RECORD AND IT FILLS THE LINE", ord05) + +/* The families themselves are ordered by the table in tmain.c rather than + * alphabetically, so smoke runs before the backends. Nothing narrates that, + * but the sort must at least be total: every test lands in exactly one place + * and no two tests claim it. */ +static void ord06(void) +{ + for (int i = 1; i < th_cnt; i++) { + if (strcmp(th_list[i].tfam, th_list[i - 1].tfam) != 0) continue; + if (th_list[i].tnum > th_list[i - 1].tnum) continue; + printf(" %s follows %s within %s\n", + th_list[i].tname, th_list[i - 1].tname, th_list[i].tfam); + nfail++; + return; + } + PASS(); +} +TH_REG("ord", 6, "numbers ascend within every family", ord06) + +/* A family's tests must be contiguous once sorted, or --fam and the headers + * disagree about what the family contains. */ +static void ord07(void) +{ + for (int i = 0; i < th_cnt; i++) { + for (int j = i + 2; j < th_cnt; j++) { + if (strcmp(th_list[i].tfam, th_list[j].tfam) != 0) continue; + if (strcmp(th_list[j - 1].tfam, th_list[i].tfam) == 0) continue; + printf(" family %s is split: %s then %s then %s\n", + th_list[i].tfam, th_list[i].tname, + th_list[j - 1].tname, th_list[j].tname); + nfail++; + return; + } + } + PASS(); +} +TH_REG("ord", 7, "each family occupies one run", ord07) diff --git a/tests/tphase.c b/tests/tphase.c index 9b8f2d8..80cb4c9 100644 --- a/tests/tphase.c +++ b/tests/tphase.c @@ -7,16 +7,16 @@ static char obuf[TH_BUFSZ]; /* ---- phase: preprocessor ---- */ -static void pha_pp(void) +static void pha01(void) { int rc = th_run(BC_BIN " --pp tests/vector_add.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); CHECK(strlen(obuf) > 0); PASS(); } -TH_REG("phase", pha_pp) +TH_REG("pha", 1, "preprocessor runs", pha01) -static void pha_pp_comment_opaque(void) +static void pha02(void) { int rc = th_run(BC_BIN " --pp tests/comment_macro_skip.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -24,30 +24,30 @@ static void pha_pp_comment_opaque(void) CHECK(strstr(obuf, "23 should remain literal") == NULL); PASS(); } -TH_REG("phase", pha_pp_comment_opaque) +TH_REG("pha", 2, "preprocessor comment opaque", pha02) -static void pha_lex_comment_quotes(void) +static void pha03(void) { int rc = th_run(BC_BIN " --lex tests/comment_quotes.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); CHECK(strstr(obuf, "0 error(s)") != NULL); PASS(); } -TH_REG("phase", pha_lex_comment_quotes) +TH_REG("pha", 3, "lexer comment quotes", pha03) /* ---- phase: parser (AST dump) ---- */ -static void pha_ast(void) +static void pha04(void) { int rc = th_run(BC_BIN " --parse tests/vector_add.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); PASS(); } -TH_REG("phase", pha_ast) +TH_REG("pha", 4, "parser dumps an AST", pha04) /* A function pointer declarator used to parse as a call expression, which cost nothing at parse time and everything later. */ -static void pha_fnptr(void) +static void pha05(void) { int rc = th_run(BC_BIN " --parse tests/fnptr.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -57,11 +57,11 @@ static void pha_fnptr(void) CHECK(strstr(obuf, "call") == NULL); PASS(); } -TH_REG("phase", pha_fnptr) +TH_REG("pha", 5, "function pointers survive the parser", pha05) /* Nested structs matter here: the enclosing name is saved and restored, so inner's constructor must not be mistaken for outer's. */ -static void pha_ctor(void) +static void pha06(void) { int rc = th_run(BC_BIN " --parse tests/ctor.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -71,25 +71,25 @@ static void pha_ctor(void) CHECK(strstr(obuf, "(ident sum)") != NULL); PASS(); } -TH_REG("phase", pha_ctor) +TH_REG("pha", 6, "constructors survive the parser", pha06) /* ---- phase: IR ---- */ -static void pha_ir(void) +static void pha07(void) { int rc = th_run(BC_BIN " --ir tests/vector_add.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); CHECK(strstr(obuf, "func") != NULL); PASS(); } -TH_REG("phase", pha_ir) +TH_REG("pha", 7, "IR comes out", pha07) /* ---- phase: semantic analysis ---- */ -static void pha_sema(void) +static void pha08(void) { int rc = th_run(BC_BIN " --sema tests/vector_add.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); PASS(); } -TH_REG("phase", pha_sema) +TH_REG("pha", 8, "semantic analysis runs", pha08) diff --git a/tests/tra_ssa.c b/tests/tra_ssa.c index 6a339c1..b7766f5 100644 --- a/tests/tra_ssa.c +++ b/tests/tra_ssa.c @@ -7,7 +7,7 @@ * declares. A kernel that reads past its declared count reads whatever the * previous wave left there. * - * Six fixtures still leak vregs (see ssa_ra_rejects_cleanly). They are pinned + * Six fixtures still leak vregs (see rss05). They are pinned * here rather than skipped, so fixing the allocator fails this file and makes * whoever fixes it move the fixture into the clean list. */ @@ -129,7 +129,7 @@ static const char *sr_leaky[] = { /* Every vreg must be gone by the time RA is done. verify.c says so too, but * asserting it here names the actual failure instead of an exit code. */ -static void ssa_ra_no_vreg_leak(void) +static void rss01(void) { int i; for (i = 0; sr_clean[i]; i++) { @@ -145,11 +145,11 @@ static void ssa_ra_no_vreg_leak(void) } PASS(); } -TH_REG("ra_ssa", ssa_ra_no_vreg_leak) +TH_REG("rss", 1, "no vreg leaks into the output", rss01) /* ---- Declared registers bound the ones actually used ---- */ -static void ssa_ra_within_declared(void) +static void rss02(void) { int i; for (i = 0; sr_clean[i]; i++) { @@ -161,10 +161,10 @@ static void ssa_ra_within_declared(void) } PASS(); } -TH_REG("ra_ssa", ssa_ra_within_declared) +TH_REG("rss", 2, "within declared", rss02) /* Same check with the budget squeezed enough to force spilling. */ -static void ssa_ra_within_declared_spilling(void) +static void rss03(void) { int i; for (i = 0; sr_clean[i]; i++) { @@ -181,12 +181,12 @@ static void ssa_ra_within_declared_spilling(void) } PASS(); } -TH_REG("ra_ssa", ssa_ra_within_declared_spilling) +TH_REG("rss", 3, "within declared spilling", rss03) /* Squeeze harder and the spill path leaks vregs too, on kernels the default * allocator handles down to --max-vgprs 2. Same deal as sr_leaky: pinned so a * fix trips this and gets folded into the test above. */ -static void ssa_ra_tight_cap_rejects_cleanly(void) +static void rss04(void) { const char *caps[] = { "--max-vgprs 4", "--max-vgprs 2", NULL }; int c; @@ -204,14 +204,14 @@ static void ssa_ra_tight_cap_rejects_cleanly(void) } PASS(); } -TH_REG("ra_ssa", ssa_ra_tight_cap_rejects_cleanly) +TH_REG("rss", 4, "tight cap rejects cleanly", rss04) /* ---- Known-broken fixtures fail loudly ---- */ /* These leak vregs. What matters until that is fixed is that verify catches * it and we exit non-zero, rather than quietly emitting a broken kernel. * Fix the allocator and this test fails: move the fixture to sr_clean. */ -static void ssa_ra_rejects_cleanly(void) +static void rss05(void) { int i; for (i = 0; sr_leaky[i]; i++) { @@ -227,13 +227,13 @@ static void ssa_ra_rejects_cleanly(void) } PASS(); } -TH_REG("ra_ssa", ssa_ra_rejects_cleanly) +TH_REG("rss", 5, "rejects cleanly", rss05) /* ---- The allocator doesn't take the compiler down ---- */ /* Bad output is one thing, a crash is another. Everything above runs through * here with the budget squeezed hard, and a clean rejection is fine. */ -static void ssa_ra_survives_pressure(void) +static void rss06(void) { const char *caps[] = { "", "--max-vgprs 4", "--max-vgprs 2", NULL }; int c, i; @@ -258,13 +258,13 @@ static void ssa_ra_survives_pressure(void) } PASS(); } -TH_REG("ra_ssa", ssa_ra_survives_pressure) +TH_REG("rss", 6, "survives pressure", rss06) /* ---- Both allocators agree on the shape of what they emit ---- */ /* Not a text diff: the whole point is that they allocate differently. But a * kernel is still a kernel, so the entry label and terminator must survive. */ -static void ssa_ra_emits_a_kernel(void) +static void rss07(void) { char cmd[TH_BUFSZ]; int i; @@ -280,4 +280,4 @@ static void ssa_ra_emits_a_kernel(void) } PASS(); } -TH_REG("ra_ssa", ssa_ra_emits_a_kernel) +TH_REG("rss", 7, "the allocator emits a kernel", rss07) diff --git a/tests/tregalloc.c b/tests/tregalloc.c index ccab5b6..869cd0a 100644 --- a/tests/tregalloc.c +++ b/tests/tregalloc.c @@ -107,21 +107,24 @@ static int cmp_ra(const char *cu, const char *extra) /* ---- regalloc: representative .cu files ---- */ -#define RA_TEST(name, cu, extra) \ - static void name(void) { \ +/* fn and num are both spelled out rather than pasted together, because ## on a + * padded number hands TH_REG an octal constant and 09 does not even compile. + * th_check confirms the two agree at startup. */ +#define RA_TEST(fn, num, desc, cu, extra) \ + static void fn(void) { \ CHECK(cmp_ra(cu, extra) == 0); \ PASS(); \ } \ - TH_REG("regalloc", name) + TH_REG("ral", num, desc, fn) -RA_TEST(ra_vecadd, "tests/vector_add.cu", "") -RA_TEST(ra_canon, "tests/canonical.cu", "") -RA_TEST(ra_notgpt, "tests/notgpt.cu", "") -RA_TEST(ra_stress, "tests/stress.cu", "") +RA_TEST(ral01, 1, "colouring beats scan: vector add", "tests/vector_add.cu", "") +RA_TEST(ral02, 2, "colouring beats scan: canonical", "tests/canonical.cu", "") +RA_TEST(ral03, 3, "colouring beats scan: notgpt", "tests/notgpt.cu", "") +RA_TEST(ral04, 4, "colouring beats scan: stress", "tests/stress.cu", "") /* ---- Spill path: force low VGPR cap to exercise spill code ---- */ -static void ra_spill(void) +static void ral05(void) { char cmd[TH_BUFSZ]; int rc; @@ -144,4 +147,4 @@ static void ra_spill(void) PASS(); } -TH_REG("regalloc", ra_spill) +TH_REG("ral", 5, "spills cleanly under a low VGPR cap", ral05) diff --git a/tests/trpi.c b/tests/trpi.c new file mode 100644 index 0000000..b399887 --- /dev/null +++ b/tests/trpi.c @@ -0,0 +1,53 @@ +/* trpi.c -- one test per bug that got away + * + * z390 files these as RPI1540, RPI2001A and so on, named for the problem report + * they came from, so years later you can still tell why a test exists at all. + * Ours are numbered in sequence like every other family and carry the issue in + * the description instead, which greps just as well and means there is only one + * naming rule to remember. + * + * The bar for landing here is that the bug shipped. If a fix has no test, the + * fix is a coincidence waiting to be undone. */ + +#include "tharns.h" + +static char obuf[1 << 16]; + +/* Reads the "; line N" annotation bir_print puts on the instruction matching + * needle, searching only within the named function. */ +static int line_of(const char *fn, const char *needle) +{ + const char *f = strstr(obuf, fn); + if (!f) return -1; + const char *end = strstr(f, "\n}"); + if (!end) return -1; + + const char *p = strstr(f, needle); + if (!p || p > end) return -1; + + const char *tag = strstr(p, "; line "); + if (!tag || tag > end) return -1; + return atoi(tag + 7); +} + +/* #160: DCE and mem2reg shuffled instructions down over the top of a deletion + * without moving inst_lines[] along with them, so every instruction past the + * first thing deleted reported whatever line its old neighbour had. Four sites + * were fixed and none of them got a test, which is what this is. + * + * dce_chain in test_dce.cu is the shape that catches it. Two dead instructions + * on lines 10 and 11 go away, and the store after them is on line 12. Get the + * line table wrong and the store starts claiming line 10 or 11. */ +static void rpi01(void) +{ + int rc = th_run(BC_BIN " --ir tests/test_dce.cu", obuf, (int)sizeof obuf); + CHEQ(rc, 0); + + /* int live = a + b; */ + CHEQ(line_of("@dce_chain", "= add "), 9); + /* out[0] = live; sits two deleted instructions later */ + CHEQ(line_of("@dce_chain", "store "), 12); + + PASS(); +} +TH_REG("rpi", 1, "#160 line numbers survive DCE", rpi01) diff --git a/tests/trv_buf.c b/tests/trv_buf.c index 8cb6f73..abb0ee9 100644 --- a/tests/trv_buf.c +++ b/tests/trv_buf.c @@ -8,7 +8,7 @@ static rv_buf_t B; /* ---- init zeros out the buffer ---- */ -static void rv_buf_init_zero(void) +static void rvb01(void) { memset(&B, 0xAB, sizeof(B)); rv_buf_init(&B); @@ -16,11 +16,11 @@ static void rv_buf_init_zero(void) CHEQ(rv_buf_nbytes(&B), 0u); PASS(); } -TH_REG("rv_enc", rv_buf_init_zero); +TH_REG("rvb", 1, "init zero", rvb01); /* ---- emit appends and returns the slot index ---- */ -static void rv_buf_emit_seq(void) +static void rvb02(void) { rv_buf_init(&B); int i0 = rv_buf_emit(&B, rv_addi(RV_X1, RV_X0, 1)); @@ -33,11 +33,11 @@ static void rv_buf_emit_seq(void) CHEQ(rv_buf_nbytes(&B), 12u); PASS(); } -TH_REG("rv_enc", rv_buf_emit_seq); +TH_REG("rvb", 2, "emit seq", rvb02); /* ---- patch overwrites a previously emitted word ---- */ -static void rv_buf_patch_basic(void) +static void rvb03(void) { rv_buf_init(&B); uint32_t orig = rv_addi(RV_X1, RV_X0, 0); @@ -49,42 +49,42 @@ static void rv_buf_patch_basic(void) CHEQ(rv_buf_data(&B)[0], want); PASS(); } -TH_REG("rv_enc", rv_buf_patch_basic); +TH_REG("rvb", 3, "patch basic", rvb03); /* ---- patch past end is rejected ---- */ -static void rv_buf_patch_oob(void) +static void rvb04(void) { rv_buf_init(&B); rv_buf_emit(&B, rv_nop()); CHEQ(rv_buf_patch(&B, 5u, rv_nop()), -1); PASS(); } -TH_REG("rv_enc", rv_buf_patch_oob); +TH_REG("rvb", 4, "patch out of bounds", rvb04); /* ---- offset arithmetic for branches ---- */ -static void rv_buf_offset_forward(void) +static void rvb05(void) { /* Branch at word 0 jumping to word 4 is +16 bytes. */ rv_buf_init(&B); CHEQ(rv_buf_offset(&B, 0u, 4u), 16); PASS(); } -TH_REG("rv_enc", rv_buf_offset_forward); +TH_REG("rvb", 5, "offset forward", rvb05); -static void rv_buf_offset_backward(void) +static void rvb06(void) { /* Branch at word 10 jumping back to word 3 is -28 bytes. */ rv_buf_init(&B); CHEQ(rv_buf_offset(&B, 10u, 3u), -28); PASS(); } -TH_REG("rv_enc", rv_buf_offset_backward); +TH_REG("rvb", 6, "offset backward", rvb06); /* ---- overflow returns -1 rather than corrupting memory ---- */ -static void rv_buf_overflow(void) +static void rvb07(void) { rv_buf_init(&B); for (uint32_t i = 0; i < RV_BUF_MAX_WORDS; i++) { @@ -95,4 +95,4 @@ static void rv_buf_overflow(void) CHEQ(rv_buf_n_words(&B), RV_BUF_MAX_WORDS); PASS(); } -TH_REG("rv_enc", rv_buf_overflow); +TH_REG("rvb", 7, "overflow", rvb07); diff --git a/tests/trv_elf.c b/tests/trv_elf.c index c1600aa..2c701a0 100644 --- a/tests/trv_elf.c +++ b/tests/trv_elf.c @@ -50,7 +50,7 @@ static void build_kernel(void) /* ---- magic bytes are present ---- */ -static void rv_elf_magic(void) +static void rvl01(void) { build_kernel(); CHEQ(rv_elf_write(&B, ELF_OUT), BC_OK); @@ -62,11 +62,11 @@ static void rv_elf_magic(void) CHEQ(rd[3], (uint8_t)'F'); PASS(); } -TH_REG("rv_enc", rv_elf_magic); +TH_REG("rvl", 1, "magic", rvl01); /* ---- e_ident says 32-bit LE, EV_CURRENT, no OSABI ---- */ -static void rv_elf_ident(void) +static void rvl02(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -77,11 +77,11 @@ static void rv_elf_ident(void) CHEQ(rd[7], 0u); /* ELFOSABI_NONE */ PASS(); } -TH_REG("rv_enc", rv_elf_ident); +TH_REG("rvl", 2, "ident", rvl02); /* ---- e_type, e_machine, e_version ---- */ -static void rv_elf_type_machine(void) +static void rvl03(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -91,11 +91,11 @@ static void rv_elf_type_machine(void) CHEQ(rd32(20), 1u); /* e_version */ PASS(); } -TH_REG("rv_enc", rv_elf_type_machine); +TH_REG("rvl", 3, "type machine", rvl03); /* ---- e_flags is zero (soft-float ABI, no compressed) ---- */ -static void rv_elf_flags_soft(void) +static void rvl04(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -103,11 +103,11 @@ static void rv_elf_flags_soft(void) CHEQ(rd32(36), 0u); PASS(); } -TH_REG("rv_enc", rv_elf_flags_soft); +TH_REG("rvl", 4, "flags soft", rvl04); /* ---- e_entry matches the documented load address ---- */ -static void rv_elf_entry_addr(void) +static void rvl05(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -115,11 +115,11 @@ static void rv_elf_entry_addr(void) CHEQ(rd32(24), RV_ELF_LOAD_ADDR); PASS(); } -TH_REG("rv_enc", rv_elf_entry_addr); +TH_REG("rvl", 5, "entry addr", rvl05); /* ---- PT_LOAD program header points at the code bytes ---- */ -static void rv_elf_pt_load(void) +static void rvl06(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -141,11 +141,11 @@ static void rv_elf_pt_load(void) CHEQ(rd32(e_phoff + 24), 5u); PASS(); } -TH_REG("rv_enc", rv_elf_pt_load); +TH_REG("rvl", 6, "pt load", rvl06); /* ---- code bytes appear at the offset the program header claims ---- */ -static void rv_elf_code_bytes(void) +static void rvl07(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -159,11 +159,11 @@ static void rv_elf_code_bytes(void) CHEQ(w2, rv_add (RV_A0, RV_A0, RV_A1)); PASS(); } -TH_REG("rv_enc", rv_elf_code_bytes); +TH_REG("rvl", 7, "code bytes", rvl07); /* ---- section header string table contains the section names ---- */ -static void rv_elf_shstrtab(void) +static void rvl08(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -181,17 +181,17 @@ static void rv_elf_shstrtab(void) CHECK(found_shstr); PASS(); } -TH_REG("rv_enc", rv_elf_shstrtab); +TH_REG("rvl", 8, "shstrtab", rvl08); /* ---- empty buffer refused ---- */ -static void rv_elf_empty(void) +static void rvl09(void) { rv_buf_init(&B); CHEQ(rv_elf_write(&B, ELF_OUT), BC_ERR_IO); PASS(); } -TH_REG("rv_enc", rv_elf_empty); +TH_REG("rvl", 9, "empty", rvl09); /* * The tests below encode tt-metal's loader contract from @@ -219,7 +219,7 @@ static uint32_t sfind(const char *n) } /* ReadImage:498 -- "first loadable segment is not text". */ -static void rv_elf_entry_is_first_seg(void) +static void rvl10(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -227,10 +227,10 @@ static void rv_elf_entry_is_first_seg(void) CHEQ(rd32(24), rd32(rd32(28) + 8)); PASS(); } -TH_REG("rv_enc", rv_elf_entry_is_first_seg); +TH_REG("rvl", 10, "entry is first segment", rvl10); /* ReadImage:472 -- p_offset, p_vaddr and p_paddr share 4-byte alignment. */ -static void rv_elf_phdr_aligned(void) +static void rvl11(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -239,10 +239,10 @@ static void rv_elf_phdr_aligned(void) CHEQ((rd32(ph + 4) | rd32(ph + 8) | rd32(ph + 12)) & 3u, 0u); PASS(); } -TH_REG("rv_enc", rv_elf_phdr_aligned); +TH_REG("rvl", 11, "phdr aligned", rvl11); /* ReadImage:435 -- sections and a valid nonzero shstrndx are mandatory. */ -static void rv_elf_has_sections(void) +static void rvl12(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -253,12 +253,12 @@ static void rv_elf_has_sections(void) CHECK(rd16(50) < rd16(48)); PASS(); } -TH_REG("rv_enc", rv_elf_has_sections); +TH_REG("rvl", 12, "has sections", rvl12); /* XIPify:1039 -- "there are no relocation sections". The section must be * SHT_RELA, and sh_info must name an alloc section that lies inside a * segment, else the loader skips it and the count stays zero. */ -static void rv_elf_rela_present(void) +static void rvl13(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -272,11 +272,11 @@ static void rv_elf_rela_present(void) CHEQ(rd32(shdr(r) + 24), sfind(".symtab")); /* sh_link */ PASS(); } -TH_REG("rv_enc", rv_elf_rela_present); +TH_REG("rvl", 13, "rela present", rvl13); /* TrimSegments:545 -- matched by name, SHT_PROGBITS and NOT SHF_ALLOC, * holding one (vma, trim_bound, size_limit) triple per segment. */ -static void rv_elf_segments_meta(void) +static void rvl14(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -292,11 +292,11 @@ static void rv_elf_segments_meta(void) CHEQ(rd32(off + 8), td_txtmax(td_chip())); PASS(); } -TH_REG("rv_enc", rv_elf_segments_meta); +TH_REG("rvl", 14, "segments meta", rvl14); /* XIPify:779 resolves relocation symbols through the symtab, so it must * exist and be non-alloc (WeakenDataSymbols skips alloc symtabs). */ -static void rv_elf_symtab(void) +static void rvl15(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -319,10 +319,10 @@ static void rv_elf_symtab(void) CHEQ(rd16(off + 16 + 14), sfind(".text")); PASS(); } -TH_REG("rv_enc", rv_elf_symtab); +TH_REG("rvl", 15, "symtab", rvl15); /* ReadImage:520 -- every SHF_ALLOC section must fall inside a segment. */ -static void rv_elf_alloc_in_seg(void) +static void rvl16(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -338,7 +338,7 @@ static void rv_elf_alloc_in_seg(void) } PASS(); } -TH_REG("rv_enc", rv_elf_alloc_in_seg); +TH_REG("rvl", 16, "alloc in segment", rvl16); /* * Structural invariants. The field-by-field tests above check that each @@ -348,7 +348,7 @@ TH_REG("rv_enc", rv_elf_alloc_in_seg); /* Every section body lies inside the file. SHT_NULL and SHT_NOBITS occupy * no file space and are exempt. */ -static void rv_elf_sections_in_file(void) +static void rvl17(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -362,10 +362,10 @@ static void rv_elf_sections_in_file(void) } PASS(); } -TH_REG("rv_enc", rv_elf_sections_in_file); +TH_REG("rvl", 17, "sections in file", rvl17); /* The section header table itself lies inside the file. */ -static void rv_elf_shdrs_in_file(void) +static void rvl18(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -373,11 +373,11 @@ static void rv_elf_shdrs_in_file(void) CHECK((long)(rd32(32) + rd16(48) * 40u) <= n); PASS(); } -TH_REG("rv_enc", rv_elf_shdrs_in_file); +TH_REG("rvl", 18, "section headers in file", rvl18); /* No two non-empty section bodies overlap on disk. A planner that * miscomputes one offset usually shows up here first. */ -static void rv_elf_no_overlap(void) +static void rvl19(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -394,11 +394,11 @@ static void rv_elf_no_overlap(void) } PASS(); } -TH_REG("rv_enc", rv_elf_no_overlap); +TH_REG("rvl", 19, "no overlap", rvl19); /* ReadImage:509 -- alloc, rela and symtab sections need sh_offset and * sh_addr 4-byte aligned. */ -static void rv_elf_section_align(void) +static void rvl20(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -412,10 +412,10 @@ static void rv_elf_section_align(void) } PASS(); } -TH_REG("rv_enc", rv_elf_section_align); +TH_REG("rvl", 20, "section align", rvl20); /* Every section name resolves inside .shstrtab and is NUL-terminated. */ -static void rv_elf_names_resolve(void) +static void rvl21(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -429,11 +429,11 @@ static void rv_elf_names_resolve(void) } PASS(); } -TH_REG("rv_enc", rv_elf_names_resolve); +TH_REG("rvl", 21, "names resolve", rvl21); /* .text's section header and its program header must describe the same * bytes, since the loader validates coverage across the two. */ -static void rv_elf_text_matches_phdr(void) +static void rvl22(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -444,13 +444,13 @@ static void rv_elf_text_matches_phdr(void) CHEQ(rd32(shdr(t) + 20), rd32(ph + 16)); /* sh_size == p_filesz */ PASS(); } -TH_REG("rv_enc", rv_elf_text_matches_phdr); +TH_REG("rvl", 22, "text matches phdr", rvl22); /* The load address is named in three places and the loader checks each * against a different thing, so they have to agree. TrimSegments runs * during ReadImage and matches the link-time address, before XIPify * rezeros it. */ -static void rv_elf_load_addr_agrees(void) +static void rvl23(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -460,11 +460,11 @@ static void rv_elf_load_addr_agrees(void) CHEQ(rd32(24), rd32(segs)); /* e_entry == segments vma */ PASS(); } -TH_REG("rv_enc", rv_elf_load_addr_agrees); +TH_REG("rvl", 23, "load addr agrees", rvl23); /* .symtab size must be a whole number of entries, and sh_info must name a * real one. */ -static void rv_elf_symtab_wellformed(void) +static void rvl24(void) { build_kernel(); rv_elf_write(&B, ELF_OUT); @@ -476,11 +476,11 @@ static void rv_elf_symtab_wellformed(void) CHECK(rd32(shdr(y) + 28) <= sz / es); PASS(); } -TH_REG("rv_enc", rv_elf_symtab_wellformed); +TH_REG("rvl", 24, "symtab well formed", rvl24); /* A one-instruction kernel still has to satisfy every alignment rule; the * smallest input is where padding bugs hide. */ -static void rv_elf_minimal_kernel(void) +static void rvl25(void) { rv_buf_init(&B); rv_buf_emit(&B, rv_nop()); @@ -492,10 +492,10 @@ static void rv_elf_minimal_kernel(void) CHECK((long)(rd32(32) + rd16(48) * 40u) <= n); PASS(); } -TH_REG("rv_enc", rv_elf_minimal_kernel); +TH_REG("rvl", 25, "minimal kernel", rvl25); /* Text size tracks the buffer rather than a fixed guess. */ -static void rv_elf_size_tracks_code(void) +static void rvl26(void) { rv_buf_init(&B); for (int i = 0; i < 64; i++) rv_buf_emit(&B, rv_nop()); @@ -505,7 +505,7 @@ static void rv_elf_size_tracks_code(void) CHEQ(rd32(shdr(sfind(".text")) + 20), 256u); PASS(); } -TH_REG("rv_enc", rv_elf_size_tracks_code); +TH_REG("rvl", 26, "size tracks code", rvl26); /* * End-to-end: drive the real compiler binary over a real .cu and validate @@ -545,7 +545,7 @@ static void text_span(uint32_t *off, uint32_t *nwords) /* A kernel with __device__ helpers compiles, and the result satisfies the * same loader contract the synthetic tests check. */ -static void e2e_rv_device_calls(void) +static void rvl27(void) { CHEQ(compile_cu("device_calls.cu"), 0); long n = slurp(E2E_OUT); @@ -558,12 +558,12 @@ static void e2e_rv_device_calls(void) CHEQ(rd32(24), rd32(rd32(28) + 8)); /* e_entry == p_vaddr */ PASS(); } -TH_REG("rv_enc", e2e_rv_device_calls); +TH_REG("rvl", 27, "device calls", rvl27); /* An unpatched call placeholder is a zero word, which is an illegal RV32 * instruction. main.c used to call rv_isel_func, which records call * patches but never resolves them. */ -static void e2e_rv_calls_patched(void) +static void rvl28(void) { CHEQ(compile_cu("device_calls.cu"), 0); slurp(E2E_OUT); @@ -579,7 +579,7 @@ static void e2e_rv_calls_patched(void) CHECK(jal > 0); PASS(); } -TH_REG("rv_enc", e2e_rv_calls_patched); +TH_REG("rvl", 28, "calls patched", rvl28); /* * The entry point must be the __global__ kernel. Source order puts the @@ -590,7 +590,7 @@ TH_REG("rv_enc", e2e_rv_calls_patched); * device_calls.cu's kernel calls helpers and every helper is a leaf, so * "the first function contains a JAL" distinguishes them. */ -static void e2e_rv_kernel_is_entry(void) +static void rvl29(void) { CHEQ(compile_cu("device_calls.cu"), 0); slurp(E2E_OUT); @@ -607,10 +607,10 @@ static void e2e_rv_kernel_is_entry(void) CHECK(jal > 0); /* and it is not a leaf */ PASS(); } -TH_REG("rv_enc", e2e_rv_kernel_is_entry); +TH_REG("rvl", 29, "kernel is entry", rvl29); /* Text is a whole number of 4-byte instructions and fits a baby core. */ -static void e2e_rv_text_sane(void) +static void rvl30(void) { CHEQ(compile_cu("device_calls.cu"), 0); slurp(E2E_OUT); @@ -620,11 +620,11 @@ static void e2e_rv_text_sane(void) CHECK(sz > 0u && sz <= td_txtmax(td_chip())); PASS(); } -TH_REG("rv_enc", e2e_rv_text_sane); +TH_REG("rvl", 30, "text sane", rvl30); /* Switch lowers to a compare chain, one BEQ per case, falling through to an * unconditional jump to the default block. */ -static void e2e_rv_switch_chain(void) +static void rvl31(void) { CHEQ(compile_cu("rv_switch.cu"), 0); slurp(E2E_OUT); @@ -641,11 +641,11 @@ static void e2e_rv_switch_chain(void) CHECK(jal > 0); PASS(); } -TH_REG("rv_enc", e2e_rv_switch_chain); +TH_REG("rvl", 31, "switch chain", rvl31); /* __shared__ resolves to an absolute address in the L1 slab at the top of * memory, materialised with a LUI whose upper immediate is the slab base. */ -static void e2e_rv_shared_addr(void) +static void rvl32(void) { CHEQ(compile_cu("rv_shared.cu"), 0); slurp(E2E_OUT); @@ -660,14 +660,14 @@ static void e2e_rv_shared_addr(void) CHECK(found); PASS(); } -TH_REG("rv_enc", e2e_rv_shared_addr); +TH_REG("rvl", 32, "shared addr", rvl32); /* The shared slab and the circular buffers must not overlap, which is only * true while the placer's ceiling is the slab base. */ -static void rv_l1_shared_below_end(void) +static void rvl33(void) { CHECK(td_shbase(td_chip()) + TD_L1_SHARED_SIZE == td_l1end(td_chip())); CHECK(TD_L1_CB_BASE < td_shbase(td_chip())); PASS(); } -TH_REG("rv_enc", rv_l1_shared_below_end); +TH_REG("rvl", 33, "L1 shared below end", rvl33); diff --git a/tests/trv_enc.c b/tests/trv_enc.c index 7a6543d..18f7c1c 100644 --- a/tests/trv_enc.c +++ b/tests/trv_enc.c @@ -11,217 +11,217 @@ * funct7=0, rs2=7, rs1=6, funct3=0, rd=5, opcode=0x33 packs to * (7<<20) | (6<<15) | (5<<7) | 0x33 = 0x007302B3. */ -static void rv_add_zero(void) +static void rve01(void) { CHEQ(rv_add(RV_X0, RV_X0, RV_X0), 0x00000033u); PASS(); } -TH_REG("rv_enc", rv_add_zero); +TH_REG("rve", 1, "add zero", rve01); -static void rv_add_x5_x6_x7(void) +static void rve02(void) { CHEQ(rv_add(RV_X5, RV_X6, RV_X7), 0x007302B3u); PASS(); } -TH_REG("rv_enc", rv_add_x5_x6_x7); +TH_REG("rve", 2, "add x5 x6 x7", rve02); /* sub differs from add only in funct7 bit 5 (funct7 = 0x20), so * sub x5, x6, x7 = 0x40000000 | 0x007302B3 = 0x407302B3. */ -static void rv_sub_x5_x6_x7(void) +static void rve03(void) { CHEQ(rv_sub(RV_X5, RV_X6, RV_X7), 0x407302B3u); PASS(); } -TH_REG("rv_enc", rv_sub_x5_x6_x7); +TH_REG("rve", 3, "sub x5 x6 x7", rve03); /* sra differs from srl only in funct7 bit 5. */ -static void rv_sra_vs_srl(void) +static void rve04(void) { uint32_t srl = rv_srl(RV_X1, RV_X2, RV_X3); uint32_t sra = rv_sra(RV_X1, RV_X2, RV_X3); CHEQ(sra ^ srl, 0x40000000u); /* differ only in bit 30 */ PASS(); } -TH_REG("rv_enc", rv_sra_vs_srl); +TH_REG("rve", 4, "SRA vs SRL", rve04); /* ---- I-type ALU ---- */ /* addi x1, x0, 5: imm=5 in bits [31:20], rd=1, opcode 0x13, giving * (5<<20) | (1<<7) | 0x13 = 0x00500093. */ -static void rv_addi_pos(void) +static void rve05(void) { CHEQ(rv_addi(RV_X1, RV_X0, 5), 0x00500093u); PASS(); } -TH_REG("rv_enc", rv_addi_pos); +TH_REG("rve", 5, "ADDI pos", rve05); /* Negative immediates sign-extend through the 12-bit field. For * addi x1, x0, -1 the imm masks to 0xFFF, so (0xFFF<<20) | (1<<7) | 0x13 * = 0xFFF00093. */ -static void rv_addi_neg(void) +static void rve06(void) { CHEQ(rv_addi(RV_X1, RV_X0, -1), 0xFFF00093u); PASS(); } -TH_REG("rv_enc", rv_addi_neg); +TH_REG("rve", 6, "ADDI neg", rve06); /* SLLI uses a 5-bit shamt in bits [24:20] with funct7=0x00, so * slli x1, x2, 4 = (4<<20) | (2<<15) | (1<<12) | (1<<7) | 0x13 = 0x00411093. */ -static void rv_slli_basic(void) +static void rve07(void) { CHEQ(rv_slli(RV_X1, RV_X2, 4), 0x00411093u); PASS(); } -TH_REG("rv_enc", rv_slli_basic); +TH_REG("rve", 7, "SLLI basic", rve07); /* SRAI funct7=0x20. */ -static void rv_srai_basic(void) +static void rve08(void) { uint32_t srli = rv_srli(RV_X1, RV_X2, 4); uint32_t srai = rv_srai(RV_X1, RV_X2, 4); CHEQ(srai ^ srli, 0x40000000u); PASS(); } -TH_REG("rv_enc", rv_srai_basic); +TH_REG("rve", 8, "SRAI basic", rve08); /* ---- Loads ---- */ /* lw x1, 0(x2): rs1=2, funct3=2 (LW), rd=1, opcode 0x03, giving * (2<<15) | (2<<12) | (1<<7) | 0x03 = 0x00012083. */ -static void rv_lw_basic(void) +static void rve09(void) { CHEQ(rv_lw(RV_X1, RV_X2, 0), 0x00012083u); PASS(); } -TH_REG("rv_enc", rv_lw_basic); +TH_REG("rve", 9, "LW basic", rve09); /* lw with a non-zero, sign-extended offset: for lw x1, -4(x2) the * imm masks to 0xFFC, giving 0xFFC00000 | 0x00012083 = 0xFFC12083. */ -static void rv_lw_neg_off(void) +static void rve10(void) { CHEQ(rv_lw(RV_X1, RV_X2, -4), 0xFFC12083u); PASS(); } -TH_REG("rv_enc", rv_lw_neg_off); +TH_REG("rve", 10, "LW neg offset", rve10); /* ---- Stores ---- */ /* sw x2, 0(x1): rs2=2, rs1=1, funct3=2, opcode 0x23, and no S-type * bit-scrambling matters since imm=0, giving (2<<20) | (1<<15) | (2<<12) * | 0x23 = 0x0020A023. */ -static void rv_sw_basic(void) +static void rve11(void) { CHEQ(rv_sw(RV_X2, RV_X1, 0), 0x0020A023u); PASS(); } -TH_REG("rv_enc", rv_sw_basic); +TH_REG("rve", 11, "SW basic", rve11); /* sw with a scrambled immediate verifies the S-type bit split. For * sw x2, 24(x1), 24 splits to imm[11:5]=0 and imm[4:0]=0x18, so * (2<<20) | (1<<15) | (2<<12) | (0x18<<7) | 0x23 = 0x0020AC23. */ -static void rv_sw_off24(void) +static void rve12(void) { CHEQ(rv_sw(RV_X2, RV_X1, 24), 0x0020AC23u); PASS(); } -TH_REG("rv_enc", rv_sw_off24); +TH_REG("rve", 12, "SW off24", rve12); /* ---- Branches ---- */ /* beq x0, x0, 0: opcode 0x63, funct3 0, all registers and immediate * bits zero, giving 0x00000063. */ -static void rv_beq_zero(void) +static void rve13(void) { CHEQ(rv_beq(RV_X0, RV_X0, 0), 0x00000063u); PASS(); } -TH_REG("rv_enc", rv_beq_zero); +TH_REG("rve", 13, "BEQ zero", rve13); /* beq x5, x6, 8: offset 8 puts imm[4:1]=0b0100 into bits [11:8], so * (6<<20) | (5<<15) | (4<<8) | 0x63 = 0x00628463. */ -static void rv_beq_off8(void) +static void rve14(void) { CHEQ(rv_beq(RV_X5, RV_X6, 8), 0x00628463u); PASS(); } -TH_REG("rv_enc", rv_beq_off8); +TH_REG("rve", 14, "BEQ off8", rve14); /* A negative branch tests the sign bit landing in inst[31]. Offset -4 * is 0x1FFC in the 13-bit field, scattering to 0x80000000 | 0x7E000000 * | 0xE00 | 0x80 | 0x63 = 0xFE000EE3. */ -static void rv_beq_back(void) +static void rve15(void) { CHEQ(rv_beq(RV_X0, RV_X0, -4), 0xFE000EE3u); PASS(); } -TH_REG("rv_enc", rv_beq_back); +TH_REG("rve", 15, "BEQ back", rve15); /* ---- Jumps and uppers ---- */ /* jal x1, 0: opcode 0x6F, rd=1, offset 0, giving (1<<7) | 0x6F = 0x000000EF. */ -static void rv_jal_zero(void) +static void rve16(void) { CHEQ(rv_jal(RV_X1, 0), 0x000000EFu); PASS(); } -TH_REG("rv_enc", rv_jal_zero); +TH_REG("rve", 16, "JAL zero", rve16); /* lui x1, 0x12345: imm[31:12]=0x12345, rd=1, opcode 0x37, giving * 0x12345000 | 0xB7 = 0x123450B7. */ -static void rv_lui_basic(void) +static void rve17(void) { CHEQ(rv_lui(RV_X1, 0x12345u), 0x123450B7u); PASS(); } -TH_REG("rv_enc", rv_lui_basic); +TH_REG("rve", 17, "LUI basic", rve17); /* auipc differs from lui only in opcode (0x17 vs 0x37). */ -static void rv_auipc_vs_lui(void) +static void rve18(void) { uint32_t lui = rv_lui(RV_X1, 0x12345u); uint32_t auipc = rv_auipc(RV_X1, 0x12345u); CHEQ(lui ^ auipc, 0x20u); /* opcode bits 6:0 differ by 0x20 */ PASS(); } -TH_REG("rv_enc", rv_auipc_vs_lui); +TH_REG("rve", 18, "AUIPC vs LUI", rve18); /* ---- M extension ---- */ /* mul x5, x6, x7: funct7=0x01, rs2=7, rs1=6, rd=5, opcode 0x33, giving * 0x02000000 | (7<<20) | (6<<15) | (5<<7) | 0x33 = 0x027302B3. */ -static void rv_mul_basic(void) +static void rve19(void) { CHEQ(rv_mul(RV_X5, RV_X6, RV_X7), 0x027302B3u); PASS(); } -TH_REG("rv_enc", rv_mul_basic); +TH_REG("rve", 19, "MUL basic", rve19); /* div and divu share the funct7=0x01 prefix with mul; differ only * in funct3. div: funct3=4, mul: funct3=0. */ -static void rv_div_vs_mul(void) +static void rve20(void) { uint32_t mul = rv_mul(RV_X5, RV_X6, RV_X7); uint32_t div = rv_div(RV_X5, RV_X6, RV_X7); CHEQ(div ^ mul, 0x4000u); /* funct3 bit 2 differs */ PASS(); } -TH_REG("rv_enc", rv_div_vs_mul); +TH_REG("rve", 20, "DIV vs MUL", rve20); /* All 8 M-extension instructions must have funct7 bit 25 set. */ -static void rv_m_ext_funct7(void) +static void rve21(void) { CHEQ((rv_mul (RV_X1, RV_X2, RV_X3) >> 25) & 0x7Fu, 0x01u); CHEQ((rv_mulh (RV_X1, RV_X2, RV_X3) >> 25) & 0x7Fu, 0x01u); @@ -233,39 +233,39 @@ static void rv_m_ext_funct7(void) CHEQ((rv_remu (RV_X1, RV_X2, RV_X3) >> 25) & 0x7Fu, 0x01u); PASS(); } -TH_REG("rv_enc", rv_m_ext_funct7); +TH_REG("rve", 21, "m ext funct7", rve21); /* ---- System ---- */ -static void rv_ecall_const(void) +static void rve22(void) { CHEQ(rv_ecall(), 0x00000073u); PASS(); } -TH_REG("rv_enc", rv_ecall_const); +TH_REG("rve", 22, "ECALL const", rve22); -static void rv_ebreak_const(void) +static void rve23(void) { CHEQ(rv_ebreak(), 0x00100073u); PASS(); } -TH_REG("rv_enc", rv_ebreak_const); +TH_REG("rve", 23, "EBREAK const", rve23); /* nop is addi x0, x0, 0 = 0x00000013 */ -static void rv_nop_canonical(void) +static void rve24(void) { CHEQ(rv_nop(), 0x00000013u); PASS(); } -TH_REG("rv_enc", rv_nop_canonical); +TH_REG("rve", 24, "NOP canonical", rve24); /* fence rw, rw with pred=0xF and succ=0xF gives (0xF<<24) | (0xF<<20) * | 0x0F = 0x0FF0000F. */ -static void rv_fence_rwrw(void) +static void rve25(void) { CHEQ(rv_fence(0xF, 0xF), 0x0FF0000Fu); PASS(); } -TH_REG("rv_enc", rv_fence_rwrw); +TH_REG("rve", 25, "FENCE rw,rw", rve25); diff --git a/tests/trv_isel.c b/tests/trv_isel.c index a69f70e..9f68fb9 100644 --- a/tests/trv_isel.c +++ b/tests/trv_isel.c @@ -94,7 +94,7 @@ static void build_bir_add(void) /* ---- isel produces some code without erroring ---- */ -static void rv_isel_smoke(void) +static void rvi01(void) { build_bir_add(); rv_buf_init(&code); @@ -102,11 +102,11 @@ static void rv_isel_smoke(void) CHECK(rv_buf_n_words(&code) > 0u); PASS(); } -TH_REG("rv_enc", rv_isel_smoke); +TH_REG("rvi", 1, "smoke", rvi01); /* ---- prologue is the documented one: drop sp, save ra ---- */ -static void rv_isel_prologue(void) +static void rvi02(void) { build_bir_add(); rv_buf_init(&code); @@ -122,11 +122,11 @@ static void rv_isel_prologue(void) CHEQ(w1, rv_sw(RV_RA, RV_SP, 0)); PASS(); } -TH_REG("rv_enc", rv_isel_prologue); +TH_REG("rvi", 2, "prologue", rvi02); /* ---- ADD instruction shows up in the body ---- */ -static void rv_isel_has_add(void) +static void rvi03(void) { build_bir_add(); rv_buf_init(&code); @@ -140,11 +140,11 @@ static void rv_isel_has_add(void) CHECK(found); PASS(); } -TH_REG("rv_enc", rv_isel_has_add); +TH_REG("rvi", 3, "has add", rvi03); /* ---- final SW writes through the pointer parameter ---- */ -static void rv_isel_has_store(void) +static void rvi04(void) { build_bir_add(); rv_buf_init(&code); @@ -158,11 +158,11 @@ static void rv_isel_has_store(void) CHECK(found); PASS(); } -TH_REG("rv_enc", rv_isel_has_store); +TH_REG("rvi", 4, "has store", rvi04); /* ---- epilogue restores ra, jalr to return ---- */ -static void rv_isel_epilogue(void) +static void rvi05(void) { build_bir_add(); rv_buf_init(&code); @@ -180,11 +180,11 @@ static void rv_isel_epilogue(void) CHEQ(rv_buf_data(&code)[n - 3], rv_lw(RV_RA, RV_SP, 0)); PASS(); } -TH_REG("rv_enc", rv_isel_epilogue); +TH_REG("rvi", 5, "epilogue", rvi05); /* ---- end-to-end: isel + ELF write produces a valid file ---- */ -static void rv_isel_to_elf(void) +static void rvi06(void) { build_bir_add(); rv_buf_init(&code); @@ -203,11 +203,11 @@ static void rv_isel_to_elf(void) CHEQ(hdr[3], (uint8_t)'F'); PASS(); } -TH_REG("rv_enc", rv_isel_to_elf); +TH_REG("rvi", 6, "to ELF", rvi06); /* ---- unsupported op produces a clean refusal ---- */ -static void rv_isel_unsupported(void) +static void rvi07(void) { /* Replace inst 4 with a SHARED_ALLOC, which still has no * lowering (the L1 region management for __shared__ memory @@ -220,7 +220,7 @@ static void rv_isel_unsupported(void) CHEQ(rv_isel_func(&fake_bir, 0, &code), BC_ERR_TDF); PASS(); } -TH_REG("rv_enc", rv_isel_unsupported); +TH_REG("rvi", 7, "unsupported", rvi07); /* ---- THREAD_ID etc. lower to LW from L1 runtime args ---- */ /* Each coordinate intrinsic emits two instructions: a LUI to materialise @@ -229,7 +229,7 @@ TH_REG("rv_enc", rv_isel_unsupported); * surrounding code shifts as other lowerings change. No stub here: the * kernel really reads from L1, and what's in that slot is the launcher's * job, not the compiler's. */ -static void rv_isel_thread_id_x_via_l1(void) +static void rvi08(void) { build_bir_add(); fake_bir.insts[3].op = BIR_THREAD_ID; @@ -243,9 +243,9 @@ static void rv_isel_thread_id_x_via_l1(void) (int16_t)RT_ARG_OFF_TID_X))); PASS(); } -TH_REG("rv_enc", rv_isel_thread_id_x_via_l1); +TH_REG("rvi", 8, "thread id x via L1", rvi08); -static void rv_isel_block_id_z_via_l1(void) +static void rvi09(void) { build_bir_add(); fake_bir.insts[3].op = BIR_BLOCK_ID; @@ -258,9 +258,9 @@ static void rv_isel_block_id_z_via_l1(void) (int16_t)RT_ARG_OFF_BID_Z))); PASS(); } -TH_REG("rv_enc", rv_isel_block_id_z_via_l1); +TH_REG("rvi", 9, "block id z via L1", rvi09); -static void rv_isel_block_dim_y_via_l1(void) +static void rvi10(void) { build_bir_add(); fake_bir.insts[3].op = BIR_BLOCK_DIM; @@ -272,9 +272,9 @@ static void rv_isel_block_dim_y_via_l1(void) (int16_t)RT_ARG_OFF_BDIM_Y))); PASS(); } -TH_REG("rv_enc", rv_isel_block_dim_y_via_l1); +TH_REG("rvi", 10, "block dim y via L1", rvi10); -static void rv_isel_grid_dim_x_via_l1(void) +static void rvi11(void) { build_bir_add(); fake_bir.insts[3].op = BIR_GRID_DIM; @@ -286,10 +286,10 @@ static void rv_isel_grid_dim_x_via_l1(void) (int16_t)RT_ARG_OFF_GDIM_X))); PASS(); } -TH_REG("rv_enc", rv_isel_grid_dim_x_via_l1); +TH_REG("rvi", 11, "grid dim x via L1", rvi11); /* Bad dimension still refuses cleanly. */ -static void rv_isel_intrinsic_bad_dim(void) +static void rvi12(void) { build_bir_add(); fake_bir.insts[3].op = BIR_THREAD_ID; @@ -299,10 +299,10 @@ static void rv_isel_intrinsic_bad_dim(void) CHEQ(rv_isel_func(&fake_bir, 0, &code), BC_ERR_TDF); PASS(); } -TH_REG("rv_enc", rv_isel_intrinsic_bad_dim); +TH_REG("rvi", 12, "intrinsic bad dim", rvi12); /* PARAM now also reads from L1 runtime args, not from a0..a7. */ -static void rv_isel_param_via_l1(void) +static void rvi13(void) { build_bir_add(); rv_buf_init(&code); @@ -318,13 +318,13 @@ static void rv_isel_param_via_l1(void) (int16_t)RT_ARG_OFF_KARG(2)))); PASS(); } -TH_REG("rv_enc", rv_isel_param_via_l1); +TH_REG("rvi", 13, "param via L1", rvi13); /* A baby core is a single hardware thread with no SIMT, so __syncthreads() * has nothing to wait on and lowers to nothing. FENCE would be wrong to * emit: it is a documented no-op on this core and cannot order anything, * so it would only suggest a guarantee we are not making. */ -static void rv_isel_barrier_is_noop(void) +static void rvi14(void) { build_bir_add(); rv_buf_init(&code); @@ -342,7 +342,7 @@ static void rv_isel_barrier_is_noop(void) CHECK(!buf_contains(rv_fence(0xFu, 0xFu))); PASS(); } -TH_REG("rv_enc", rv_isel_barrier_is_noop); +TH_REG("rvi", 14, "barrier is noop", rvi14); /* ---- Helper: build a minimal "binop kernel" template ---- * @@ -363,7 +363,7 @@ static void build_bir_binop(uint16_t op) * so the only post-prologue change is the single arithmetic word * between the operand loads and the result spill. */ -static void rv_isel_binop_and(void) +static void rvi15(void) { build_bir_binop(BIR_AND); rv_buf_init(&code); @@ -371,9 +371,9 @@ static void rv_isel_binop_and(void) CHECK(buf_contains(rv_and(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_binop_and); +TH_REG("rvi", 15, "binop and", rvi15); -static void rv_isel_binop_or(void) +static void rvi16(void) { build_bir_binop(BIR_OR); rv_buf_init(&code); @@ -381,9 +381,9 @@ static void rv_isel_binop_or(void) CHECK(buf_contains(rv_or(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_binop_or); +TH_REG("rvi", 16, "binop or", rvi16); -static void rv_isel_binop_xor(void) +static void rvi17(void) { build_bir_binop(BIR_XOR); rv_buf_init(&code); @@ -391,9 +391,9 @@ static void rv_isel_binop_xor(void) CHECK(buf_contains(rv_xor(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_binop_xor); +TH_REG("rvi", 17, "binop xor", rvi17); -static void rv_isel_binop_shl(void) +static void rvi18(void) { build_bir_binop(BIR_SHL); rv_buf_init(&code); @@ -401,9 +401,9 @@ static void rv_isel_binop_shl(void) CHECK(buf_contains(rv_sll(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_binop_shl); +TH_REG("rvi", 18, "binop shl", rvi18); -static void rv_isel_binop_lshr(void) +static void rvi19(void) { build_bir_binop(BIR_LSHR); rv_buf_init(&code); @@ -411,9 +411,9 @@ static void rv_isel_binop_lshr(void) CHECK(buf_contains(rv_srl(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_binop_lshr); +TH_REG("rvi", 19, "binop lshr", rvi19); -static void rv_isel_binop_ashr(void) +static void rvi20(void) { build_bir_binop(BIR_ASHR); rv_buf_init(&code); @@ -421,9 +421,9 @@ static void rv_isel_binop_ashr(void) CHECK(buf_contains(rv_sra(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_binop_ashr); +TH_REG("rvi", 20, "binop ashr", rvi20); -static void rv_isel_binop_sdiv(void) +static void rvi21(void) { build_bir_binop(BIR_SDIV); rv_buf_init(&code); @@ -431,9 +431,9 @@ static void rv_isel_binop_sdiv(void) CHECK(buf_contains(rv_div(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_binop_sdiv); +TH_REG("rvi", 21, "binop sdiv", rvi21); -static void rv_isel_binop_udiv(void) +static void rvi22(void) { build_bir_binop(BIR_UDIV); rv_buf_init(&code); @@ -441,9 +441,9 @@ static void rv_isel_binop_udiv(void) CHECK(buf_contains(rv_divu(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_binop_udiv); +TH_REG("rvi", 22, "binop udiv", rvi22); -static void rv_isel_binop_srem(void) +static void rvi23(void) { build_bir_binop(BIR_SREM); rv_buf_init(&code); @@ -451,9 +451,9 @@ static void rv_isel_binop_srem(void) CHECK(buf_contains(rv_rem(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_binop_srem); +TH_REG("rvi", 23, "binop srem", rvi23); -static void rv_isel_binop_urem(void) +static void rvi24(void) { build_bir_binop(BIR_UREM); rv_buf_init(&code); @@ -461,7 +461,7 @@ static void rv_isel_binop_urem(void) CHECK(buf_contains(rv_remu(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_binop_urem); +TH_REG("rvi", 24, "binop urem", rvi24); /* ---- ICMP variants ---- * @@ -475,7 +475,7 @@ static void build_bir_icmp(uint8_t kind) fake_bir.insts[3].subop = kind; } -static void rv_isel_icmp_eq(void) +static void rvi25(void) { build_bir_icmp(BIR_ICMP_EQ); rv_buf_init(&code); @@ -484,9 +484,9 @@ static void rv_isel_icmp_eq(void) CHECK(buf_contains(rv_sltiu(RV_T0, RV_T0, 1))); PASS(); } -TH_REG("rv_enc", rv_isel_icmp_eq); +TH_REG("rvi", 25, "icmp eq", rvi25); -static void rv_isel_icmp_ne(void) +static void rvi26(void) { build_bir_icmp(BIR_ICMP_NE); rv_buf_init(&code); @@ -495,9 +495,9 @@ static void rv_isel_icmp_ne(void) CHECK(buf_contains(rv_sltu(RV_T0, RV_ZERO, RV_T0))); PASS(); } -TH_REG("rv_enc", rv_isel_icmp_ne); +TH_REG("rvi", 26, "icmp ne", rvi26); -static void rv_isel_icmp_slt(void) +static void rvi27(void) { build_bir_icmp(BIR_ICMP_SLT); rv_buf_init(&code); @@ -505,9 +505,9 @@ static void rv_isel_icmp_slt(void) CHECK(buf_contains(rv_slt(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_icmp_slt); +TH_REG("rvi", 27, "icmp slt", rvi27); -static void rv_isel_icmp_sgt(void) +static void rvi28(void) { /* a > b implemented as b < a, so the operands are swapped. */ build_bir_icmp(BIR_ICMP_SGT); @@ -516,9 +516,9 @@ static void rv_isel_icmp_sgt(void) CHECK(buf_contains(rv_slt(RV_T0, RV_T1, RV_T0))); PASS(); } -TH_REG("rv_enc", rv_isel_icmp_sgt); +TH_REG("rvi", 28, "icmp sgt", rvi28); -static void rv_isel_icmp_sle(void) +static void rvi29(void) { /* a <= b == NOT(b < a), so slt with swapped operands plus xori 1. */ build_bir_icmp(BIR_ICMP_SLE); @@ -528,9 +528,9 @@ static void rv_isel_icmp_sle(void) CHECK(buf_contains(rv_xori(RV_T0, RV_T0, 1))); PASS(); } -TH_REG("rv_enc", rv_isel_icmp_sle); +TH_REG("rvi", 29, "icmp sle", rvi29); -static void rv_isel_icmp_ult(void) +static void rvi30(void) { build_bir_icmp(BIR_ICMP_ULT); rv_buf_init(&code); @@ -538,7 +538,7 @@ static void rv_isel_icmp_ult(void) CHECK(buf_contains(rv_sltu(RV_T0, RV_T0, RV_T1))); PASS(); } -TH_REG("rv_enc", rv_isel_icmp_ult); +TH_REG("rvi", 30, "icmp ult", rvi30); /* ---- load_imm32: small constant ---- */ /* Below ADDI's 12-bit signed range, materialisation is a single ADDI @@ -574,7 +574,7 @@ static void build_bir_return_const(int32_t v) fake_bir.insts[0].operands[0] = BIR_MAKE_CONST(0); } -static void rv_isel_const_small_pos(void) +static void rvi31(void) { build_bir_return_const(42); rv_buf_init(&code); @@ -583,9 +583,9 @@ static void rv_isel_const_small_pos(void) CHECK(buf_contains(rv_addi(RV_A0, RV_ZERO, 42))); PASS(); } -TH_REG("rv_enc", rv_isel_const_small_pos); +TH_REG("rvi", 31, "const small pos", rvi31); -static void rv_isel_const_small_neg(void) +static void rvi32(void) { build_bir_return_const(-1); rv_buf_init(&code); @@ -593,9 +593,9 @@ static void rv_isel_const_small_neg(void) CHECK(buf_contains(rv_addi(RV_A0, RV_ZERO, -1))); PASS(); } -TH_REG("rv_enc", rv_isel_const_small_neg); +TH_REG("rvi", 32, "const small neg", rvi32); -static void rv_isel_const_large_positive(void) +static void rvi33(void) { /* 0x12345678: low 12 = 0x678 (bit 11 clear), so just LUI+ADDI * with hi=0x12345 and lo=0x678 (no carry correction). */ @@ -606,9 +606,9 @@ static void rv_isel_const_large_positive(void) CHECK(buf_contains(rv_addi(RV_A0, RV_A0, 0x678))); PASS(); } -TH_REG("rv_enc", rv_isel_const_large_positive); +TH_REG("rvi", 33, "const large positive", rvi33); -static void rv_isel_const_large_signbump(void) +static void rvi34(void) { /* 0x12345fff: low 12 = 0xfff (bit 11 set, lo treated as -1), * so hi must bump from 0x12345 to 0x12346 and ADDI = -1. */ @@ -619,11 +619,11 @@ static void rv_isel_const_large_signbump(void) CHECK(buf_contains(rv_addi(RV_A0, RV_A0, -1))); PASS(); } -TH_REG("rv_enc", rv_isel_const_large_signbump); +TH_REG("rvi", 34, "const large signbump", rvi34); /* ---- UNREACHABLE emits EBREAK ---- */ -static void rv_isel_unreachable_ebreak(void) +static void rvi35(void) { memset(&fake_bir, 0, sizeof(fake_bir)); fake_bir.num_types = 1; @@ -645,7 +645,7 @@ static void rv_isel_unreachable_ebreak(void) CHECK(buf_contains(rv_ebreak())); PASS(); } -TH_REG("rv_enc", rv_isel_unreachable_ebreak); +TH_REG("rvi", 35, "unreachable EBREAK", rvi35); /* ---- Identity casts: PTRTOINT / INTTOPTR / BITCAST ---- */ /* @@ -655,7 +655,7 @@ TH_REG("rv_enc", rv_isel_unreachable_ebreak); * between the load+store pair should match what we'd get with * a no-op transformation. */ -static void rv_isel_cast_id_bitcast(void) +static void rvi36(void) { build_bir_add(); fake_bir.insts[3].op = BIR_BITCAST; @@ -667,7 +667,7 @@ static void rv_isel_cast_id_bitcast(void) CHECK(buf_contains(rv_lw(RV_T0, RV_SP, 12))); /* load %1 slot */ PASS(); } -TH_REG("rv_enc", rv_isel_cast_id_bitcast); +TH_REG("rvi", 36, "cast id bitcast", rvi36); /* ---- Width conversions ---- */ /* @@ -711,7 +711,7 @@ static void build_bir_zext_i8_to_i32(void) fake_bir.insts[2].operands[0] = BIR_MAKE_VAL(1); } -static void rv_isel_zext_i8(void) +static void rvi37(void) { build_bir_zext_i8_to_i32(); rv_buf_init(&code); @@ -720,9 +720,9 @@ static void rv_isel_zext_i8(void) CHECK(buf_contains(rv_srli(RV_T0, RV_T0, 24))); PASS(); } -TH_REG("rv_enc", rv_isel_zext_i8); +TH_REG("rvi", 37, "zext i8", rvi37); -static void rv_isel_sext_i8(void) +static void rvi38(void) { build_bir_zext_i8_to_i32(); fake_bir.insts[1].op = BIR_SEXT; @@ -732,7 +732,7 @@ static void rv_isel_sext_i8(void) CHECK(buf_contains(rv_srai(RV_T0, RV_T0, 24))); PASS(); } -TH_REG("rv_enc", rv_isel_sext_i8); +TH_REG("rvi", 38, "sext i8", rvi38); /* ---- Multi-function module: kernel calls a helper ---- * @@ -842,7 +842,7 @@ static void build_bir_kern_calls_helper(void) fake_bir.insts[9].operands[0] = BIR_MAKE_VAL(8); } -static void rv_isel_module_two_funcs(void) +static void rvi39(void) { build_bir_kern_calls_helper(); rv_buf_init(&code); @@ -859,13 +859,13 @@ static void rv_isel_module_two_funcs(void) CHEQ(count, 2); PASS(); } -TH_REG("rv_enc", rv_isel_module_two_funcs); +TH_REG("rvi", 39, "module two funcs", rvi39); /* The JAL for the kern->helper call should appear with a positive * offset (helper is emitted after kern). The exact value depends * on kern's body length, so we just check that some JAL with a * positive J-immediate exists. */ -static void rv_isel_module_jal_present(void) +static void rvi40(void) { build_bir_kern_calls_helper(); rv_buf_init(&code); @@ -883,12 +883,12 @@ static void rv_isel_module_jal_present(void) CHECK(found_forward_jal); PASS(); } -TH_REG("rv_enc", rv_isel_module_jal_present); +TH_REG("rvi", 40, "module JAL present", rvi40); /* Args land in a0, a1, etc. The call sequence loads the first * two arg values from the kern's local slots into a0 and a1 just * before the JAL. */ -static void rv_isel_module_args_in_regs(void) +static void rvi41(void) { build_bir_kern_calls_helper(); rv_buf_init(&code); @@ -903,14 +903,14 @@ static void rv_isel_module_args_in_regs(void) CHECK(buf_contains(rv_lw(RV_A1, RV_SP, 16))); PASS(); } -TH_REG("rv_enc", rv_isel_module_args_in_regs); +TH_REG("rvi", 41, "module args in regs", rvi41); /* Direct recursion is refused cleanly. A function calling itself * needs more sophisticated prologue/epilogue handling than we have * today; the alternative is to silently miscompile, which is the * sort of decision that turns up later in a debugger as a value * that was perfectly correct an instant ago. */ -static void rv_isel_module_refuses_recursion(void) +static void rvi42(void) { build_bir_kern_calls_helper(); /* Edit function 1 (helper) to call itself instead of just @@ -926,7 +926,7 @@ static void rv_isel_module_refuses_recursion(void) CHEQ(rv_isel_module(&fake_bir, &code), BC_ERR_TDF); PASS(); } -TH_REG("rv_enc", rv_isel_module_refuses_recursion); +TH_REG("rvi", 42, "module refuses recursion", rvi42); /* ---- Forward + backward branches ---- * @@ -975,7 +975,7 @@ static void build_bir_three_blocks(void) fake_bir.insts[2].num_operands = 0; } -static void rv_isel_br_forward_and_backward(void) +static void rvi43(void) { build_bir_three_blocks(); rv_buf_init(&code); @@ -1008,13 +1008,13 @@ static void rv_isel_br_forward_and_backward(void) CHECK(found_back); PASS(); } -TH_REG("rv_enc", rv_isel_br_forward_and_backward); +TH_REG("rvi", 43, "br forward and backward", rvi43); /* ---- BR_COND emits BNE + JAL ---- */ /* `if (cond) goto X; else goto Y;` lowers to loading cond into T0, a * BNE T0, zero, X, then a JAL zero, Y. We construct a two-arm BIR and * check both branch flavours appear in the output. */ -static void rv_isel_br_cond_shape(void) +static void rvi44(void) { memset(&fake_bir, 0, sizeof(fake_bir)); fake_bir.num_types = 2; @@ -1080,7 +1080,7 @@ static void rv_isel_br_cond_shape(void) } PASS(); } -TH_REG("rv_enc", rv_isel_br_cond_shape); +TH_REG("rvi", 44, "br cond shape", rvi44); /* ---- SELECT: BEQ skip + value loads ---- */ /* cond ? a : b lowers to loading cond into T2 and b into T0, a @@ -1089,7 +1089,7 @@ TH_REG("rv_enc", rv_isel_br_cond_shape); * depending on every downstream instruction count, but we can check a * BEQ appears and the buffer parses past it without unresolved * placeholders. */ -static void rv_isel_select_shape(void) +static void rvi45(void) { memset(&fake_bir, 0, sizeof(fake_bir)); fake_bir.num_types = 2; @@ -1141,7 +1141,7 @@ static void rv_isel_select_shape(void) CHECK(found_beq); PASS(); } -TH_REG("rv_enc", rv_isel_select_shape); +TH_REG("rvi", 45, "select shape", rvi45); /* ---- Beyond a0..a7 still works because all params go via L1 ---- * @@ -1151,7 +1151,7 @@ TH_REG("rv_enc", rv_isel_select_shape); * RT_ARG_KERNEL_BASE + i*4 in L1. We construct a kernel with 10 * params and check that the 9th (subop=8) and 10th (subop=9) * read from slot 8 and slot 9 of the kernel-arg block. */ -static void rv_isel_param_9_via_l1(void) +static void rvi46(void) { memset(&fake_bir, 0, sizeof(fake_bir)); fake_bir.num_types = 2; @@ -1187,10 +1187,10 @@ static void rv_isel_param_9_via_l1(void) (int16_t)RT_ARG_OFF_KARG(9)))); PASS(); } -TH_REG("rv_enc", rv_isel_param_9_via_l1); +TH_REG("rvi", 46, "param 9 via L1", rvi46); /* Past RT_ARG_N_KERNEL_SLOTS the isel refuses cleanly. */ -static void rv_isel_param_overflow_refused(void) +static void rvi47(void) { build_bir_add(); fake_bir.insts[1].subop = RT_ARG_N_KERNEL_SLOTS; @@ -1198,7 +1198,7 @@ static void rv_isel_param_overflow_refused(void) CHEQ(rv_isel_func(&fake_bir, 0, &code), BC_ERR_TDF); PASS(); } -TH_REG("rv_enc", rv_isel_param_overflow_refused); +TH_REG("rvi", 47, "param overflow refused", rvi47); /* ---- Struct GEP: field index becomes byte offset ---- * @@ -1210,7 +1210,7 @@ TH_REG("rv_enc", rv_isel_param_overflow_refused); * The BIR uses types 0=i32, 1=struct{i32,i32,f32,i32}, 2=ptr, 3=ptr, 4=void and 5=f32, one i32 const of 1, and * insts PARAM 0, a no-op GEP on idx 0, a GEP to field 1, then RET void. */ -static void rv_isel_struct_gep_field1(void) +static void rvi48(void) { memset(&fake_bir, 0, sizeof(fake_bir)); fake_bir.num_types = 6; @@ -1281,7 +1281,7 @@ static void rv_isel_struct_gep_field1(void) CHECK(buf_contains(rv_addi(RV_T0, RV_T0, 4))); PASS(); } -TH_REG("rv_enc", rv_isel_struct_gep_field1); +TH_REG("rvi", 48, "struct GEP field1", rvi48); /* ---- type_bytes for nested types ---- * @@ -1292,7 +1292,7 @@ TH_REG("rv_enc", rv_isel_struct_gep_field1); * For struct { i8, [4 x i32], i32 } the i8 sits at offset 0, the [4 x i32] * starts at 4 (aligned) with size 16, and the trailing i32 lands at offset * 20, so GEP-to-field-2 should produce ADDI 20. */ -static void rv_isel_struct_gep_with_array(void) +static void rvi49(void) { memset(&fake_bir, 0, sizeof(fake_bir)); fake_bir.num_types = 6; @@ -1358,7 +1358,7 @@ static void rv_isel_struct_gep_with_array(void) CHECK(buf_contains(rv_addi(RV_T0, RV_T0, 20))); PASS(); } -TH_REG("rv_enc", rv_isel_struct_gep_with_array); +TH_REG("rvi", 49, "struct GEP with array", rvi49); /* ---- Array-of-struct GEP: stride access (the bug regression) ---- * @@ -1371,7 +1371,7 @@ TH_REG("rv_enc", rv_isel_struct_gep_with_array); * struct> param and a runtime i32 idx, then GEPs base by idx to yield the * same pointee before RET. The codegen loads base and idx, multiplies idx * by 8 and adds it to base, so we look for the LI of 8 and the MUL. */ -static void rv_isel_array_of_struct_stride(void) +static void rvi50(void) { memset(&fake_bir, 0, sizeof(fake_bir)); fake_bir.num_types = 5; @@ -1425,7 +1425,7 @@ static void rv_isel_array_of_struct_stride(void) CHECK(buf_contains(rv_mul (RV_T1, RV_T1, RV_T2))); PASS(); } -TH_REG("rv_enc", rv_isel_array_of_struct_stride); +TH_REG("rvi", 50, "array of struct stride", rvi50); /* ---- ALLOCA: pointer into the frame's alloca region ---- */ /* @@ -1433,7 +1433,7 @@ TH_REG("rv_enc", rv_isel_array_of_struct_stride); * should be sp + (ISEL_LOCALS_BASE + total_insts*4 + 0). * total_insts here is 2 (the ALLOCA itself and the RET). So * the absolute offset is 8 + 8 = 16. We expect ADDI t0, sp, 16. */ -static void rv_isel_alloca_offset(void) +static void rvi51(void) { memset(&fake_bir, 0, sizeof(fake_bir)); fake_bir.num_types = 3; @@ -1463,7 +1463,7 @@ static void rv_isel_alloca_offset(void) CHECK(buf_contains(rv_addi(RV_T0, RV_SP, 16))); PASS(); } -TH_REG("rv_enc", rv_isel_alloca_offset); +TH_REG("rvi", 51, "alloca offset", rvi51); /* * Regressions for the four silent miscompiles found in the July 2026 @@ -1488,7 +1488,7 @@ static int sp_slot_off(uint32_t w) * the function's own instruction count, so function 1 addressed past its * frame and into the caller's. Every slot must sit inside the largest * frame this module allocates, which is 32 bytes here. */ -static void rv_isel_slots_are_frame_local(void) +static void rvi52(void) { build_bir_kern_calls_helper(); rv_buf_init(&code); @@ -1500,12 +1500,12 @@ static void rv_isel_slots_are_frame_local(void) } PASS(); } -TH_REG("rv_enc", rv_isel_slots_are_frame_local); +TH_REG("rvi", 52, "slots are frame local", rvi52); /* Past roughly 510 instructions the slot offset overflowed the 12-bit * signed immediate and wrapped negative, reading below the frame. It has * to refuse instead. */ -static void rv_isel_refuses_huge_frame(void) +static void rvi53(void) { memset(&fake_bir, 0, sizeof(fake_bir)); fake_bir.num_types = 2; @@ -1536,7 +1536,7 @@ static void rv_isel_refuses_huge_frame(void) CHEQ(rv_isel_func(&fake_bir, 0, &code), BC_ERR_TDF); PASS(); } -TH_REG("rv_enc", rv_isel_refuses_huge_frame); +TH_REG("rvi", 53, "refuses huge frame", rvi53); /* Build a kernel storing through a pointer whose pointee is `w` bits * wide, then loading it back: PARAM ptr, LOAD, STORE, RET. */ @@ -1581,7 +1581,7 @@ static void build_bir_narrow(uint16_t w) /* Sub-word accesses used to emit LW/SW regardless of pointee width, so a * char store clobbered the three neighbouring bytes. */ -static void rv_isel_byte_access(void) +static void rvi54(void) { build_bir_narrow(8); rv_buf_init(&code); @@ -1592,9 +1592,9 @@ static void rv_isel_byte_access(void) CHECK(!buf_contains(rv_sw(RV_T0, RV_T1, 0))); PASS(); } -TH_REG("rv_enc", rv_isel_byte_access); +TH_REG("rvi", 54, "byte access", rvi54); -static void rv_isel_half_access(void) +static void rvi55(void) { build_bir_narrow(16); rv_buf_init(&code); @@ -1603,9 +1603,9 @@ static void rv_isel_half_access(void) CHECK(buf_contains(rv_sh(RV_T0, RV_T1, 0))); PASS(); } -TH_REG("rv_enc", rv_isel_half_access); +TH_REG("rvi", 55, "half access", rvi55); -static void rv_isel_word_access_unchanged(void) +static void rvi56(void) { build_bir_narrow(32); rv_buf_init(&code); @@ -1614,11 +1614,11 @@ static void rv_isel_word_access_unchanged(void) CHECK(buf_contains(rv_sw(RV_T0, RV_T1, 0))); PASS(); } -TH_REG("rv_enc", rv_isel_word_access_unchanged); +TH_REG("rvi", 56, "word access unchanged", rvi56); /* tt-metal reads a0 as the stack watermark, where zero means "not * computed". A void return used to leave it holding whatever was there. */ -static void rv_isel_void_ret_zeroes_a0(void) +static void rvi57(void) { build_bir_narrow(32); rv_buf_init(&code); @@ -1626,7 +1626,7 @@ static void rv_isel_void_ret_zeroes_a0(void) CHECK(buf_contains(rv_addi(RV_A0, RV_ZERO, 0))); PASS(); } -TH_REG("rv_enc", rv_isel_void_ret_zeroes_a0); +TH_REG("rvi", 57, "void ret zeroes a0", rvi57); /* Build two functions where the second one allocas, so the alloca table is * indexed by a nonzero module-global instruction index. */ @@ -1680,7 +1680,7 @@ static void build_bir_helper_allocas(void) * the region was sized per function, so a second function's allocas landed * outside its own frame. Function 1's frame is 8 + 3*4 + 8 = 28, rounded to * 32, with its allocas at 20 and 24. */ -static void rv_isel_alloca_in_second_func(void) +static void rvi58(void) { build_bir_helper_allocas(); rv_buf_init(&code); @@ -1689,12 +1689,12 @@ static void rv_isel_alloca_in_second_func(void) CHECK(buf_contains(rv_addi(RV_T0, RV_SP, 24))); PASS(); } -TH_REG("rv_enc", rv_isel_alloca_in_second_func); +TH_REG("rvi", 58, "alloca in second func", rvi58); /* An unpatched call placeholder is a zero word, which is an illegal RV32 * instruction that traps with no diagnostic. Nothing the module path emits * should ever be zero. */ -static void rv_isel_no_illegal_words(void) +static void rvi59(void) { build_bir_kern_calls_helper(); rv_buf_init(&code); @@ -1703,21 +1703,21 @@ static void rv_isel_no_illegal_words(void) CHECK(!buf_contains(0u)); PASS(); } -TH_REG("rv_enc", rv_isel_no_illegal_words); +TH_REG("rvi", 59, "no illegal words", rvi59); /* An i64 access has no single RV32 instruction. It used to compile to LW and * silently move four of the eight bytes. */ -static void rv_isel_refuses_i64_access(void) +static void rvi60(void) { build_bir_narrow(64); rv_buf_init(&code); CHEQ(rv_isel_func(&fake_bir, 0, &code), BC_ERR_TDF); PASS(); } -TH_REG("rv_enc", rv_isel_refuses_i64_access); +TH_REG("rvi", 60, "refuses i64 access", rvi60); /* Same for an aggregate pointee. */ -static void rv_isel_refuses_struct_access(void) +static void rvi61(void) { build_bir_narrow(32); /* Repoint the pointer at a two-field struct instead of a bare i32. */ @@ -1734,7 +1734,7 @@ static void rv_isel_refuses_struct_access(void) CHEQ(rv_isel_func(&fake_bir, 0, &code), BC_ERR_TDF); PASS(); } -TH_REG("rv_enc", rv_isel_refuses_struct_access); +TH_REG("rvi", 61, "refuses struct access", rvi61); /* Build a single function of n instructions: n-1 PARAMs then a void RET. */ static void build_bir_n_insts(uint32_t n) @@ -1766,7 +1766,7 @@ static void build_bir_n_insts(uint32_t n) /* The frame is 8 + 4n rounded to 16, and the prologue adjusts sp with one * ADDI, so 2032 bytes is the last frame that fits the 12-bit immediate. * That puts the boundary at 506 instructions. */ -static void rv_isel_frame_boundary_ok(void) +static void rvi62(void) { build_bir_n_insts(506); rv_buf_init(&code); @@ -1775,20 +1775,20 @@ static void rv_isel_frame_boundary_ok(void) CHECK(buf_contains(rv_addi(RV_SP, RV_SP, 2032))); PASS(); } -TH_REG("rv_enc", rv_isel_frame_boundary_ok); +TH_REG("rvi", 62, "frame boundary ok", rvi62); -static void rv_isel_frame_boundary_refused(void) +static void rvi63(void) { build_bir_n_insts(507); rv_buf_init(&code); CHEQ(rv_isel_func(&fake_bir, 0, &code), BC_ERR_TDF); PASS(); } -TH_REG("rv_enc", rv_isel_frame_boundary_refused); +TH_REG("rvi", 63, "frame boundary refused", rvi63); /* Every slot access in the largest accepted function stays inside the frame * and within the positive half of the 12-bit immediate. */ -static void rv_isel_max_frame_slots_in_range(void) +static void rvi64(void) { build_bir_n_insts(506); rv_buf_init(&code); @@ -1801,11 +1801,11 @@ static void rv_isel_max_frame_slots_in_range(void) } PASS(); } -TH_REG("rv_enc", rv_isel_max_frame_slots_in_range); +TH_REG("rvi", 64, "max frame slots in range", rvi64); /* i64 has no RV32 representation. This used to compile, with the mul becoming * 32-bit and the shift by 32 becoming a shift by zero. No diagnostic. */ -static void rv_isel_refuses_i64_arith(void) +static void rvi65(void) { memset(&fake_bir, 0, sizeof(fake_bir)); fake_bir.num_types = 3; @@ -1842,4 +1842,4 @@ static void rv_isel_refuses_i64_arith(void) CHEQ(rv_isel_func(&fake_bir, 0, &code), BC_ERR_TDF); PASS(); } -TH_REG("rv_enc", rv_isel_refuses_i64_arith); +TH_REG("rvi", 65, "refuses i64 arith", rvi65); diff --git a/tests/tsched.c b/tests/tsched.c index cd69921..de88005 100644 --- a/tests/tsched.c +++ b/tests/tsched.c @@ -7,7 +7,7 @@ static char obuf[TH_BUFSZ]; /* ---- sched: loads grouped (no wait between two global_load_dword) ---- */ -static void sch_loadpair(void) +static void sch01(void) { int rc = th_run(BC_BIN " --amdgpu tests/test_sched.cu", obuf, TH_BUFSZ); CHEQ(rc, 0); @@ -31,11 +31,11 @@ static void sch_loadpair(void) PASS(); } -TH_REG("sched", sch_loadpair) +TH_REG("sch", 1, "loads group with no wait between them", sch01) /* ---- sched: --no-sched still produces correct output ---- */ -static void sch_nosched(void) +static void sch02(void) { int rc = th_run(BC_BIN " --amdgpu --no-sched tests/test_sched.cu", obuf, TH_BUFSZ); @@ -45,11 +45,11 @@ static void sch_nosched(void) CHECK(strstr(obuf, "global_load_") != NULL); PASS(); } -TH_REG("sched", sch_nosched) +TH_REG("sch", 2, "--no-sched is still correct", sch02) /* ---- sched: compiles to ELF with scheduling ---- */ -static void sch_elf(void) +static void sch03(void) { const char *out = "test_sched_out.hsaco"; char cmd[TH_BUFSZ]; @@ -61,11 +61,11 @@ static void sch_elf(void) remove(out); PASS(); } -TH_REG("sched", sch_elf) +TH_REG("sch", 3, "compiles to ELF with scheduling", sch03) /* ---- sched: all targets compile with scheduling ---- */ -static void sch_targets(void) +static void sch04(void) { static const char *targets[] = { "--gfx1030", "", "--gfx1200" }; static const char *tnames[] = { "gfx1030", "gfx1100", "gfx1200" }; @@ -86,4 +86,4 @@ static void sch_targets(void) } PASS(); } -TH_REG("sched", sch_targets) +TH_REG("sch", 4, "every target compiles with scheduling", sch04) diff --git a/tests/tsmoke.c b/tests/tsmoke.c index 26a2765..28556c8 100644 --- a/tests/tsmoke.c +++ b/tests/tsmoke.c @@ -7,42 +7,42 @@ static char obuf[TH_BUFSZ]; /* ---- smoke: help ---- */ -static void smk_help(void) +static void smk01(void) { int rc = th_run(BC_BIN " --help", obuf, TH_BUFSZ); CHEQ(rc, 0); CHECK(strstr(obuf, "Usage") != NULL); PASS(); } -TH_REG("smoke", smk_help) +TH_REG("smk", 1, "--help prints usage", smk01) /* ---- smoke: no args ---- */ -static void smk_noarg(void) +static void smk02(void) { int rc = th_run(BC_BIN, obuf, TH_BUFSZ); CHNE(rc, 0); PASS(); } -TH_REG("smoke", smk_noarg) +TH_REG("smk", 2, "no args is an error", smk02) /* ---- smoke: version ---- */ -static void smk_vers(void) +static void smk03(void) { SKIP("not implemented"); } -TH_REG("smoke", smk_vers) +TH_REG("smk", 3, "--version", smk03) /* ---- smoke: bad flag ---- */ -static void smk_badf(void) +static void smk04(void) { int rc = th_run(BC_BIN " --nonsense", obuf, TH_BUFSZ); CHNE(rc, 0); PASS(); } -TH_REG("smoke", smk_badf) +TH_REG("smk", 4, "a bad flag is an error", smk04) /* ---- smoke: backend gates ---- */ @@ -65,40 +65,40 @@ static int smk_emit(const char *args, const char *out) return rc == 0 && got; } -static void smk_tcpu(void) +static void smk05(void) { CHECK(smk_emit("--triton --cpu tests/tri_vadd.py", "smk_tc.o")); PASS(); } -TH_REG("smoke", smk_tcpu) +TH_REG("smk", 5, "Triton to CPU object", smk05) -static void smk_trv64(void) +static void smk06(void) { CHECK(smk_emit("--triton --rv64 tests/tri_vadd.py", "smk_tr.o")); PASS(); } -TH_REG("smoke", smk_trv64) +TH_REG("smk", 6, "Triton to RISC-V object", smk06) -static void smk_tptx(void) +static void smk07(void) { CHECK(smk_emit("--triton --nvidia-ptx tests/tri_vadd.py", "smk_tp.ptx")); PASS(); } -TH_REG("smoke", smk_tptx) +TH_REG("smk", 7, "Triton to PTX", smk07) -static void smk_ccpu(void) +static void smk08(void) { CHECK(smk_emit("--cpu tests/canonical.cu", "smk_cc.o")); PASS(); } -TH_REG("smoke", smk_ccpu) +TH_REG("smk", 8, "CUDA to CPU object", smk08) /* The Triton fallthrough returned the lexer's status, so a mode it could * not honour still exited 0. --pp on Python is nonsense and must say so. */ -static void smk_tbad(void) +static void smk09(void) { int rc = th_run(BC_BIN " --triton --pp tests/tri_vadd.py", obuf, TH_BUFSZ); CHNE(rc, 0); PASS(); } -TH_REG("smoke", smk_tbad) +TH_REG("smk", 9, "--pp on Python must not exit 0", smk09) diff --git a/tests/tsoft_fp.c b/tests/tsoft_fp.c index bf04114..1e34068 100644 --- a/tests/tsoft_fp.c +++ b/tests/tsoft_fp.c @@ -45,49 +45,49 @@ static uint32_t bbits(float f) /* ---- Negation ---- */ -static void sfp_neg_one(void) +static void sfp01(void) { CHEQ(bbits(__negsf2(fbits(F_POS_ONE))), F_NEG_ONE); CHEQ(bbits(__negsf2(fbits(F_NEG_ONE))), F_POS_ONE); PASS(); } -TH_REG("soft_fp", sfp_neg_one); +TH_REG("sfp", 1, "neg one", sfp01); -static void sfp_neg_zero(void) +static void sfp02(void) { CHEQ(bbits(__negsf2(fbits(F_POS_ZERO))), F_NEG_ZERO); CHEQ(bbits(__negsf2(fbits(F_NEG_ZERO))), F_POS_ZERO); PASS(); } -TH_REG("soft_fp", sfp_neg_zero); +TH_REG("sfp", 2, "neg zero", sfp02); -static void sfp_neg_inf(void) +static void sfp03(void) { CHEQ(bbits(__negsf2(fbits(F_POS_INF))), F_NEG_INF); CHEQ(bbits(__negsf2(fbits(F_NEG_INF))), F_POS_INF); PASS(); } -TH_REG("soft_fp", sfp_neg_inf); +TH_REG("sfp", 3, "neg inf", sfp03); /* ---- Addition: canonical values ---- */ -static void sfp_add_one_plus_one(void) +static void sfp04(void) { /* 1.0 + 1.0 = 2.0 */ CHEQ(bbits(__addsf3(fbits(F_POS_ONE), fbits(F_POS_ONE))), F_POS_TWO); PASS(); } -TH_REG("soft_fp", sfp_add_one_plus_one); +TH_REG("sfp", 4, "add one plus one", sfp04); -static void sfp_add_one_minus_one(void) +static void sfp05(void) { /* 1.0 + (-1.0) = +0.0 per round-to-nearest. */ CHEQ(bbits(__addsf3(fbits(F_POS_ONE), fbits(F_NEG_ONE))), F_POS_ZERO); PASS(); } -TH_REG("soft_fp", sfp_add_one_minus_one); +TH_REG("sfp", 5, "add one minus one", sfp05); -static void sfp_add_zero_zero(void) +static void sfp06(void) { /* +0 + +0 = +0; -0 + -0 = -0; +0 + -0 = +0. */ CHEQ(bbits(__addsf3(fbits(F_POS_ZERO), fbits(F_POS_ZERO))), F_POS_ZERO); @@ -95,9 +95,9 @@ static void sfp_add_zero_zero(void) CHEQ(bbits(__addsf3(fbits(F_POS_ZERO), fbits(F_NEG_ZERO))), F_POS_ZERO); PASS(); } -TH_REG("soft_fp", sfp_add_zero_zero); +TH_REG("sfp", 6, "add zero zero", sfp06); -static void sfp_add_inf(void) +static void sfp07(void) { /* Inf + Inf = Inf; Inf - Inf = NaN. */ CHEQ(bbits(__addsf3(fbits(F_POS_INF), fbits(F_POS_INF))), F_POS_INF); @@ -105,16 +105,16 @@ static void sfp_add_inf(void) CHEQ(bbits(__addsf3(fbits(F_POS_INF), fbits(F_POS_ONE))), F_POS_INF); PASS(); } -TH_REG("soft_fp", sfp_add_inf); +TH_REG("sfp", 7, "add inf", sfp07); -static void sfp_add_nan(void) +static void sfp08(void) { /* NaN + anything = NaN. */ CHEQ(bbits(__addsf3(fbits(F_QNAN), fbits(F_POS_ONE))), F_QNAN); CHEQ(bbits(__addsf3(fbits(F_POS_ONE), fbits(F_QNAN))), F_QNAN); PASS(); } -TH_REG("soft_fp", sfp_add_nan); +TH_REG("sfp", 8, "add NaN", sfp08); /* ---- Addition vs host FPU on random normal values ---- */ @@ -139,7 +139,7 @@ static float random_normal_float(void) return fbits(bits); } -static void sfp_add_vs_host(void) +static void sfp09(void) { prng_state = 0xC0FFEEu; for (int i = 0; i < 1000; i++) { @@ -162,53 +162,53 @@ static void sfp_add_vs_host(void) } PASS(); } -TH_REG("soft_fp", sfp_add_vs_host); +TH_REG("sfp", 9, "add vs host", sfp09); /* ---- Multiplication ---- */ -static void sfp_mul_one_times_two(void) +static void sfp10(void) { CHEQ(bbits(__mulsf3(fbits(F_POS_ONE), fbits(F_POS_TWO))), F_POS_TWO); PASS(); } -TH_REG("soft_fp", sfp_mul_one_times_two); +TH_REG("sfp", 10, "MUL one times two", sfp10); -static void sfp_mul_sign(void) +static void sfp11(void) { CHEQ(bbits(__mulsf3(fbits(F_NEG_ONE), fbits(F_POS_TWO))), F_NEG_TWO); CHEQ(bbits(__mulsf3(fbits(F_NEG_ONE), fbits(F_NEG_ONE))), F_POS_ONE); PASS(); } -TH_REG("soft_fp", sfp_mul_sign); +TH_REG("sfp", 11, "MUL sign", sfp11); -static void sfp_mul_inf_zero(void) +static void sfp12(void) { /* Inf * 0 is NaN per IEEE. */ CHEQ(bbits(__mulsf3(fbits(F_POS_INF), fbits(F_POS_ZERO))), F_QNAN); CHEQ(bbits(__mulsf3(fbits(F_POS_ZERO), fbits(F_POS_INF))), F_QNAN); PASS(); } -TH_REG("soft_fp", sfp_mul_inf_zero); +TH_REG("sfp", 12, "MUL inf zero", sfp12); -static void sfp_mul_inf_finite(void) +static void sfp13(void) { /* Inf * finite_positive = Inf; Inf * finite_negative = -Inf. */ CHEQ(bbits(__mulsf3(fbits(F_POS_INF), fbits(F_POS_ONE))), F_POS_INF); CHEQ(bbits(__mulsf3(fbits(F_POS_INF), fbits(F_NEG_ONE))), F_NEG_INF); PASS(); } -TH_REG("soft_fp", sfp_mul_inf_finite); +TH_REG("sfp", 13, "MUL inf finite", sfp13); -static void sfp_mul_zero_sign(void) +static void sfp14(void) { /* Sign-of-zero rule: +0 * -1 = -0. */ CHEQ(bbits(__mulsf3(fbits(F_POS_ZERO), fbits(F_NEG_ONE))), F_NEG_ZERO); CHEQ(bbits(__mulsf3(fbits(F_NEG_ZERO), fbits(F_NEG_ONE))), F_POS_ZERO); PASS(); } -TH_REG("soft_fp", sfp_mul_zero_sign); +TH_REG("sfp", 14, "MUL zero sign", sfp14); -static void sfp_mul_vs_host(void) +static void sfp15(void) { prng_state = 0xBADBEEFu; for (int i = 0; i < 1000; i++) { @@ -227,19 +227,19 @@ static void sfp_mul_vs_host(void) } PASS(); } -TH_REG("soft_fp", sfp_mul_vs_host); +TH_REG("sfp", 15, "MUL vs host", sfp15); /* ---- Division ---- */ -static void sfp_div_one_by_two(void) +static void sfp16(void) { /* 1.0 / 2.0 = 0.5 */ CHEQ(bbits(__divsf3(fbits(F_POS_ONE), fbits(F_POS_TWO))), 0x3F000000u); PASS(); } -TH_REG("soft_fp", sfp_div_one_by_two); +TH_REG("sfp", 16, "DIV one by two", sfp16); -static void sfp_div_by_zero(void) +static void sfp17(void) { /* x / 0 = signed Inf for finite x; 0/0 = NaN. */ CHEQ(bbits(__divsf3(fbits(F_POS_ONE), fbits(F_POS_ZERO))), F_POS_INF); @@ -247,17 +247,17 @@ static void sfp_div_by_zero(void) CHEQ(bbits(__divsf3(fbits(F_POS_ZERO), fbits(F_POS_ZERO))), F_QNAN); PASS(); } -TH_REG("soft_fp", sfp_div_by_zero); +TH_REG("sfp", 17, "DIV by zero", sfp17); -static void sfp_div_inf_inf(void) +static void sfp18(void) { /* Inf / Inf = NaN. */ CHEQ(bbits(__divsf3(fbits(F_POS_INF), fbits(F_POS_INF))), F_QNAN); PASS(); } -TH_REG("soft_fp", sfp_div_inf_inf); +TH_REG("sfp", 18, "DIV inf inf", sfp18); -static void sfp_div_vs_host(void) +static void sfp19(void) { prng_state = 0xDEADBABEu; for (int i = 0; i < 1000; i++) { @@ -276,11 +276,11 @@ static void sfp_div_vs_host(void) } PASS(); } -TH_REG("soft_fp", sfp_div_vs_host); +TH_REG("sfp", 19, "DIV vs host", sfp19); /* ---- Comparisons ---- */ -static void sfp_cmp_equal(void) +static void sfp20(void) { CHEQ(__eqsf2(fbits(F_POS_ONE), fbits(F_POS_ONE)), 0); CHEQ(__eqsf2(fbits(F_POS_ONE), fbits(F_POS_TWO)), 1); @@ -288,18 +288,18 @@ static void sfp_cmp_equal(void) CHEQ(__eqsf2(fbits(F_POS_ZERO), fbits(F_NEG_ZERO)), 0); PASS(); } -TH_REG("soft_fp", sfp_cmp_equal); +TH_REG("sfp", 20, "compare equal", sfp20); -static void sfp_cmp_less(void) +static void sfp21(void) { CHECK(__ltsf2(fbits(F_NEG_ONE), fbits(F_POS_ONE)) < 0); CHECK(__ltsf2(fbits(F_POS_ONE), fbits(F_NEG_ONE)) > 0); CHEQ(__ltsf2(fbits(F_POS_ONE), fbits(F_POS_ONE)), 0); PASS(); } -TH_REG("soft_fp", sfp_cmp_less); +TH_REG("sfp", 21, "compare less", sfp21); -static void sfp_cmp_nan(void) +static void sfp22(void) { /* NaN comparisons are unordered; the magnitude comparisons * return a positive value per libgcc convention. */ @@ -308,11 +308,11 @@ static void sfp_cmp_nan(void) CHECK(__lesf2 (fbits(F_QNAN), fbits(F_POS_ONE)) > 0); PASS(); } -TH_REG("soft_fp", sfp_cmp_nan); +TH_REG("sfp", 22, "compare NaN", sfp22); /* ---- Integer to float ---- */ -static void sfp_int_to_float(void) +static void sfp23(void) { CHEQ(bbits(__floatsisf(0)), F_POS_ZERO); CHEQ(bbits(__floatsisf(1)), F_POS_ONE); @@ -320,20 +320,20 @@ static void sfp_int_to_float(void) CHEQ(bbits(__floatsisf(2)), F_POS_TWO); PASS(); } -TH_REG("soft_fp", sfp_int_to_float); +TH_REG("sfp", 23, "int to float", sfp23); -static void sfp_int_to_float_large(void) +static void sfp24(void) { /* 16777216 = 2^24 is exactly representable. */ CHEQ(bbits(__floatsisf(16777216)), 0x4B800000u); CHEQ(bbits(__floatsisf(-16777216)), 0xCB800000u); PASS(); } -TH_REG("soft_fp", sfp_int_to_float_large); +TH_REG("sfp", 24, "int to float large", sfp24); /* ---- Float to integer ---- */ -static void sfp_float_to_int(void) +static void sfp25(void) { CHEQ(__fixsfsi(fbits(F_POS_ZERO)), 0); CHEQ(__fixsfsi(fbits(F_POS_ONE)), 1); @@ -341,9 +341,9 @@ static void sfp_float_to_int(void) CHEQ(__fixsfsi(fbits(F_POS_TWO)), 2); PASS(); } -TH_REG("soft_fp", sfp_float_to_int); +TH_REG("sfp", 25, "float to int", sfp25); -static void sfp_float_to_int_truncates(void) +static void sfp26(void) { /* 1.5 = 0x3FC00000 should truncate to 1. */ CHEQ(__fixsfsi(fbits(0x3FC00000u)), 1); @@ -351,9 +351,9 @@ static void sfp_float_to_int_truncates(void) CHEQ(__fixsfsi(fbits(0xBFC00000u)), -1); PASS(); } -TH_REG("soft_fp", sfp_float_to_int_truncates); +TH_REG("sfp", 26, "float to int truncates", sfp26); -static void sfp_float_to_int_overflow(void) +static void sfp27(void) { /* Inf clamps to INT_MAX / INT_MIN. */ CHEQ(__fixsfsi(fbits(F_POS_INF)), (int32_t)0x7FFFFFFF); @@ -362,7 +362,7 @@ static void sfp_float_to_int_overflow(void) CHEQ(__fixsfsi(fbits(F_QNAN)), 0); PASS(); } -TH_REG("soft_fp", sfp_float_to_int_overflow); +TH_REG("sfp", 27, "float to int overflow", sfp27); /* ---- Subnormal handling (strict-only; SKIP by default) ---- * @@ -372,7 +372,7 @@ TH_REG("soft_fp", sfp_float_to_int_overflow); * we keep them visible so the strict-mode work has a known test * surface to revive. */ -static void sfp_subnormal_unpacks_to_zero(void) +static void sfp28(void) { #if SFP_STRICT_IEEE SKIP("strict-mode subnormal handling not yet implemented"); @@ -385,4 +385,4 @@ static void sfp_subnormal_unpacks_to_zero(void) PASS(); #endif } -TH_REG("soft_fp", sfp_subnormal_unpacks_to_zero); +TH_REG("sfp", 28, "subnormal unpacks to zero", sfp28); diff --git a/tests/tsroa.c b/tests/tsroa.c index b06bb55..52efc21 100644 --- a/tests/tsroa.c +++ b/tests/tsroa.c @@ -13,7 +13,7 @@ /* ---- sroa: a 2-field struct splits, accesses repoint ---- */ -static void sroa_splits(void) +static void sro01(void) { bir_module_t *M = malloc(sizeof(*M)); uint32_t i32, f32, st, pst, pi32, pf32, vt, c0, c1, flds[2]; @@ -95,11 +95,11 @@ static void sroa_splits(void) free(M); PASS(); } -TH_REG("sroa", sroa_splits) +TH_REG("sro", 1, "a two-field struct splits and repoints", sro01) /* ---- sroa: an escaping struct pointer is left untouched ---- */ -static void sroa_escape_bails(void) +static void sro02(void) { bir_module_t *M = malloc(sizeof(*M)); uint32_t i32, f32, st, pst, pi32, vt, c0, flds[2]; @@ -150,4 +150,4 @@ static void sroa_escape_bails(void) free(M); PASS(); } -TH_REG("sroa", sroa_escape_bails) +TH_REG("sro", 2, "an escaping struct pointer is untouched", sro02) diff --git a/tests/tstruct.c b/tests/tstruct.c index 65e4374..4e69bd3 100644 --- a/tests/tstruct.c +++ b/tests/tstruct.c @@ -38,7 +38,7 @@ static int emit_struct_msl(char *out, int cap) } /* The headline: not one goto, not one block label, in the whole file. */ -static void struct_no_gotos(void) +static void str01(void) { int n = emit_struct_msl(msl, sizeof(msl)); CHECK(n > 0); @@ -47,11 +47,11 @@ static void struct_no_gotos(void) CHECK(strstr(msl, ": ;") == NULL); PASS(); } -TH_REG("struct", struct_no_gotos) +TH_REG("str", 1, "structurises without gotos", str01) /* The if and the if/else both turn into real ifs, and the else turns up * exactly where an else is wanted. */ -static void struct_conditionals(void) +static void str02(void) { int n = emit_struct_msl(msl, sizeof(msl)); CHECK(n > 0); @@ -63,12 +63,12 @@ static void struct_conditionals(void) CHECK(strstr(k, "} else {") != NULL && strstr(k, "} else {") < end); PASS(); } -TH_REG("struct", struct_conditionals) +TH_REG("str", 2, "conditionals structurise", str02) /* The counted for becomes a while(true), and both loop-carried values, the * counter and the accumulator, get settled by copies before the lap ends. * If either copy were missing the loop would compute rubbish. */ -static void struct_for_loop(void) +static void str03(void) { int n = emit_struct_msl(msl, sizeof(msl)); CHECK(n > 0); @@ -80,12 +80,12 @@ static void struct_for_loop(void) CHECK(strstr(k, "continue;") != NULL && strstr(k, "continue;") < end); PASS(); } -TH_REG("struct", struct_for_loop) +TH_REG("str", 3, "for loops structurise", str03) /* The while with the break and the continue keeps both: the break leaves * the loop, the continue takes another lap, and neither is a goto in a * trench coat. */ -static void struct_while_break(void) +static void str04(void) { int n = emit_struct_msl(msl, sizeof(msl)); CHECK(n > 0); @@ -98,4 +98,4 @@ static void struct_while_break(void) CHECK(strstr(k, "continue;") != NULL && strstr(k, "continue;") < end); PASS(); } -TH_REG("struct", struct_while_break) +TH_REG("str", 4, "while with break structurises", str04) diff --git a/tests/tsysprint.c b/tests/tsysprint.c index f344621..fdbc17c 100644 --- a/tests/tsysprint.c +++ b/tests/tsysprint.c @@ -63,7 +63,7 @@ static void trace_sink(uint32_t cid, const char *cname, /* ---- Tests ---- */ -static void sp_intern_returns_same_id(void) +static void spr01(void) { bc_sp_reset_globals(); uint32_t a = bc_sp_intern("STEP1.TRACE"); @@ -74,9 +74,9 @@ static void sp_intern_returns_same_id(void) CHNE(a, c); PASS(); } -TH_REG("sysprint", sp_intern_returns_same_id) +TH_REG("spr", 1, "intern returns same id", spr01) -static void sp_class_name_roundtrip(void) +static void spr02(void) { bc_sp_reset_globals(); uint32_t id = bc_sp_intern("STEP2.ERROR"); @@ -85,9 +85,9 @@ static void sp_class_name_roundtrip(void) CHSTR(name, "STEP2.ERROR"); PASS(); } -TH_REG("sysprint", sp_class_name_roundtrip) +TH_REG("spr", 2, "class name roundtrip", spr02) -static void sp_intern_rejects_empty_or_long(void) +static void spr03(void) { bc_sp_reset_globals(); CHEQ(bc_sp_intern(""), BC_SP_CLASS_NONE); @@ -97,9 +97,9 @@ static void sp_intern_rejects_empty_or_long(void) CHEQ(bc_sp_intern(too_long), BC_SP_CLASS_NONE); PASS(); } -TH_REG("sysprint", sp_intern_rejects_empty_or_long) +TH_REG("spr", 3, "intern rejects empty or long", spr03) -static void sp_emit_drain_roundtrip(void) +static void spr04(void) { bc_sp_reset_globals(); cap_reset(); @@ -117,9 +117,9 @@ static void sp_emit_drain_roundtrip(void) CHSTR(cap_payload, "hello from the kernel"); PASS(); } -TH_REG("sysprint", sp_emit_drain_roundtrip) +TH_REG("spr", 4, "emit drain roundtrip", spr04) -static void sp_emitf_formats(void) +static void spr05(void) { bc_sp_reset_globals(); cap_reset(); @@ -135,9 +135,9 @@ static void sp_emitf_formats(void) CHSTR(cap_payload, "tid=7 val=3.14"); PASS(); } -TH_REG("sysprint", sp_emitf_formats) +TH_REG("spr", 5, "emitf formats", spr05) -static void sp_prefix_pattern_routes(void) +static void spr06(void) { /* The Sysprint People's Front and the People's Front of * Sysprint shall both match STEP1.*, which is to say neither @@ -167,9 +167,9 @@ static void sp_prefix_pattern_routes(void) CHSTR(trace_last, "t"); PASS(); } -TH_REG("sysprint", sp_prefix_pattern_routes) +TH_REG("spr", 6, "prefix pattern routes", spr06) -static void sp_star_matches_everything(void) +static void spr07(void) { bc_sp_reset_globals(); cap_reset(); @@ -186,9 +186,9 @@ static void sp_star_matches_everything(void) CHEQ(cap_count, 2); PASS(); } -TH_REG("sysprint", sp_star_matches_everything) +TH_REG("spr", 7, "star matches everything", spr07) -static void sp_overflow_drops(void) +static void spr08(void) { /* Telemetry that exceeds the buffer is silently lost. We do * not allocate more storage on the kernel's behalf. The @@ -210,9 +210,9 @@ static void sp_overflow_drops(void) CHEQ(buf.dropped, 1); PASS(); } -TH_REG("sysprint", sp_overflow_drops) +TH_REG("spr", 8, "overflow drops", spr08) -static void sp_unmatched_class_silently_dropped(void) +static void spr09(void) { bc_sp_reset_globals(); err_count = 0; trace_count = 0; @@ -228,4 +228,4 @@ static void sp_unmatched_class_silently_dropped(void) CHEQ(trace_count, 0); PASS(); } -TH_REG("sysprint", sp_unmatched_class_silently_dropped) +TH_REG("spr", 9, "unmatched class silently dropped", spr09) diff --git a/tests/ttabs.c b/tests/ttabs.c index ad9ffa0..759470f 100644 --- a/tests/ttabs.c +++ b/tests/ttabs.c @@ -4,10 +4,11 @@ #include "tharns.h" #include "amdgpu.h" +#include "lexer.h" /* ---- tables: GFX11 completeness ---- */ -static void tab_gx11(void) +static void tab01(void) { int miss = 0; for (int i = 0; i < AMD_OP_COUNT; i++) { @@ -27,11 +28,11 @@ static void tab_gx11(void) CHEQ(miss, 0); PASS(); } -TH_REG("tables", tab_gx11) +TH_REG("tab", 1, "GFX11 table has no gaps", tab01) /* ---- tables: GFX10 completeness ---- */ -static void tab_gx10(void) +static void tab02(void) { int miss = 0; for (int i = 0; i < AMD_OP_COUNT; i++) { @@ -47,11 +48,11 @@ static void tab_gx10(void) CHEQ(miss, 0); PASS(); } -TH_REG("tables", tab_gx10) +TH_REG("tab", 2, "GFX10 table has no gaps", tab02) /* ---- tables: GFX10 DS mnemonic is ds_read, not ds_load ---- */ -static void tab_mnem(void) +static void tab03(void) { CHECK(amd_enc_table_gfx10[AMD_DS_READ_B32].mnemonic != NULL); CHSTR(amd_enc_table_gfx10[AMD_DS_READ_B32].mnemonic, "ds_read_b32"); @@ -59,11 +60,11 @@ static void tab_mnem(void) CHSTR(amd_enc_table_gfx10[AMD_DS_WRITE_B32].mnemonic, "ds_write_b32"); PASS(); } -TH_REG("tables", tab_mnem) +TH_REG("tab", 3, "GFX10 DS ops are read and write", tab03) /* ---- tables: known opcode differences ---- */ -static void tab_diff(void) +static void tab04(void) { /* s_and_b32: GFX10=0x0E, GFX11=0x16 */ CHEQ(amd_enc_table[AMD_S_AND_B32].hw_opcode, 0x16); @@ -79,4 +80,50 @@ static void tab_diff(void) PASS(); } -TH_REG("tables", tab_diff) +TH_REG("tab", 4, "GFX10 and GFX11 opcodes differ where known", tab04) + +/* ---- tables: keyword table order ---- */ + +/* Nothing checked this until now, which was optimistic of us. A misfiled entry + * lexes its keyword as an identifier and the mess turns up as a parse error + * somewhere downstream, a long way from the typo that caused it. */ + +static void tab05(void) +{ + int n = lexer_kw_count(); + CHECK(n > 0); + for (int i = 1; i < n; i++) { + const char *prev = lexer_kw_at(i - 1), *cur = lexer_kw_at(i); + CHECK(prev != NULL && cur != NULL); + if (strcmp(prev, cur) >= 0) { + printf(" keyword %d (%s) is not after %d (%s)\n", + i, cur, i - 1, prev); + nfail++; + return; + } + } + PASS(); +} +TH_REG("tab", 5, "keyword table is strictly ascending", tab05) + +/* The order matters because of what the search does with it, so check the + * search too rather than only the invariant it relies on. */ +static void tab06(void) +{ + token_t toks[8]; + int n = lexer_kw_count(); + + for (int i = 0; i < n; i++) { + const char *kw = lexer_kw_at(i); + lexer_t L; + lexer_init(&L, kw, (uint32_t)strlen(kw), toks, 8); + lexer_tokenize(&L); + if (L.num_tokens == 0 || toks[0].type == TOK_IDENT) { + printf(" keyword %s (%d) lexes as an identifier\n", kw, i); + nfail++; + return; + } + } + PASS(); +} +TH_REG("tab", 6, "every keyword still lexes as a keyword", tab06) diff --git a/tests/ttdf.c b/tests/ttdf.c index 42d1431..d0ce1d0 100644 --- a/tests/ttdf.c +++ b/tests/ttdf.c @@ -69,7 +69,7 @@ static int dump_to_buf(const td_mod_t *M, char *buf, int bufsz) /* ---- init: zero state, target set ---- */ -static void tdf_init_zero(void) +static void tdf01(void) { /* M is the shared static above */ /* dirty the struct first so we can tell init really wiped it */ @@ -81,11 +81,11 @@ static void tdf_init_zero(void) CHEQ(M.target, TD_TGT_AMD); PASS(); } -TH_REG("tdf", tdf_init_zero); +TH_REG("tdf", 1, "init zero", tdf01); /* ---- solo region: the AMD/NVIDIA degenerate shape ---- */ -static void tdf_solo_rgn(void) +static void tdf02(void) { /* M is the shared static above */ td_init(&M, TD_TGT_NVIDIA); @@ -98,11 +98,11 @@ static void tdf_solo_rgn(void) CHEQ(M.narc, 0); PASS(); } -TH_REG("tdf", tdf_solo_rgn); +TH_REG("tdf", 2, "solo region", tdf02); /* ---- three-region fission: reader/compute/writer ---- */ -static void tdf_three_rgns(void) +static void tdf03(void) { /* M is the shared static above */ td_init(&M, TD_TGT_TENSIX); @@ -119,11 +119,11 @@ static void tdf_three_rgns(void) CHEQ(M.rgns[2].role, TD_RG_WRT); PASS(); } -TH_REG("tdf", tdf_three_rgns); +TH_REG("tdf", 3, "three regions", tdf03); /* ---- channels link two regions with tile shape and depth ---- */ -static void tdf_channel_link(void) +static void tdf04(void) { /* M is the shared static above */ td_init(&M, TD_TGT_TENSIX); @@ -143,11 +143,11 @@ static void tdf_channel_link(void) CHEQ(M.chans[0].tag.layout, TD_LAY_INTRL); PASS(); } -TH_REG("tdf", tdf_channel_link); +TH_REG("tdf", 4, "channel link", tdf04); /* ---- link with bogus producer rejected ---- */ -static void tdf_link_bad_rgn(void) +static void tdf05(void) { /* M is the shared static above */ td_init(&M, TD_TGT_TENSIX); @@ -158,11 +158,11 @@ static void tdf_link_bad_rgn(void) CHEQ(M.ncha, 0); PASS(); } -TH_REG("tdf", tdf_link_bad_rgn); +TH_REG("tdf", 5, "link bad region", tdf05); /* ---- arcs: full CB pipeline (push, wait, pop) plus a NoC read ---- */ -static void tdf_arcs_pipeline(void) +static void tdf06(void) { /* M is the shared static above */ td_init(&M, TD_TGT_TENSIX); @@ -190,11 +190,11 @@ static void tdf_arcs_pipeline(void) CHEQ(M.arcs[3].kind, TD_AR_POP); PASS(); } -TH_REG("tdf", tdf_arcs_pipeline); +TH_REG("tdf", 6, "arcs pipeline", tdf06); /* ---- CB arc with unknown channel rejected ---- */ -static void tdf_arc_bad_chan(void) +static void tdf07(void) { /* M is the shared static above */ td_init(&M, TD_TGT_TENSIX); @@ -204,11 +204,11 @@ static void tdf_arc_bad_chan(void) CHEQ(M.narc, 0); PASS(); } -TH_REG("tdf", tdf_arc_bad_chan); +TH_REG("tdf", 7, "arc bad channel", tdf07); /* ---- lookups return NULL for out-of-range ids ---- */ -static void tdf_lookup_oor(void) +static void tdf08(void) { /* M is the shared static above */ td_init(&M, TD_TGT_AMD); @@ -219,11 +219,11 @@ static void tdf_lookup_oor(void) CHECK(td_arc(&M, 0) == NULL); PASS(); } -TH_REG("tdf", tdf_lookup_oor); +TH_REG("tdf", 8, "lookup out of range", tdf08); /* ---- dump emits something readable for both shapes ---- */ -static void tdf_dump_solo(void) +static void tdf09(void) { /* M is the shared static above */ td_init(&M, TD_TGT_AMD); @@ -236,9 +236,9 @@ static void tdf_dump_solo(void) CHECK(strstr(obuf, "region 0: SOLO") != NULL); PASS(); } -TH_REG("tdf", tdf_dump_solo); +TH_REG("tdf", 9, "dump solo", tdf09); -static void tdf_dump_pipeline(void) +static void tdf10(void) { /* M is the shared static above */ td_init(&M, TD_TGT_TENSIX); @@ -260,11 +260,11 @@ static void tdf_dump_pipeline(void) CHECK(strstr(obuf, "rgn1 WAIT") != NULL); PASS(); } -TH_REG("tdf", tdf_dump_pipeline); +TH_REG("tdf", 10, "dump pipeline", tdf10); /* ---- lower: AMD/NVIDIA solo passthrough ---- */ -static void tdf_lower_solo_amd(void) +static void tdf11(void) { /* M is the shared static above */ td_init(&M, TD_TGT_AMD); @@ -280,9 +280,9 @@ static void tdf_lower_solo_amd(void) CHEQ(out.owns[0], 0); PASS(); } -TH_REG("tdf", tdf_lower_solo_amd); +TH_REG("tdf", 11, "lower solo AMD", tdf11); -static void tdf_lower_solo_nv(void) +static void tdf12(void) { /* M is the shared static above */ td_init(&M, TD_TGT_NVIDIA); @@ -295,11 +295,11 @@ static void tdf_lower_solo_nv(void) CHECK(out.mods[0] == &fake_mod_a); PASS(); } -TH_REG("tdf", tdf_lower_solo_nv); +TH_REG("tdf", 12, "lower solo NVIDIA", tdf12); /* ---- lower: SOLO without a body is rejected ---- */ -static void tdf_lower_solo_no_body(void) +static void tdf13(void) { /* M is the shared static above */ td_init(&M, TD_TGT_AMD); @@ -308,11 +308,11 @@ static void tdf_lower_solo_no_body(void) CHEQ(td_lower(&M, &out), BC_ERR_TDF); PASS(); } -TH_REG("tdf", tdf_lower_solo_no_body); +TH_REG("tdf", 13, "lower solo no body", tdf13); /* ---- lower: AMD with the wrong region role is rejected ---- */ -static void tdf_lower_wrong_role(void) +static void tdf14(void) { /* M is the shared static above */ td_init(&M, TD_TGT_AMD); @@ -322,11 +322,11 @@ static void tdf_lower_wrong_role(void) CHEQ(td_lower(&M, &out), BC_ERR_TDF); PASS(); } -TH_REG("tdf", tdf_lower_wrong_role); +TH_REG("tdf", 14, "lower wrong role", tdf14); /* ---- lower: AMD with channels present is rejected ---- */ -static void tdf_lower_solo_with_chan(void) +static void tdf15(void) { /* M is the shared static above */ td_init(&M, TD_TGT_AMD); @@ -341,11 +341,11 @@ static void tdf_lower_solo_with_chan(void) CHEQ(td_lower(&M, &out), BC_ERR_TDF); PASS(); } -TH_REG("tdf", tdf_lower_solo_with_chan); +TH_REG("tdf", 15, "lower solo with channel", tdf15); /* ---- lower: Tensix passes SOLO through for now (stub) ---- */ -static void tdf_lower_tensix_solo(void) +static void tdf16(void) { /* M is the shared static above */ td_init(&M, TD_TGT_TENSIX); @@ -357,11 +357,11 @@ static void tdf_lower_tensix_solo(void) CHECK(out.mods[0] == &fake_mod_a); PASS(); } -TH_REG("tdf", tdf_lower_tensix_solo); +TH_REG("tdf", 16, "lower tensix solo", tdf16); /* ---- lower: Tensix fission not yet implemented, must say so ---- */ -static void tdf_lower_tensix_fission_stub(void) +static void tdf17(void) { /* M is the shared static above */ td_init(&M, TD_TGT_TENSIX); @@ -373,11 +373,11 @@ static void tdf_lower_tensix_fission_stub(void) CHEQ(td_lower(&M, &out), BC_ERR_TDF); PASS(); } -TH_REG("tdf", tdf_lower_tensix_fission_stub); +TH_REG("tdf", 17, "lower tensix fission stub", tdf17); /* ---- build_solo_from_bir: happy path ---- */ -static void tdf_build_solo_ok(void) +static void tdf18(void) { /* M is the shared static above */ CHEQ(td_build_solo_from_bir(&M, TD_TGT_AMD, &fake_mod_a), BC_OK); @@ -387,22 +387,22 @@ static void tdf_build_solo_ok(void) CHECK(M.rgns[0].body == &fake_mod_a); PASS(); } -TH_REG("tdf", tdf_build_solo_ok); +TH_REG("tdf", 18, "build solo ok", tdf18); /* ---- build_solo_from_bir: NULL body refused ---- */ -static void tdf_build_solo_null(void) +static void tdf19(void) { /* M is the shared static above */ td_init(&M, TD_TGT_AMD); CHEQ(td_build_solo_from_bir(&M, TD_TGT_AMD, NULL), BC_ERR_TDF); PASS(); } -TH_REG("tdf", tdf_build_solo_null); +TH_REG("tdf", 19, "build solo null", tdf19); /* ---- build_solo + lower round-trips back to the same BIR pointer ---- */ -static void tdf_build_then_lower(void) +static void tdf20(void) { /* M is the shared static above */ CHEQ(td_build_solo_from_bir(&M, TD_TGT_NVIDIA, &fake_mod_a), BC_OK); @@ -412,22 +412,22 @@ static void tdf_build_then_lower(void) CHECK(out.mods[0] == &fake_mod_a); PASS(); } -TH_REG("tdf", tdf_build_then_lower); +TH_REG("tdf", 20, "build then lower", tdf20); /* ---- td_tile_bytes returns the right size per dtype ---- */ /* Wormhole tiles are 32x32, so 4096 bytes at fp32 and 2048 at fp16/bf16. * Those numbers feed NoC alignment, CB depth and runtime args, so a * regression here ripples everywhere. */ -static void tdf_tile_bytes_fp32(void) +static void tdf21(void) { td_tag_t t = tile_f32(32, 32, TD_LAY_INTRL); CHEQ(td_tile_bytes(t), 32u * 32u * 4u); PASS(); } -TH_REG("tdf", tdf_tile_bytes_fp32); +TH_REG("tdf", 21, "tile bytes fp32", tdf21); -static void tdf_tile_bytes_bf16(void) +static void tdf22(void) { td_tag_t t; t.rows = 32; t.cols = 32; @@ -435,7 +435,7 @@ static void tdf_tile_bytes_bf16(void) CHEQ(td_tile_bytes(t), 32u * 32u * 2u); PASS(); } -TH_REG("tdf", tdf_tile_bytes_bf16); +TH_REG("tdf", 22, "tile bytes bf16", tdf22); /* ---- placement assigns increasing offsets to every channel ---- */ /* @@ -443,7 +443,7 @@ TH_REG("tdf", tdf_tile_bytes_bf16); * at TD_L1_CB_BASE, second sits one (tile-data + FIFO) block above * it, rounded to 16-byte alignment. */ -static void tdf_place_two_chans(void) +static void tdf23(void) { /* M is the shared static above */ td_init(&M, TD_TGT_TENSIX); @@ -461,7 +461,7 @@ static void tdf_place_two_chans(void) CHEQ(M.chans[c1].l1_off, TD_L1_CB_BASE + 0x2010u); PASS(); } -TH_REG("tdf", tdf_place_two_chans); +TH_REG("tdf", 23, "place two channels", tdf23); /* ---- placement refuses when the budget is exceeded ---- */ /* @@ -469,7 +469,7 @@ TH_REG("tdf", tdf_place_two_chans); * entire CB region. Placement should refuse cleanly rather than * silently wrap or scribble outside L1. */ -static void tdf_place_budget(void) +static void tdf24(void) { /* M is the shared static above */ td_init(&M, TD_TGT_TENSIX); @@ -483,7 +483,7 @@ static void tdf_place_budget(void) CHEQ(td_place_l1(&M), BC_ERR_TDF); PASS(); } -TH_REG("tdf", tdf_place_budget); +TH_REG("tdf", 24, "place budget", tdf24); /* ---- NoC address encoder: spec bit layout ---- */ /* @@ -491,16 +491,16 @@ TH_REG("tdf", tdf_place_budget); * Pick values that exercise each field's high bit so a bad shift * shows up immediately. */ -static void tdf_noc_addr_basic(void) +static void tdf25(void) { /* (x=1, y=2, local=0x1000) packs to 0x1000 | (1<<36) | (2<<42), * which is 0x0000081000001000. */ CHEQ(td_noc_addr(1, 2, 0x1000ull), 0x0000081000001000ull); PASS(); } -TH_REG("tdf", tdf_noc_addr_basic); +TH_REG("tdf", 25, "NoC addr basic", tdf25); -static void tdf_noc_addr_coord_masking(void) +static void tdf26(void) { /* X and Y are 6-bit fields. Anything above bit 5 must be * masked off rather than spilling into the reserved zone. */ @@ -508,9 +508,9 @@ static void tdf_noc_addr_coord_masking(void) (0x3Full << 36) | (0x3Full << 42)); PASS(); } -TH_REG("tdf", tdf_noc_addr_coord_masking); +TH_REG("tdf", 26, "NoC addr coord masking", tdf26); -static void tdf_noc_addr_local_36bits(void) +static void tdf27(void) { /* Local address is 36 bits, anything above must be masked off * since the high 28 bits of the 64-bit word are reserved. */ @@ -518,7 +518,7 @@ static void tdf_noc_addr_local_36bits(void) TD_NOC_LOCAL_MASK); PASS(); } -TH_REG("tdf", tdf_noc_addr_local_36bits); +TH_REG("tdf", 27, "NoC addr local 36bits", tdf27); /* ---- orchestrator: RD goes on NoC 0, WR on NoC 1 ---- */ /* @@ -526,7 +526,7 @@ TH_REG("tdf", tdf_noc_addr_local_36bits); * write arc, run the orchestrator, check the noc_id and length * fields come out per the spec. */ -static void tdf_noc_orchestrate_dirs(void) +static void tdf28(void) { td_init(&M, TD_TGT_TENSIX); uint16_t rdr = td_mkrgn(&M, TD_RG_RDR); @@ -547,7 +547,7 @@ static void tdf_noc_orchestrate_dirs(void) CHEQ(M.arcs[wr].length, 32u * 32u * 4u); PASS(); } -TH_REG("tdf", tdf_noc_orchestrate_dirs); +TH_REG("tdf", 28, "NoC orchestrate dirs", tdf28); /* ---- orchestrator: refuses transfers above 8 KiB ---- */ /* @@ -555,7 +555,7 @@ TH_REG("tdf", tdf_noc_orchestrate_dirs); * The orchestrator must refuse rather than truncate; the splitting * pass that handles oversize transfers is not implemented yet. */ -static void tdf_noc_orchestrate_oversize(void) +static void tdf29(void) { td_init(&M, TD_TGT_TENSIX); uint16_t rdr = td_mkrgn(&M, TD_RG_RDR); @@ -568,14 +568,14 @@ static void tdf_noc_orchestrate_oversize(void) CHEQ(td_noc_orchestrate(&M), BC_ERR_TDF); PASS(); } -TH_REG("tdf", tdf_noc_orchestrate_oversize); +TH_REG("tdf", 29, "NoC orchestrate oversize", tdf29); /* ---- orchestrator leaves CB arcs untouched ---- */ /* * PUSH/WAIT/RSV/POP do not cross the NoC; their noc_id and length * fields should be left at zero, not silently mutated. */ -static void tdf_noc_skips_cb_arcs(void) +static void tdf30(void) { td_init(&M, TD_TGT_TENSIX); uint16_t rdr = td_mkrgn(&M, TD_RG_RDR); @@ -592,7 +592,7 @@ static void tdf_noc_skips_cb_arcs(void) CHEQ(M.arcs[wait].length, 0u); PASS(); } -TH_REG("tdf", tdf_noc_skips_cb_arcs); +TH_REG("tdf", 30, "NoC skips circular buffer arcs", tdf30); /* ---- integration: fission on test_dce.cu (one stored ptr) ---- */ /* @@ -602,7 +602,7 @@ TH_REG("tdf", tdf_noc_skips_cb_arcs); * AMDGPU backend purely so the C99 pipeline runs all the way to * BIR; --tdf-fission then bails before any AMDGPU work happens. */ -static void tdf_fission_store_only(void) +static void tdf31(void) { int rc = th_run(BC_BIN " --tdf-fission --amdgpu-bin -o " "/tmp/_tdf_fission_test.hsaco tests/test_dce.cu", @@ -628,7 +628,7 @@ static void tdf_fission_store_only(void) CHECK(strstr(obuf, "RD chan") == NULL); PASS(); } -TH_REG("tdf", tdf_fission_store_only); +TH_REG("tdf", 31, "fission store only", tdf31); /* ---- integration: fission on matmul (1 store + 2 loads) ---- */ /* @@ -638,7 +638,7 @@ TH_REG("tdf", tdf_fission_store_only); * (for C) plus two RDR -> CMP channels (for A and B), six arcs * total. */ -static void tdf_fission_matmul_shape(void) +static void tdf32(void) { int rc = th_run(BC_BIN " --tdf-fission --amdgpu-bin -o " "/tmp/_tdf_fission_matmul.hsaco tests/canonical.cu", @@ -663,11 +663,11 @@ static void tdf_fission_matmul_shape(void) CHECK(strstr(obuf, "noc1 len=4096") != NULL); PASS(); } -TH_REG("tdf", tdf_fission_matmul_shape); +TH_REG("tdf", 32, "fission matmul shape", tdf32); /* ---- region overflow returns TD_BAD_ID after the table fills ---- */ -static void tdf_region_overflow(void) +static void tdf33(void) { /* M is the shared static above */ td_init(&M, TD_TGT_TENSIX); @@ -680,11 +680,11 @@ static void tdf_region_overflow(void) CHEQ(M.nrgn, TD_MAX_RGNS); PASS(); } -TH_REG("tdf", tdf_region_overflow); +TH_REG("tdf", 33, "region overflow", tdf33); /* ---- CB-arc emit lowers each kind to the right primitive at placed addrs ---- */ -static void tdf_emit_cb_arcs(void) +static void tdf34(void) { td_init(&M, TD_TGT_TENSIX); uint16_t rdr = td_mkrgn(&M, TD_RG_RDR); @@ -729,11 +729,11 @@ static void tdf_emit_cb_arcs(void) CHECK(bufs_eq(&EA, &EB)); PASS(); } -TH_REG("tdf", tdf_emit_cb_arcs); +TH_REG("tdf", 34, "emit circular buffer arcs", tdf34); /* ---- the CB emitter refuses NoC arcs (those are emitted elsewhere) ---- */ -static void tdf_emit_cb_rejects_noc(void) +static void tdf35(void) { td_init(&M, TD_TGT_TENSIX); uint16_t rdr = td_mkrgn(&M, TD_RG_RDR); @@ -746,11 +746,11 @@ static void tdf_emit_cb_rejects_noc(void) CHEQ(td_emit_cb_arc(&M, &M.arcs[a_rd], &EA), BC_ERR_TDF); PASS(); } -TH_REG("tdf", tdf_emit_cb_rejects_noc); +TH_REG("tdf", 35, "emit circular buffer rejects NoC", tdf35); /* ---- reader loop: DRAM -> L1, read barrier, back-jump ---- */ -static void tdf_reader_loop_shape(void) +static void tdf36(void) { rv_buf_init(&EA); CHEQ(td_emit_dma_loop(&EA, /*is_write=*/0, @@ -773,11 +773,11 @@ static void tdf_reader_loop_shape(void) CHEQ((last >> 7) & 0x1fu, 0u); /* rd == zero */ PASS(); } -TH_REG("tdf", tdf_reader_loop_shape); +TH_REG("tdf", 36, "reader loop shape", tdf36); /* ---- writer loop: L1 -> DRAM, write-ack barrier ---- */ -static void tdf_writer_loop_shape(void) +static void tdf37(void) { rv_buf_init(&EA); CHEQ(td_emit_dma_loop(&EA, /*is_write=*/1, @@ -790,7 +790,7 @@ static void tdf_writer_loop_shape(void) CHECK(has_word(&EA, rv_lw(RV_T2, RV_T0, 0x04))); PASS(); } -TH_REG("tdf", tdf_writer_loop_shape); +TH_REG("tdf", 37, "writer loop shape", tdf37); /* * The --tdf and --tdf-fission flags have to appear in every one of main.c's @@ -823,16 +823,16 @@ static const char *run_flag(const char *flag) return buf; } -static void tdf_flag_reaches_pass(void) +static void tdf38(void) { const char *o = run_flag("--tdf"); CHECK(strstr(o, "TDF module") != NULL); CHECK(strstr(o, "translation_unit") == NULL); PASS(); } -TH_REG("tdf", tdf_flag_reaches_pass); +TH_REG("tdf", 38, "flag reaches pass", tdf38); -static void tdf_fission_flag_reaches_pass(void) +static void tdf39(void) { const char *o = run_flag("--tdf-fission"); CHECK(strstr(o, "TDF module") != NULL); @@ -843,4 +843,4 @@ static void tdf_fission_flag_reaches_pass(void) CHECK(strstr(o, "WRT") != NULL); PASS(); } -TH_REG("tdf", tdf_fission_flag_reaches_pass); +TH_REG("tdf", 39, "fission flag reaches pass", tdf39); diff --git a/tests/ttmc.c b/tests/ttmc.c index 7113cd0..41463d9 100644 --- a/tests/ttmc.c +++ b/tests/ttmc.c @@ -68,7 +68,7 @@ static int n_compute(const uint32_t *w, int n) /* One instruction per encoding format with operands zero, so each word is just * its opcode byte, and a pseudo-op in the middle that must be dropped. */ -static void ttmc_bin_opcodes(void) +static void tmc01(void) { uint32_t k = 0; set_op(k++, TT_SFPNOP); /* 0x02, FMT_C */ @@ -92,11 +92,11 @@ static void ttmc_bin_opcodes(void) CHECK(has_word(w, n, 0x8F000000u)); /* sfpwnop */ PASS(); } -TH_REG("ttmc", ttmc_bin_opcodes); +TH_REG("tmc", 1, "bin opcodes", tmc01); /* The baby core pushes each Tensix word through a custom instruction encoded as * the word rol-2'd, so both emitters must agree word-for-word under rol2. */ -static void ttmc_ttinsn_is_rol2(void) +static void tmc02(void) { uint32_t k = 0; set_op(k++, TT_SFPNOP); @@ -120,11 +120,11 @@ static void ttmc_ttinsn_is_rol2(void) CHEQX(tti[i], rol2(bin[i])); PASS(); } -TH_REG("ttmc", ttmc_ttinsn_is_rol2); +TH_REG("tmc", 2, "ttinsn is rol2", tmc02); /* rol-2 only round-trips because every real opcode is below 0xC0000000, leaving * the top two bits free. Check the ceiling holds and ror-2 restores each word. */ -static void ttmc_ttinsn_reversible(void) +static void tmc03(void) { uint32_t k = 0; set_op(k++, TT_SFPLOAD); @@ -150,12 +150,12 @@ static void ttmc_ttinsn_reversible(void) } PASS(); } -TH_REG("ttmc", ttmc_ttinsn_reversible); +TH_REG("tmc", 3, "ttinsn reversible", tmc03); /* Pin the Sync Unit field packing, not just the opcode byte, using values * distinct from the bracket's own sync words so each match is unambiguous. The * bit layouts these goldens encode are the field definitions in emit.c. */ -static void ttmc_sync_fields(void) +static void tmc04(void) { set_op(0, TT_SEMINIT); /* sem_sel 3, init 1, max 2 */ set_imm(0, 0, 3); set_imm(0, 1, 1); set_imm(0, 2, 2); @@ -179,4 +179,4 @@ static void ttmc_sync_fields(void) CHECK(has_word(w, n, 0xA4000000u | (5u << 2))); PASS(); } -TH_REG("ttmc", ttmc_sync_fields); +TH_REG("tmc", 4, "sync fields", tmc04); diff --git a/tests/ttriton.c b/tests/ttriton.c index c99167f..fa14afa 100644 --- a/tests/ttriton.c +++ b/tests/ttriton.c @@ -27,7 +27,7 @@ static int tt_run(const char *args) * Lexer * ============================================================ */ -static void tt_lex_simple(void) +static void tri01(void) { int rc = tt_run("--triton --lex tests/tri_simple.py"); CHEQ(rc, 0); @@ -42,9 +42,9 @@ static void tt_lex_simple(void) CHECK(strstr(obuf, "def") != NULL); PASS(); } -TH_REG("triton", tt_lex_simple) +TH_REG("tri", 1, "lexer simple", tri01) -static void tt_lex_slop_docstring(void) +static void tri02(void) { /* The AI slop file is enormous and the lex dump is enormous; the * 4 KB capture buffer truncates the trailing summary. We confirm @@ -59,13 +59,13 @@ static void tt_lex_slop_docstring(void) CHECK(strstr(obuf, "import") != NULL); PASS(); } -TH_REG("triton", tt_lex_slop_docstring) +TH_REG("tri", 2, "lexer slop docstring", tri02) /* ============================================================ * Parser * ============================================================ */ -static void tt_parse_funcdef(void) +static void tri03(void) { int rc = tt_run("--triton --parse tests/tri_vadd.py"); CHEQ(rc, 0); @@ -75,9 +75,9 @@ static void tt_parse_funcdef(void) CHECK(strstr(obuf, "Block") != NULL); PASS(); } -TH_REG("triton", tt_parse_funcdef) +TH_REG("tri", 3, "parse funcdef", tri03) -static void tt_parse_expressions(void) +static void tri04(void) { int rc = tt_run("--triton --parse tests/tri_vadd.py"); CHEQ(rc, 0); @@ -90,9 +90,9 @@ static void tt_parse_expressions(void) CHECK(strstr(obuf, "Keyword") != NULL); PASS(); } -TH_REG("triton", tt_parse_expressions) +TH_REG("tri", 4, "parse expressions", tri04) -static void tt_parse_goblin(void) +static void tri05(void) { /* The goblin kernel is identical in shape to vector_add but with * names the Triton tutorial people would never choose. We are @@ -105,13 +105,13 @@ static void tt_parse_goblin(void) CHECK(strstr(obuf, "Param 'hoard_ptr'") != NULL); PASS(); } -TH_REG("triton", tt_parse_goblin) +TH_REG("tri", 5, "parse goblin", tri05) /* ============================================================ * Sema * ============================================================ */ -static void tt_sema_resolves(void) +static void tri06(void) { int rc = tt_run("--triton --sema tests/tri_vadd.py"); CHEQ(rc, 0); @@ -125,9 +125,9 @@ static void tt_sema_resolves(void) CHECK(strstr(obuf, "type(constexpr)") != NULL); PASS(); } -TH_REG("triton", tt_sema_resolves) +TH_REG("tri", 6, "sema resolves", tri06) -static void tt_sema_locals_and_params(void) +static void tri07(void) { int rc = tt_run("--triton --sema tests/tri_vadd.py"); CHEQ(rc, 0); @@ -136,9 +136,9 @@ static void tt_sema_locals_and_params(void) CHECK(strstr(obuf, "Name 'BLOCK_SIZE' -> param") != NULL); PASS(); } -TH_REG("triton", tt_sema_locals_and_params) +TH_REG("tri", 7, "sema locals and params", tri07) -static void tt_sema_ai_slop_still_resolves(void) +static void tri08(void) { /* The AI slop kernel uses comically long names but every one of * them should still bind correctly. The point of this test is to @@ -154,9 +154,9 @@ static void tt_sema_ai_slop_still_resolves(void) CHECK(strstr(obuf, "0 error(s)") != NULL); PASS(); } -TH_REG("triton", tt_sema_ai_slop_still_resolves) +TH_REG("tri", 8, "sema ai slop still resolves", tri08) -static void tt_sema_shapes_vector_add(void) +static void tri09(void) { /* Sitting one of tile shape inference. The vector add kernel uses * arange to build a rank-1 tile, broadcasts a scalar against it, @@ -171,9 +171,9 @@ static void tt_sema_shapes_vector_add(void) CHECK(strstr(obuf, "vec[?]:int1") != NULL); PASS(); } -TH_REG("triton", tt_sema_shapes_vector_add) +TH_REG("tri", 9, "sema shapes vector add", tri09) -static void tt_sema_shapes_2d_broadcast(void) +static void tri10(void) { /* Rank-2 broadcasting. The matmul-shape kernel exercises both * sides of the canonical [:, None] / [None, :] reshape pattern, @@ -193,9 +193,9 @@ static void tt_sema_shapes_2d_broadcast(void) CHECK(strstr(obuf, "mat[?, ?]:float32") != NULL); PASS(); } -TH_REG("triton", tt_sema_shapes_2d_broadcast) +TH_REG("tri", 10, "sema shapes 2d broadcast", tri10) -static void tt_sema_shapes_scalar_kernel(void) +static void tri11(void) { /* The simple all-scalar kernel should have no tile shapes at all; * every expression is rank 0 with dtype int32. We check both that @@ -208,9 +208,9 @@ static void tt_sema_shapes_scalar_kernel(void) CHECK(strstr(obuf, "mat[") == NULL); PASS(); } -TH_REG("triton", tt_sema_shapes_scalar_kernel) +TH_REG("tri", 11, "sema shapes scalar kernel", tri11) -static void tt_sema_constexpr_resolves_dim(void) +static void tri12(void) { /* BLOCK: tl.constexpr = 256 should propagate so arange and the * downstream broadcast get vec[256] rather than vec[?]. */ @@ -220,9 +220,9 @@ static void tt_sema_constexpr_resolves_dim(void) CHECK(strstr(obuf, "vec[?]") == NULL); PASS(); } -TH_REG("triton", tt_sema_constexpr_resolves_dim) +TH_REG("tri", 12, "sema constexpr resolves dim", tri12) -static void tt_lower_matmul(void) +static void tri13(void) { /* Rank-2 tiles with tl.dot now lower for the CPU path: the tile is * materialized and fully unrolled (block sizes are constexpr), so a @@ -234,9 +234,9 @@ static void tt_lower_matmul(void) CHECK(strstr(obuf, "E138") == NULL); PASS(); } -TH_REG("triton", tt_lower_matmul) +TH_REG("tri", 13, "lower matmul", tri13) -static void tt_lower_matmul_kloop(void) +static void tri14(void) { /* Matmul with a runtime K-loop lowers to a counted loop: the * accumulator is scratch-backed and the loop counter is a phi, so @@ -248,13 +248,13 @@ static void tt_lower_matmul_kloop(void) CHECK(strstr(obuf, "E138") == NULL); PASS(); } -TH_REG("triton", tt_lower_matmul_kloop) +TH_REG("tri", 14, "lower matmul kloop", tri14) /* ============================================================ * BIR Lowering * ============================================================ */ -static void tt_lower_thread_model(void) +static void tri15(void) { int rc = tt_run("--triton --ir tests/tri_vadd.py"); CHEQ(rc, 0); @@ -262,9 +262,9 @@ static void tt_lower_thread_model(void) CHECK(strstr(obuf, "thread_id") != NULL); PASS(); } -TH_REG("triton", tt_lower_thread_model) +TH_REG("tri", 15, "lower thread model", tri15) -static void tt_lower_simple_arith(void) +static void tri16(void) { /* tri_simple chains a handful of integer operations and ends in * a store, so every arithmetic instruction is reachable from a @@ -282,9 +282,9 @@ static void tt_lower_simple_arith(void) CHECK(strstr(obuf, "store ") != NULL); PASS(); } -TH_REG("triton", tt_lower_simple_arith) +TH_REG("tri", 16, "lower simple arith", tri16) -static void tt_lower_arithmetic(void) +static void tri17(void) { /* The lowerer emits the icmp for the mask, but DCE removes it * downstream because mask= is currently warned-and-ignored on @@ -304,9 +304,9 @@ static void tt_lower_arithmetic(void) CHECK(strstr(obuf, "store f32") != NULL); PASS(); } -TH_REG("triton", tt_lower_arithmetic) +TH_REG("tri", 17, "lower arithmetic", tri17) -static void tt_lower_math_intrinsics(void) +static void tri18(void) { int rc = tt_run("--triton --ir tests/tri_math.py"); CHEQ(rc, 0); @@ -325,9 +325,9 @@ static void tt_lower_math_intrinsics(void) CHECK(strstr(obuf, "E134") == NULL); PASS(); } -TH_REG("triton", tt_lower_math_intrinsics) +TH_REG("tri", 18, "lower math intrinsics", tri18) -static void tt_lower_where_select(void) +static void tri19(void) { int rc = tt_run("--triton --ir tests/tri_where.py"); CHEQ(rc, 0); @@ -336,9 +336,9 @@ static void tt_lower_where_select(void) CHECK(strstr(obuf, "E134") == NULL); PASS(); } -TH_REG("triton", tt_lower_where_select) +TH_REG("tri", 19, "lower where select", tri19) -static void tt_lower_where_mixed_promotes(void) +static void tri20(void) { int rc = tt_run("--triton --ir --no-cfold tests/tri_where_mixed.py"); CHEQ(rc, 0); @@ -347,9 +347,9 @@ static void tt_lower_where_mixed_promotes(void) CHECK(strstr(obuf, "E134") == NULL); PASS(); } -TH_REG("triton", tt_lower_where_mixed_promotes) +TH_REG("tri", 20, "lower where mixed promotes", tri20) -static void tt_lower_where_i32_select(void) +static void tri21(void) { int rc = tt_run("--triton --ir --no-dce tests/tri_where_int.py"); CHEQ(rc, 0); @@ -358,9 +358,9 @@ static void tt_lower_where_i32_select(void) CHECK(strstr(obuf, "E134") == NULL); PASS(); } -TH_REG("triton", tt_lower_where_i32_select) +TH_REG("tri", 21, "lower where i32 select", tri21) -static void tt_lower_where_bad_cond(void) +static void tri22(void) { int rc = tt_run("--triton --ir tests/tri_where_bad_cond.py"); CHNE(rc, 0); @@ -368,13 +368,13 @@ static void tt_lower_where_bad_cond(void) CHECK(strstr(obuf, "condition must lower") != NULL); PASS(); } -TH_REG("triton", tt_lower_where_bad_cond) +TH_REG("tri", 22, "lower where bad cond", tri22) /* ============================================================ * Backend: AMD GFX11 * ============================================================ */ -static void tt_amd_hsaco(void) +static void tri23(void) { int rc = tt_run("--triton --amdgpu-bin tests/tri_vadd.py " "-o tri_vadd.hsaco"); @@ -386,9 +386,9 @@ static void tt_amd_hsaco(void) CHECK(strstr(obuf, "1 kernels") != NULL); PASS(); } -TH_REG("triton", tt_amd_hsaco) +TH_REG("tri", 23, "AMD hsaco", tri23) -static void tt_amd_goblin_hsaco(void) +static void tri24(void) { /* The goblin variant should produce the same shape of output as * vector_add. The point is to confirm the backend is name-agnostic. */ @@ -398,13 +398,13 @@ static void tt_amd_goblin_hsaco(void) CHECK(th_exist("tri_goblin.hsaco")); PASS(); } -TH_REG("triton", tt_amd_goblin_hsaco) +TH_REG("tri", 24, "AMD goblin hsaco", tri24) /* ============================================================ * Backend: NVIDIA PTX * ============================================================ */ -static void tt_nvidia_ptx(void) +static void tri25(void) { int rc = tt_run("--triton --nvidia-ptx tests/tri_vadd.py " "-o tri_vadd.ptx"); @@ -426,13 +426,13 @@ static void tt_nvidia_ptx(void) CHECK(strstr(buf, ".version") != NULL); PASS(); } -TH_REG("triton", tt_nvidia_ptx) +TH_REG("tri", 25, "nvidia PTX", tri25) /* ============================================================ * Backend: Tensix Metalium * ============================================================ */ -static void tt_tensix_metalium(void) +static void tri26(void) { int rc = tt_run("--triton --tensix tests/tri_vadd.py " "-o tri_vadd_compute.cpp"); @@ -455,7 +455,7 @@ static void tt_tensix_metalium(void) CHECK(strstr(buf, "cb_wait_front") != NULL); PASS(); } -TH_REG("triton", tt_tensix_metalium) +TH_REG("tri", 26, "tensix metalium", tri26) /* ============================================================ * The AI Slop End-to-End Test @@ -466,7 +466,7 @@ TH_REG("triton", tt_tensix_metalium) * robust to whatever the LLM industry serves up as Triton next year. * That feels worth a regression test of its very own. */ -static void tt_amd_ai_slop_still_compiles(void) +static void tri27(void) { int rc = tt_run("--triton --amdgpu-bin tests/tri_slop.py " "-o tri_slop.hsaco"); @@ -479,4 +479,4 @@ static void tt_amd_ai_slop_still_compiles(void) * fine and any wildly different value means something is off. */ PASS(); } -TH_REG("triton", tt_amd_ai_slop_still_compiles) +TH_REG("tri", 27, "AMD ai slop still compiles", tri27) diff --git a/tests/ttypes.c b/tests/ttypes.c index edd5a9e..7c0c942 100644 --- a/tests/ttypes.c +++ b/tests/ttypes.c @@ -17,7 +17,7 @@ static void typ_setup(void) /* ---- types: interning ---- */ -static void typ_intrn(void) +static void typ01(void) { typ_setup(); @@ -73,11 +73,11 @@ static void typ_intrn(void) PASS(); } -TH_REG("types", typ_intrn) +TH_REG("typ", 1, "interning returns one id per type", typ01) /* ---- types: constants ---- */ -static void typ_const(void) +static void typ02(void) { typ_setup(); @@ -111,11 +111,11 @@ static void typ_const(void) PASS(); } -TH_REG("types", typ_const) +TH_REG("typ", 2, "constant types", typ02) /* ---- types: strings ---- */ -static void typ_strs(void) +static void typ03(void) { typ_setup(); @@ -129,11 +129,11 @@ static void typ_strs(void) PASS(); } -TH_REG("types", typ_strs) +TH_REG("typ", 3, "string table interning", typ03) /* ---- types: name tables ---- */ -static void typ_names(void) +static void typ04(void) { CHSTR(bir_op_name(BIR_ADD), "add"); CHSTR(bir_op_name(BIR_BARRIER), "barrier"); @@ -153,11 +153,11 @@ static void typ_names(void) PASS(); } -TH_REG("types", typ_names) +TH_REG("typ", 4, "name tables", typ04) /* ---- types: printer doesn't segfault ---- */ -static void typ_print(void) +static void typ05(void) { typ_setup(); @@ -208,4 +208,4 @@ static void typ_print(void) PASS(); } -TH_REG("types", typ_print) +TH_REG("typ", 5, "the printer does not segfault", typ05) diff --git a/tests/twarpsize.c b/tests/twarpsize.c index 5fa74fc..1e5b275 100644 --- a/tests/twarpsize.c +++ b/tests/twarpsize.c @@ -1,6 +1,6 @@ #include "tharns.h" -static void cmd_warpsize(void) { +static void wsz01(void) { char obuf[TH_BUFSZ]; // 1. Verify gfx942 produces Wave64 (constant 64) @@ -13,4 +13,4 @@ static void cmd_warpsize(void) { PASS(); } -TH_REG("compile", cmd_warpsize) +TH_REG("wsz", 1, "warpsize", wsz01)