Repository navigation
Conversation
The __device__ helper kernels took AdmBufferCuda, AdmFixedParametersCuda, VifBufferCuda and filter_table_stuct by value. nvcc forwards those copies to the kernel's param space, but LLVM's NVPTX backend materialises the 584 and 248 byte structs in a local memory depot (via a byte-by-byte ld.param/st.local loop) as soon as they are indexed with a runtime value such as blockIdx.z. Pass them by const reference instead, so both compilers read the kernel parameters in place. adm_cm needed a const pointer for params.i_rfactor as a result. The filter1d horizontal kernels also reduce their seven 64-bit per-thread accumulators in a plain loop. nvcc unrolls it on its own; clang keeps it rolled because the inlined warp_reduce body exceeds its unroll threshold, which forces the accumulator array into local memory. Mark those loops with #pragma unroll like the neighbouring ones. With this, clang-compiled PTX has the same local memory footprint as nvcc's for every kernel and, measured with Nsight Systems over 600 1080p frames on an RTX 5090, total kernel time is within 3% of the nvcc build (178.9 ms vs 183.8 ms); before, adm_csf ran 57x and filter1d 4.5x slower. nvcc's own output is unchanged, and integer_adm and integer_vif remain bit-identical between the two compilers.
|
The second part, the clang build path this prepares for, is #1596. |
|
Tested on master The local-memory claim reproduces.
All other kernels in adm_decouple.cu, motion_score.cu and the rest of filter1d.cu are at 0 B before and after. nvcc (the default build): Scores:
|
|
Thanks for reproducing this on your 4090. You're right about filter1d, with nvcc the resource usage and scores are unchanged, but the SASS isn't byte-identical. Most likely that's the added #pragma unroll, since the two ADM files only got the by-reference change and stayed identical. I've reworded the description! |
2 participants
The device helper kernels take AdmBufferCuda, AdmFixedParametersCuda, VifBufferCuda and filter_table_stuct by value. nvcc forwards those copies to param space, but clang's NVPTX backend copies the 584 and 248 byte structs into a local memory depot with a byte loop as soon as they are indexed with a runtime value like blockIdx.z. The filter1d horizontal kernels have the same problem with their seven int64 accumulators: nvcc unrolls that loop on its own, clang doesn't, and the array lands in local memory.
Passing the structs by const reference and adding #pragma unroll on those loops fixes both. No kernel logic changes. With nvcc, resource usage and scores are unchanged (same local memory footprint, integer_adm and integer_vif bit identical before and after). The filter1d SASS is reordered but has the same instruction count. With clang, adm_csf went from 57x slower than nvcc to parity and total kernel time is within 3% of nvcc
This is the first half of making -Denable_nvcc=false usable without the CUDA toolkit headers, see #1596