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 enable_nvcc=false path added in Netflix#1436 compiles the kernels with clang but still relies on the toolkit's cuda_runtime.h; -nocudainc was left commented out because device intrinsics broke without it. The toolkit headers do not support a MinGW-w64 host: with _WIN32 defined but not _MSC_VER, crt/math_functions.hpp takes a legacy code path whose isinf/isnan and __signbit declarations conflict with clang's own CUDA wrappers. So the clang path cannot be used on Windows without an MSVC installation, and add_languages('cuda') rules the nvcc path out there as well. Finish the -nocudainc approach instead: compile against a small compat header (src/cuda/cuda_runtime_compat.h) that includes the self-contained CUDA wrappers clang ships (builtin variables, device intrinsics and libdevice math) the way clang's own runtime wrapper does, minus the toolkit parts, and adds the few definitions that live in toolkit headers (qualifiers, vector types, warp shuffles, atomicAdd, double min/max). The kernels are emitted as PTX that cuModuleLoadData() JIT compiles for the actual GPU. The toolkit is still needed for libdevice and bin2c; its root is derived from the bin2c binary the build already requires, with -Dcudatoolkit_path as an override for layouts where bin2c lives in a shared bindir. Device code only needs the driver API types, so cuda_helper.cuh includes ffnvcodec/dynlink_cuda.h instead of dynlink_loader.h in the DEVICE_CODE pass; the loader header pulls in nvEncodeAPI.h and with it <windows.h>, which does not compile in a clang device pass. nvcc never sees the compat header and its output is unchanged. Also pass -O3 (nvcc optimizes device code by default, clang does not) and map the nvcc-only -G and -lineinfo flags to their clang equivalents. Tested with clang on Windows 11 (MSYS2 mingw64 clang 22.1.8, libstdc++ 16, CUDA 13.4) and in an Ubuntu 24.04 container (clang 18.1.3, libstdc++ 13, CUDA 12.8), both on an RTX 5090: integer_adm and integer_vif are bit-identical to the nvcc build on every frame on both platforms, and on Windows total kernel time measured with Nsight Systems is within 3% of nvcc.
|
@gedoensmax this finishes the |
|
Did you try to run some 8bit and 10/12 bit footage to check if VMAF scores are correct ? You mention some meson tests, but i think the coverage is subpar. |
|
Only 8-bit before your comment, fair point. I now ran 8, 10 and 12-bit 1080p clips (120 frames each, x265 re-encodes against y4m references) through the clang build on Windows and Linux, the nvcc build on Linux, and the CPU path. integer_adm and integer_vif are bit-identical across all of them at every depth. At 10 and 12-bit the full VMAF score matches too (90.5252 and 88.5416 on every path); at 8-bit only integer_motion differs by a little between runs, and it does that between two runs of the same nvcc binary as well, so it looks like an existing race in the motion extractor rather than anything in this PR. I'll open a separate issue for that. Happy to run any other footage you have in mind. |
|
@kylophone could you take a look at this ? This would enable compilation without any CUDA Toolkit dependencies from what i understand. Basically enabling regular builds with ffmpeg i assume. |
|
Tested on master
Not tested here: MinGW/MSYS2 and Windows, clang 18, CUDA 12.x, 12-bit input. |
|
Thanks @lusoris , that covers the configurations I couldn't. Between the two sets: Linux with clang 22 and CUDA 13.4 (yours), Linux with clang 18 and CUDA 12.8, Windows MSYS2 with clang 22 and CUDA 13.4, and 8, 10 and 12-bit input (mine). integer_adm and integer_vif show zero difference against nvcc in all of them One clarification on scope for review: this removes the need for nvcc, its host compiler and the toolkit headers. The toolkit is still needed at build time for libdevice and bin2c. Nothing changes at runtime. On the integer_motion run-to-run noise I mentioned.. that looks like the accumulator reset race #1583 addresses, with #1612 covering the uninitialized previous-frame blur, so I won't open a separate issue for it |
3 participants
Builds on #1595.
#1436 added -Denable_nvcc=false to compile the kernels with clang, but -nocudainc was left commented out because device intrinsics broke without the toolkit headers. This finishes that: a small header (src/cuda/cuda_runtime_compat.h) includes the self-contained CUDA wrappers that clang ships, the way clang's own runtime wrapper does, and adds the few things that live in toolkit headers (qualifiers, vector types, __shfl_down_sync, atomicAdd, double min/max). The kernels are emitted as PTX and JIT compiled by the driver. The toolkit is still needed for libdevice and bin2c, nothing else. Device and host code both use nv-codec-headers only.
The reason: the toolkit headers don't support a MinGW-w64 host (crt/math_functions.hpp takes a pre-2013 MSVC code path when _WIN32 is defined without _MSC_VER), and nvcc can't be configured there either, so today there is no way to build CUDA support with MSYS2 at all (#1154). With this the normal MinGW toolchain is enough.
Tested on Windows 11 (MSYS2 clang 22, CUDA 13.4) and in an Ubuntu 24.04 container (clang 18, libstdc++ 13, CUDA 12.8 with nvcc for comparison), both on an RTX 5090. integer_adm and integer_vif are bit identical between clang and nvcc on every frame on both platforms, and kernel time is within 3% of nvcc. meson test gives the same results as the nvcc build.
The header only covers what the kernels use today; a new intrinsic would fail the clang build with an undeclared identifier, which is intentional. The toolkit root is taken from the location of bin2c, with -Dcudatoolkit_path as an override.