Amd waves cap - #352
Closed
gwoltman wants to merge 304 commits into
Closed
Amd waves cap#352gwoltman wants to merge 304 commits into
gwoltman wants to merge 304 commits into
Conversation
Tailsquare now computes 2ab as (a+b)^2 - a^2 - b^2. Since a^2 and b^2 are already computed, this replaces a cmul with a csq which saves a wideMul.
… handle it -- caused local memory usage.
…lHalve the way doubles does. FracBits overhauled so FP32 hybrid FFTs can use fracBits for optional doubling and halving.
…e no longer doubled for optionalDouble and optionalHalve.
…r the larger FP32 hybrid FFTs.
…tested. I have a GPU that has twice failed proof generation and the next lower proof power succeeded. This will automate the process rather than raising an error and aborting which requires manual intervention. Up next see if we figure out how to catch this problem earlier (during write of proof residues?)
…:path.c_str() returns a wide char in Windows.
…ommand queues. Prep work eeded for trying dispatching kernels to multiple queues.
…n with -use MULTI_Q. Slower on 5070Ti due to increase L2 cache hits. Significantly faster on a 5080. No change on a 4080 which has same L2 cache size as a 5080. Needs more study.
…rs trouble (fma on floats). The code was not used anyway.
…rable L2 cache locality of working on NTT data one data type at a time has been moved to a central location. Code is much more readable. Paves the way for further improvements in L2 cache locality work.
… 1080 Gpus use asm by default.
…d one's
After clLinkProgram, compile() asks p1 for its build log again -- the same
program whose log was already printed a few lines above, right after
clCompileProgram. So the linker's own diagnostics are never shown, and a
link failure reports only
Linking 'carryfused.cl' error LINK_PROGRAM_FAILURE (-17) (args )
with nothing about why. That is not hypothetical: it is exactly what an
aarch64 POCL device printed while the kernels were being brought up there,
and the reason had to be found by other means.
The same line also repeats every compile warning. With a #warning
injected into transpose.cl, an Intel iGPU run prints it 8 times before
this change and 4 after.
Ask p2 instead, guarded, since a failed clLinkProgram may return no
program at all and then there is nothing to query.
Residues unchanged on an Intel iGPU and on POCL.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
carryFused's !OLD_FENCE carry hand-off has one lane per wavefront publish
the ready flag on behalf of its wavefront, and one lane per wavefront spin
on it, with sync() as the only thing ordering the other lanes around it.
That works on AMD, where a wavefront advances in lock-step and sync() can
be free, and on nVidia, where sync() is bar.warp.sync. On anything else
sync() compiles to nothing, the hand-off races, and the answer is silently
wrong -- on an Intel iGPU, -use OLD_FENCE=0 returns
EE 400 8f64862427ff0b6f
where every other configuration of the same run returns
0d42972a9d02005c.
A barrier is not a substitute: the sync() at carryfused.cl:224 sits inside
"if (gr < H / WMUL && me >= (WMUL-1) * G_W)", which the whole workgroup
does not enter, so a workgroup barrier deadlocks there -- tried, and the
run dies with OUT_OF_RESOURCES on kCarryFused. The protocol wants a
wavefront-scoped primitive, and off AMD and nVidia there is not one.
base.cl already defaults OLD_FENCE to 1 when not AMDGPU. Make that hold
when it is set explicitly as well, the same way MULTI_Q and GRAPHS are
overridden under -time. After this, -use OLD_FENCE=0 on an Intel iGPU
logs a line and returns the right residue.
Found by sweeping every -use knob against the invariant that none of them
may change the residue. It was the only knob that did.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The kernel smoke job runs one configuration, so it can only catch a kernel that is broken everywhere. The failures this code actually has are the other kind: correct in the shipped configuration and wrong in another one, silently, because nothing ever runs that path. There is a cheap invariant for exactly that. A -use knob is a performance setting, so it must not change the answer: after a fixed number of iterations the residue depends only on the exponent. Run the same check under a few knobs and require one residue. TAIL_KERNELS=0 swaps the double-wide tail kernels for single-wide ones and WMUL=1 changes how carryFused partitions its workgroup, so between them the three runs cover rather more of the kernel set than one did. Each run starts from scratch, since otherwise prpll resumes from the previous run's savefile and the second and third runs would have nothing to do. Verified on a POCL CPU device: the three configurations agree on 0d42972a9d02005c, and adding a knob that does change the residue (OLD_FENCE=0, which races where no wavefront is in lock-step) makes the step exit 1 with the two residues in the message. Costs two more runs of about 20 s each on a warm cache. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…cal_id
The double-wide tailMul kernels run G_H*2 threads over two lines and work
out lowMe = me % G_H for the lane within a line. Every access uses lowMe
except the MUL_LOW read, which calls
read(G_H, NH, p, a, memline * SMALL_HEIGHT);
and read() indexes with get_local_id(0) -- which in these kernels runs
0..2*G_H-1. The second half-workgroup therefore reads the line shifted by
one i step, and at i == NH-1 it reads into the next memline, past the end
of the buffer for the last one. The single-wide kernels are fine: there
me == lowMe.
This breaks proof generation outright, and only proof generation, because
tailMulLow is reached only from Gpu::exponentiate via Proof. A PRP run and
its Gerbicz check never touch it, so everything looks healthy. On a 7900
XTX, a complete PRP of 786433 on current master:
786433 proof: invalid (316d17434d034dc2 expected 0abb7aeca5890b24)
786433 Proof 'proof-tmp/786433-10.proof' verification FAILED
786433 Proof generation failed; reporting the result without a proof
at every proof power from 10 down to 4. With -use TAIL_KERNELS=0 or 1
(single-wide) the same run verifies OK, which puts it on the double-wide
path; with this fix the default verifies OK for FFT types 0, 3 and 4, and
the PRP residue is bfb6553c004222ba either way, since only the proof path
is affected.
Add an overload of read() that takes the lane explicitly and use it at the
four sites, one per number type.
A failed -verify used to exit 0, so nothing but the log could tell a good
proof from a bad one. VERIFY tasks come only from the command line, so
that is always a one-shot job: throw, which the existing handling in main()
turns into exit 1.
Add a CI step that would have caught the kernel bug: verify a known-good
proof once per FFT type. Verification runs tailSquare and tailMulLow for a
few thousand iterations, and the six types are the ones allShapes()
enumerates, so every number type and hybrid is covered. A proof is just a
list of residues, so one file serves all six. On a POCL CPU device the
whole step takes about two minutes, and on an unpatched tree it fails.
test/786433-10.proof is 1.0 MB: 11 residues of a 786433-bit exponent, the
smallest an FP64 FFT can do at its 3.0 bits-per-word minimum. Proof power
10 keeps verification short (the final exponentiation is E/1024 squarings).
It was generated on a 7900 XTX and verifies on POCL and an Intel iGPU too.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The LDS reduction in updateStats guards its barriers with
if (num_threads > WAVEFRONT) bar();
which is the assumption 0e67b41 removed from bar(WG) and barsync(): that a
group no wider than a wavefront advances in lock-step and so needs neither
an execution barrier nor a fence. True on AMD GCN; not true on nVidia
Volta and later; not true at all on an Intel GPU or a CPU device.
With G_W == 64 and a 32-lane wavefront the loop runs at num_threads = 64,
32 and 16, so two of its three steps write and read LDS with no barrier
and no fence between them. Threads then miss values written by the upper
half and the reduction returns something smaller than the real maximum.
Use bar(num_threads), which now decides that by what the hardware
guarantees: AMD returns immediately exactly as before, sm_60 and later
reconverge the warp and fence LDS, anything else takes a real barrier.
num_threads is a compile-time workgroup size (G_W and friends), so every
thread makes the same number of passes and reaches both calls.
Measured, -fft 256:2:256 -block 200 -iters 1200:
POCL CPU device master ROEmax 0.117, ROEavg 0.091
patched ROEmax 0.127, ROEavg 0.095
Intel iGPU master ROEmax 0.127, ROEavg 0.094
patched ROEmax 0.127, ROEavg 0.095
0.127 is the true figure: the iGPU reaches it either way, and the patched
CPU run now agrees with it, while master on the CPU device loses the
maximum entirely. Residues are identical throughout (0d42972a9d02005c at
400, b39c99fbd03b3454 at 1200) -- only the statistic was wrong.
That statistic is not only cosmetic: the comments in fftbpw.h say the
bits-per-word tables were computed "by targeting maxROE of ~0.35 over 1000
iterations", and one entry records lowering a limit after a run failed at
ROEmax 0.294. A maximum that reads low makes an FFT size look safer than
it is.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
… cannot work
Four problems in the LDSMUL > 1 path, all reachable only by opting in, and
none of them reported by anything today. Every change here is the identity
when SBMUL == 1, which is every configuration that does not opt in --
checked by enumerating 180 combinations of numWG, SHUFL_BYTES, WG, RADIX
and LDSPAD.
1. SBMUL need not divide numWG. The workgroups are partitioned into groups
of SBMUL sharing one region and one semaphore, so the partition needs
ceil(numWG/SBMUL) groups spanning ceil(numWG/SBMUL)*SBMUL regions, but
LDS_BYTES allocated numWG of them and LDSinit initialised numWG/SBMUL
semaphores while the highest index in use is (numWG-1)/SBMUL. With
WG=64, RADIX=4, SHUFL_BYTES=8 and -use WMUL=3,LDSMUL_W=2 that is 2144
bytes of LDS written past the end of the array, and semaphore 1 read
uninitialised -- which either hangs on a stale 1, or slips through the
spin loop without the lock.
2. The spin loop tested the wrong thing. atomic_cmpxchg returns the old
value, so only 0 means the lock was taken; "== 1" let any other value
out of the loop unlocked, after which LDStx_end cleared a semaphore this
workgroup never held. Now "!= 0".
3. The "4 semaphores" assert was the wrong invariant and compiled to
nothing without -use DEBUG. The space is now sized from the group count
instead of hardwired, so there is nothing left to assert: -use
WMUL=10,LDSMUL_W=2 wrote 5 semaphores into room for 4.
4. Two combinations cannot work and now fail to build rather than quietly
misbehave:
- variant 2 with LDSMUL > 1. partial_tabMul4/8 read and write the
shared region outside the LDStx lock with only a bar(WG), which does
not cover the other workgroups sharing that memory. This is what the
"partitioned_LDS is a nightmare" note was about; refusing it also puts
LDSptr's truncating divide out of reach, since variant 2 is its only
caller and the truncation only occurs when sharing leaves
LDS_SHUFL_BYTES off a 16-byte boundary.
- LDSMUL >= 3 with SHUFL_BYTES == 4, where SBMUL*SHUFL_BYTES is 12 and
shufl's >= 16 / == 8 / == 4 chain has no branch: it would return with
the data unexchanged and no diagnostic.
I have no nVidia hardware, and LDS sharing calls barsync(), which exists
only where HAS_PTX >= 200, so none of this can run on an AMD, Intel or CPU
device here. What I could check: the arithmetic, enumerated above; that
the sharing path still compiles, with clang -fsyntax-only for an
nvptx64 target across LDSMUL 1/2/4 and SHUFL_BYTES 4/8/16; that the two
guards fire for exactly the two rejected combinations and nothing else; and
that residues are unchanged on an Intel iGPU and POCL, with defaults and
with -use TAIL_KERNELS=2,WMUL=4.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The LDSPAD "second RADIX == 8" case of the 32-bit shufl, in both the
8-byte and the 4-byte paths, picks between two index expressions on
WG == 64. The else arm lays the data out as exactly eight padded rows,
which only inverts when WG/64 == 8, and reads a row of 64 regardless of
WG. Enumerating the resulting permutation against the generic shufl:
WG = 64 0 of 512 wrong
WG = 128 896 of 1024 wrong (48 reads of slots never written)
WG = 256 1792 of 2048 wrong (32 reads of slots never written)
WG = 512 0 of 4096 wrong
So it happens to be right at exactly the two sizes that reach it today --
512 (WG 64) and 4096 (WG 512) -- and wrong at everything between. Not a
bug now: FP32 and GF31 route width/height 1024 through shufl_and_fft2
rather than plain shufl(f=8), so WG 128 never gets here, and no RADIX-8
shape produces WG 256. It goes live the moment a new shape or a different
1K dispatch does.
The 64-bit sibling a few hundred lines above solves the same problem with
an expression that is general in WG, and at WG == 64 its two arms are the
same expression, so one form serves every size:
write lds[i * (WG + 8) + lowMe]
read lds[i * (WG / 64) * 8 + (lowMe / 64) * 8 + ((lowMe / 8) & 7) * (WG + 8) + (lowMe & 7)]
Enumerated at WG 64, 128, 256 and 512: a bijection every time, agreeing
with the generic shufl element for element, no unwritten slot read, and
the highest index one below the allocation ((WG*RADIX + 56), which these
paths use) in all four cases. The WG == 64 branch goes away with it --
that branch is what the whole family of these bugs is made of.
The same shape is in the LDSPAD RADIX == 4 reads, whose else arm balances
only at WG == 256. Those are correct at both workgroup sizes RADIX 4 can
have today (WG 64 from a 256-wide shape, WG 256 from a 1024 one should the
commented-out clause in nW()/nH() be re-enabled), and generalising them
measured as no faster and no slower while adding 16 instructions to
fftw.cl, so they are left as they are and guarded instead: a RADIX-4
workgroup of any other size now fails to build, with a message saying what
to do about it, rather than silently reading slots nothing wrote. The
guard fires for exactly that case -- WG 64 and 256 at RADIX 4 still build,
as does every RADIX 8 size.
The LDSSWIZ masks have the same character one size lower: enumerating all
eight swizzle cases, five of them fold rows together at WG 32 and return
64 to 192 wrong elements of 256, while every one is exact from WG 64 up.
Nothing reaches them at WG 32 either (that branch of fft_common asks for
f=1,r=4 at RADIX 8, and f=4,r=8, which no swizzle case matches), so they
get a guard too rather than a rewrite: LDSSWIZ now requires WG >= 64.
Residues unchanged on an Intel iGPU and on POCL at 256:2:256, 512:2:256
and 1K:2:256, and with -use SHUFL_BYTES_W=4 and SHUFL_BYTES_H=4.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Since ae8aa0c added the double-wide and split-kernel variants to tailMul, TAIL_KERNELS selects the shape of both tail kernels, not just tailSquare. The help text still described it as tailSquare only. The TAIL_TRIGS line just below is left as it is: it says tailSquare, and for FP64 and FP32 that remains exactly right, because tailMul has no TAIL_TRIGS branches -- it always computes the trig values from scratch. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
-use TAIL_KERNELS governs tailMul as well now
updateStats: barrier the roundoff reduction by hardware, not by size
tailMul: read the "low" line with the half-workgroup lane, not get_local_id
shufl: the 4-byte path must honour r, like the 8- and 16-byte paths
shufl: fix the LDSSWIZ read index for workgroups other than 64
Disabling MULTI_Q under -time must reach the kernels too
TrigBufCache: give each key field its own bits instead of adding them
Refuse OLD_FENCE=0 where no wavefront runs in lock-step
LDSbar: drop the lock-step early return, as bar(WG) and barsync() did
KernelCompiler: print the linked program's build log, not the compiled one's
CI: require the residue to be the same under several -use settings
LDS sharing: fix the region and semaphore arithmetic, and refuse what cannot work
shufl: make the FP32/GF31 "second RADIX == 8" index workgroup-general
…_RESTRICT=1 Commit 65d5f08 marked the Trig* and BigTab* typedefs restrict. On nVidia that turns the table loads into ld.global.nc and lets the compiler batch them early, which raised the register need (FP64 carryFused 72 -> 96, tailSquare 80 -> 118). The default carryFused register cap (80) then spilled to local memory, about 12% slower on a TITAN V. Restrict is now off by default and enabled with the (undocumented) -use ENABLE_RESTRICT=1. Measured: default carryFused fits with no spill and matches the Sep 1 build; restrict plus REGCF64=96 is about 1% faster on a TITAN V; no measurable effect on Radeon VII / MI50 or on the 1:512:8:512:202 NTT. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
The known-keys list in Gpu.cpp is the set of officially supported -use settings, i.e. the ones users are encouraged to try. ENABLE_RESTRICT is dangerous without intimate knowledge of its register implications, so like REGCF64 it stays an undocumented setting. It still works; it just draws the "unrecognized -use key" warning. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…ce middle kernels on Vega The in-place fftMiddleIn / fftMiddleOut kernels compile to 133 / 134 VGPRs on gfx906. Above 128 VGPRs a SIMD holds only one wave, and with 256-thread workgroups that leaves one workgroup per CU, so every barrier of the 16x16 LDS transpose stalls the CU. This, not the memory layout, is why INPLACE=1 was slower on AMD. Add an AMD analog of the CUDA register cap: Gpu::amdWavesPerEu() passes -DAMD_WAVES_PER_EU=n to a kernel and the new KERNEL_CAP macro turns it into amdgpu_waves_per_eu(n). It is applied to the fftMiddleIn/Out, tailSquare/Mul and carryFused kernels. The default is 2 waves per SIMD for the middle kernels when the transform is in place on Vega class GPUs (gfx900/902/904/906/909/90c); everything else is unchanged. Overrides (undocumented, like REGCF64): -use WPE_MIDIN=n, WPE_MIDOUT=n, WPE_TAIL=n, WPE_CARRY=n, where 0 means the compiler default. Measured at 512:16:512:202, workers=1, INPLACE=1: Radeon VII 1232 -> 1080 us/iter, MI50 1020 -> 924 us/iter (out of place: 994 and 906, unchanged). The cap costs 2 spilled VGPRs in fftMiddleIn. Capping to 3 or more waves, or capping tailSquare or carryFused, was slower, so those are not defaulted. CUDA PTX is identical for all kernels. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…tc.) options instead of new WPE_* keys
numCudaRegisters() already selects the per-kernel, per-FFT-type option key (REGMI64, REGMO31, REGTS61, REGCF64, ...). Let that selection
run on every backend and hand the option value to the new Gpu::amdRegisterOption(), which follows the CUDA conventions:
0 = not specified (default), -1 = compiler default (no cap),
1..10 = minimum waves per SIMD (the AMD counterpart of CUDA launch bounds; 10 is the GCN maximum),
more than 10 = explicit VGPR count (amdgpu_num_vgpr; rocm generates its usual code and then spills to fit).
The default is unchanged: 2 waves per SIMD for the in-place middle kernels on Vega class GPUs. The WPE_* keys are gone.
KERNEL_CAP now understands both -DAMD_WAVES_PER_EU=n and -DAMD_NUM_VGPR=n. The CUDA path is unchanged (identical -v register reports).
Radeon VII, INPLACE=1: default 1079 us/iter (128/127 VGPRs), REGMI64=-1,REGMO64=-1 1231 (133/134), REGMI64=3,REGMO64=3 1076 (84/84),
REGMI64=128,REGMO64=128 1080 (128/127). MI50 default 922. All residues match.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
The function selects the per-kernel REGxxxx option for every backend: CUDA turns the value into --maxrregcount / launch bounds, AMD (amdRegisterOption) into a waves-per-SIMD request or a VGPR count. Rename it and update the comments accordingly. The per-kernel switch has to stay compiled on non-CUDA backends: besides the CUDA default register counts it selects which REGxxxx option key applies to the kernel. Guarding it with #if CUDA_BACKEND left the key empty on AMD and silently disabled all the REG* overrides there (REGMI64=-1,REGMO64=-1 still compiled the in-place middle kernels at 128/127 VGPRs instead of 133/134). Only the CUDA formatting tail is guarded. Verified: AMD default 128/127, REGMI64=-1 133/134, =3 84/84, =128 128/127; CUDA -v register reports identical to the previous build. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…kernels The tailSquare (and tailMul) .cl file is compiled once per K() instance and its code object holds both the main kernel and the Zero kernel that handles lines 0 and H/2. A REGTS64 request therefore also capped the Zero kernel, which naturally needs 126 VGPRs (the main kernel needs 84-90), so a 3-wave request forced 72 spilled VGPRs into a kernel run by two workgroups. Keep the Zero kernels on plain KERNEL(); the cap applies to the main kernels only. With the cap confined to the main kernel, -use REGTS64=3 (3 waves per SIMD) is a small win where tailSquare sits just over the 84-VGPR boundary: MI50 out of place 85 -> 84 VGPRs (5 spilled): 912 -> 891 us/iter (p6) and 998 -> 982 (p4); Radeon VII in place 90 -> 80: 1078 -> 1057. Not made a default: it depends on the kernel being within a few VGPRs of the boundary. CUDA register reports are identical. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Collaborator
Author
|
Oops. I did not mean to create a pull request for merging into GPUowl. I'm still getting a handle on the Claude workflow. |
Collaborator
Author
|
Withdrawing the pull request. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
6 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Added option to expose amd_waves_per_gpu and num_vgpr. Default amd_waves_per_eu to 2 for fttMiddleIn/Out when INPLACE=1. This improves INPLACE=1 greatly, but is still slower than INPLACE=0.