| Message ID | 20190221035753.27525-2-philipl@overt.org |
|---|---|
| State | New |
| Headers | show |
On 21.02.2019 04:57, Philip Langdale wrote: > 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 > built. > > 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. > From a technical standpoint, this whole series looks fine to me. However, it really does not fell non-nonfree to me that the only way to build these filters remains to register with nvidia, accept their various EULAs and download their SDK for nvcc and the libs around it. Even moving them out of ffmpeg and loading them at runtime (which would be a major ABI, API and usability break) does not solve that. You still need to either get said SDK and build them yourself, or get someone else to do it an run into the same issue with redistributability there. If the general majority agrees that this series way of doing things is conforming, it is is good to go. But I myself won't make that call. I'd still like to get this merged asap, so if in doubt, keep nvcc in the nonfree list for now, and discuss its removal from there in a separate patch.
Hello, On Tue, 26 Feb 2019, at 14:51, Timo Rothenpieler wrote: > > 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. But there are alternatives to MSVC to compile C or C++ code. Are there for nvcc? > However, it really does not fell non-nonfree to me that the only way to > build these filters remains to register with nvidia, accept their > various EULAs and download their SDK for nvcc and the libs around it. > > Even moving them out of ffmpeg and loading them at runtime (which would > be a major ABI, API and usability break) does not solve that. > You still need to either get said SDK and build them yourself, or get > someone else to do it an run into the same issue with redistributability > there. I agree with you here. This feels as non-free to me too. GPLv3 says: 'The “Corresponding Source” for a work in object code form means all the source code needed to generate, install, and (for an executable work) run the object code and to modify the work, including scripts to control those activities. However, it does not include the work's System Libraries, or general-purpose tools or generally available free programs which are used unmodified in performing those activities but which are not part of the work.' I don't think nvcc fit the "System Libraries" or the "general-purpose tools" or "generally available free programs". GPLv2 says: ' For an executable work, complete source code means all the source code for all modules it contains, plus any associated interface definition files, plus the scripts used to control compilation and installation of the executable. 'However, as a special exception, the source code distributed need not include anything that is normally distributed (in either source or binary form) with the major components (compiler, kernel, and so on) of the operating system on which the executable runs, unless that component itself accompanies the executable.' I don't think nvcc fit the"normally distributed with the operating system". LGPLv2 and v3 follow GPLv2 and v3, respectively. -- Jean-Baptiste Kempf - President +33 672 704 734
> From a technical standpoint, this whole series looks fine to me. > > However, it really does not fell non-nonfree to me that the only way to > build these filters remains to register with nvidia, accept their > various EULAs and download their SDK for nvcc and the libs around it. > > Even moving them out of ffmpeg and loading them at runtime (which would > be a major ABI, API and usability break) does not solve that. > You still need to either get said SDK and build them yourself, or get > someone else to do it an run into the same issue with redistributability > there. I would suggest to have both options: 1. Compiling it as part of ffmpeg with option non-free 2. Compiling without non-free option and loading the kernels from a file For the latter case: It's an external and separate executable. It is not a violation of the GPL if ffmpeg checks for existence and executes that in a separate process. When you need to distribute pre-compiled ffmpeg versions, then it's not a matter of usability. It's rather about whether there's a way to use it or not. :-) --- softworkz
> From: Jean-Baptiste Kempf > Sent: Dienstag, 26. Februar 2019 15:01 > > [...] > > I don't think nvcc fit the"normally distributed with the operating system". I'm not sure if the role of nvcc has been fully understood. nvcc is some kind of 'pre-compiler' which is not distributed or linked to. Distributing GPL code that was compiled with MSVC doesn't require the MSVC compiler to be open source as well. --- softworkz
On Tue, 26 Feb 2019, at 21:26, Soft Works wrote: > It's an external and separate executable. It is not a violation of > the GPL if ffmpeg checks for existence and executes that in a separate > process. Believing that executing in a separate process removes the obligation of the GPL is an very-much-shared but totally-wrong idea.
On Tue, 26 Feb 2019, at 21:36, Soft Works wrote: > I'm not sure if the role of nvcc has been fully understood. > nvcc is some kind of 'pre-compiler' which is not distributed or linked to. > > Distributing GPL code that was compiled with MSVC doesn't require the > MSVC compiler to be open source as well. Unfortunately, this is not clear at all, and has been debated quite a few time, and most projects have a different understanding of this issue. Especially since MSVC is not shipped with Windows.
> From: Jean-Baptiste Kempf > Sent: Dienstag, 26. Februar 2019 23:07 > To: ffmpeg-devel@ffmpeg.org > Subject: Re: [FFmpeg-devel] [PATCH 1/5] configure: Add an explicit check and option for nvcc > > On Tue, 26 Feb 2019, at 21:26, Soft Works wrote: > > It's an external and separate executable. It is not a violation of > > the GPL if ffmpeg checks for existence and executes that in a separate > > process. > > Believing that executing in a separate process removes the obligation of the GPL is an very-much-shared but totally-wrong idea. Even if you're right - this doesn't apply here because in this case we're talking about a binary that is NOT distributed as part of the software. --- softworkz
On 26.02.2019 21:36, Soft Works wrote: > >> From: Jean-Baptiste Kempf >> Sent: Dienstag, 26. Februar 2019 15:01 >> >> [...] >> >> I don't think nvcc fit the"normally distributed with the operating system". > > I'm not sure if the role of nvcc has been fully understood. > nvcc is some kind of 'pre-compiler' which is not distributed or linked to. > > Distributing GPL code that was compiled with MSVC doesn't require the > MSVC compiler to be open source as well. Yes, but there are alternative OSS compilers that implement the same programming language (at least to a very large extent). This doesn't seem to be the case here. So even if this scenario is not explicitly excluded in the GPL, if doesn't feel very free to me, too. Regards, Tobias
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 } +nvcc_o(){ + eval printf '%s\\n' $NVCC_O +} + test_cc(){ 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 } +test_nvcc(){ + log test_nvcc "$@" + cat > $TMPCU + log_file $TMPCU + test_cmd $nvcc -ptx $NVCCFLAGS "$@" $NVCC_C $(nvcc_o $TMPO) $TMPCU +} + test_cpp(){ log test_cpp "$@" cat > $TMPC @@ -1786,6 +1798,7 @@ HWACCEL_AUTODETECT_LIBRARY_LIST=" audiotoolbox crystalhd cuda + cuda_nvcc cuvid d3d11va dxva2 @@ -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" fi +check_nvcc() { + log check_nvcc "$@" + disable cuda_nvcc + test_nvcc <<EOF && enable cuda_nvcc +extern "C" { + __global__ void hello(unsigned char *data) {} +} +EOF +} + +if ! disabled cuda_nvcc; then + check_nvcc +fi + check_optflags(){ check_cflags "$@" enabled lto && check_ldflags "$@"
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 built. 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@overt.org> --- configure | 28 ++++++++++++++++++++++++++++ 1 file changed, 28 insertions(+)