[FFmpeg-devel] [PATCH 1/5] configure: Add an explicit check and option for nvcc

Philip Langdale philipl at overt.org
Thu Feb 21 05:57:49 EET 2019

The use of nvcc to compile cuda kernels is distinct from the use of
cuda sdk libraries and linking against those libraries. We have
previously not bothered to distinguish these two cases because all
the filters that used cuda kernels also used the sdk. In the following
changes, I'm going to remove the sdk dependency from those filters,
but we need a way to ensure that nvcc is present and functioning, and
also a way to explicitly disable its use so that the filters are not

Note that, unlike the cuda_sdk dependency, using nvcc to compile
a kernel does not cause a build to become non-free. Although nvcc
is distributed with the cuda sdk, and is EULA encumbered, the
compilation process we use does not introduce any EULA covered
code or libraries into the build. In this sense, using nvcc is just
like using any other proprietary compiler like msvc - compiling free
code doesn't suddently make it non-free.

There was previously some confusion on this topic, but the important
distinction is that we use nvcc to generate ptx files - these are
not compiled GPU binaries, but rather an intermediate assembly
representation that is JIT compiled (and I think linked with certain
nvidia library code) when you actually try and run the kernel. nvidia
use this technique to relax machine code compatibility between
hardware generations.

>From here, we can make two observations:
* The ptx files that we include in libavfilter are aggregated rather
  than linked, from the perspective of the (L)GPL
* No proprietary code is included with the ptx files. That code is
  only linked in at the final compilation step at runtime.

Signed-off-by: Philip Langdale <philipl at overt.org>
 configure | 28 ++++++++++++++++++++++++++++
 1 file changed, 28 insertions(+)

diff --git a/configure b/configure
index bf40c1dcb9..2219eb1515 100755
--- a/configure
+++ b/configure
@@ -322,6 +322,7 @@ External library support:
   --disable-amf            disable AMF video encoding code [autodetect]
   --disable-audiotoolbox   disable Apple AudioToolbox code [autodetect]
   --enable-cuda-sdk        enable CUDA features that require the CUDA SDK [no]
+  --disable-cuda-nvcc      disable Nvidia CUDA compiler [autodetect]
   --disable-cuvid          disable Nvidia CUVID support [autodetect]
   --disable-d3d11va        disable Microsoft Direct3D 11 video acceleration code [autodetect]
   --disable-dxva2          disable Microsoft DirectX 9 video acceleration code [autodetect]
@@ -1001,6 +1002,10 @@ hostcc_o(){
     eval printf '%s\\n' $HOSTCC_O
+    eval printf '%s\\n' $NVCC_O
     log test_cc "$@"
     cat > $TMPC
@@ -1022,6 +1027,13 @@ test_objcc(){
     test_cmd $objcc -Werror=missing-prototypes $CPPFLAGS $CFLAGS $OBJCFLAGS "$@" $OBJCC_C $(cc_o $TMPO) $TMPM
+    log test_nvcc "$@"
+    cat > $TMPCU
+    log_file $TMPCU
+    test_cmd $nvcc -ptx $NVCCFLAGS "$@" $NVCC_C $(nvcc_o $TMPO) $TMPCU
     log test_cpp "$@"
     cat > $TMPC
+    cuda_nvcc
@@ -4238,6 +4251,7 @@ tmpfile TMPCPP .cpp
 tmpfile TMPE   $EXESUF
 tmpfile TMPH   .h
 tmpfile TMPM   .m
+tmpfile TMPCU  .cu
 tmpfile TMPO   .o
 tmpfile TMPS   .S
 tmpfile TMPSH  .sh
@@ -6641,6 +6655,20 @@ else
     nvccflags="$nvccflags -m32"
+check_nvcc() {
+    log check_nvcc "$@"
+    disable cuda_nvcc
+    test_nvcc <<EOF && enable cuda_nvcc
+extern "C" {
+    __global__ void hello(unsigned char *data) {}
+if ! disabled cuda_nvcc; then
+    check_nvcc
     check_cflags "$@"
     enabled lto && check_ldflags "$@"

More information about the ffmpeg-devel mailing list