add matmul custom impl - #82
Conversation
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: dec7e3ed9c
ℹ️ About Codex in GitHub
Codex has been enabled to automatically review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
When you sign up for Codex through ChatGPT, Codex can also answer questions or update the PR, like "@codex address that feedback".
| # SELF_VERIFY (tM=32 tN=32 tK=64), full check | ||
| make TESTCASE=matmul TYPE=MATMUL_TEST tM=32 tN=32 tK=64 self_verify=on verify_elems=65535 | ||
| # RES_CHECK (tM=16 tN=16 tK=16) | ||
| make TESTCASE=matmul TYPE=MATMUL_TEST tM=16 tN=16 tK=16 res_check=on |
There was a problem hiding this comment.
The new res_check=on commands cannot perform the advertised host-side verification: the inspected src/gfrun_matmul.py::parse_matmul_shape only accepts names matching matmul_MASK_..._M<digits>, so this matmul_MATMUL_TEST_... artifact is skipped, and its input generator also never creates the required shape.txt. Add MATMUL_TEST parsing and shape-file generation, or provide a dedicated runner, before including these commands in the result-check workflow.
Useful? React with 👍 / 👎.
| ifneq ($(NAME), ) | ||
| TARGET = $(ELF_HEAD)/$(NAME)_MATMUL_TEST_tM$(tM)_tN$(tN)_tK$(tK).elf | ||
| else | ||
| TARGET = $(ELF_HEAD)/$(TESTCASE)_MATMUL_TEST_tM$(tM)_tN$(tN)_tK$(tK).elf |
There was a problem hiding this comment.
When compile.all builds a given tile configuration first with self_verify=on and then with res_check=on, both invocations resolve to this identical target path even though their compiled definitions and behavior differ. The later RES_CHECK build therefore overwrites the SELF_VERIFY artifact, so the script does not retain the self-verifying binaries it claims to build; encode the mode (and preferably VERIFY_ELEMS) in the target name or emit only one intended variant.
Useful? React with 👍 / 👎.
| float v = (float)dst[i]; | ||
| float err = fabsf(v - golden); | ||
| if (err > max_err) max_err = err; | ||
| if (err > tol) ++mismatches; |
There was a problem hiding this comment.
If a kernel defect produces NaN in any checked output, fabsf(v - golden) is also NaN, and both err > max_err and this err > tol test evaluate false. A result containing only such failures can therefore print PASS and return zero; explicitly treat non-finite v or err as a mismatch.
Useful? React with 👍 / 👎.
| __half *src0 = (__half *)malloc(gM * gK * sizeof(__half)); | ||
| __half *src1 = (__half *)malloc(gK * gN * sizeof(__half)); | ||
| __half *dst = (__half *)malloc(gM * gN * sizeof(__half)); |
There was a problem hiding this comment.
On the Linx target these pointers are passed directly to TLOAD/TSTORE, whose global-memory buffers must be 4 KiB aligned, but ordinary malloc only guarantees fundamental alignment. The existing FP16 matmul harness explicitly rounds all three buffers to a 4 KiB boundary; these allocations, repeated in every new mode, can therefore cause misaligned DMA accesses or incorrect results unless an aligned allocator or the repository's padded ALIGN_MASK + ALIGN pattern is used.
Useful? React with 👍 / 👎.
| const float tol = atol + rtol * fabsf(golden); | ||
|
|
||
| int total = gM * gN; | ||
| int check_count = total < VERIFY_ELEMS ? total : VERIFY_ELEMS; |
There was a problem hiding this comment.
The supposedly full-check configuration sets VERIFY_ELEMS=65535, but this cap still restricts verification to a prefix whenever the runtime matrix is larger. For example, shape 257 257 K has 66,049 outputs, so the final two rows—including the dynamic M-tail and corner most likely to expose boundary bugs—are never checked and the run can still report PASS; represent full verification explicitly or use total rather than a fixed cap in that mode.
Useful? React with 👍 / 👎.
1 participant
自定义matmul算子开发