Minimum CUDA version supported by GCC?

Tobias Burnus tobias@codesourcery.com
Thu Nov 19 12:09:00 GMT 2020


On 19.11.20 12:03, Thomas Schwinge wrote:

> As far as I can tell, we (GCC) don't currently state the minimum CUDA
> version supported: for nvptx target, and especially then for
> OpenACC/OpenMP nvptx offloading.

The nvptx target generates ".version 3.1" (= PTX ISA version) code
with ".target sm_30" and ".target sm_35".

If one looks at
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#release-notes__ptx-release-history
that means that we need at least CUDA 5 (which introduced PTX ISA 3.1).

CUDA 5 also added sm_30/sm_35 support - i.e. the sm_30/sm_35 GPUs.

[Hence, without bumping the .version, we cannot support newer sm_XX
and I do not see us adding support for sm_20 or other older hardware.]


In theory, newer CUDA versions still support the oldest GPUs down to sm_10.

There is an (unofficial) list which maps sm_* to GPU names at
https://arnon.dk/matching-sm-architectures-arch-and-gencode-for-various-nvidia-cards/

Contrary to the official CUDA documentation, that page states that
CUDA 10 dropped sm_20 support in CUDA 10.

And Tom's testing showed that CUDA 11.1 no longer supports sm_30
– at least its "ptxas" (which we use in JIT mode at program startup).
That is the reason that GCC 11 now defaults to sm_35.
(Use -misa=sm_30 if you have an older card.)

  * * *

In terms of supported features, moving to PTX ISA 6.3 (= CUDA 10)
makes most sense as that adds a lot of new features, in particular
'.alias'. – There are a few new features between after ISA 3.1 and
before 6.3 but they are not really worthwhile.

(The .alias issue mostly pops up for C++ code with constructors
or in special cases. There are some PRs about this.)

> Do distribution maintainers have an opinion on this maybe, what they need
> to support (for upcoming GCC 11, especially)?  (Matthias, others?)

I think for license reasons, CUDA itself is not shipped by the
distributions but I might be wrong.
In any case, the question is which CUDA versions are installed in
practice and whether requiring an update (e.g. to CUDA 10+) is
sensible and possible.

CUDA 10 was released in September 2018 – thus, it is rather new.

On 19.11.20 12:18, Jakub Jelinek via Gcc wrote:
> I think it would be nice not to kill support for older devices
> unnecessarily.  If there is a strong reason, sure.  And ditto for older
> CUDA versions.
> But I'm afraid many people have older HW and don't have anything newer.
>
> I'm testing offloading only from time to time, but the only remaining hw
> I have has 3.5 capability (Quadro K6000, GK110GL core).  And I use that for
> both upstream development and distro offloading testing.

If I look at the list above, 3.5 are still supported by CUDA 11.1,
hence, CUDA 10's PTX ISA 6.3 including .alias should still work with
sm_30 (w/ CUDA < 11.1) and sm_35.

If we add support for PTX ISA 6.3 (i.e. CUDA >= 10) to use newer
PTX features (both generic and, possibly, those for newer GPUS/sm_XX),
the question is whether we still want to support older PTX ISA/CUDA
or not.

If so, we need an additional -m* flag.

In terms of userfriendliness, adding -misaversion={3.1,6.3} makes
sense – and then adding .alias support (with warning/error with version 3.1)
plus modifying the shfl vs. shfl.sync depending on the version.


IMHO:

Especially due to the .alias issue, I am very much in favor of supporting
.version 6.3 in GCC 11; about adding -misaversion I am a bit unsure, but
it probably makes sense given that CUDA 10 is still relatively new.


Continued: On 19.11.20 12:03, Thomas Schwinge wrote:

> "CUDA" here primarily means libcuda (CUDA Driver API and corresponding
> Nvidia Linux kernel driver).
>
> A related topic then is which Nvidia GPU devices are supported by GCC.
> Due to the generic nature of PTX code generation as well as pretty
> generic hardware access via the CUDA Driver API in libgomp, this is
> primarily dictated by whichever Nvidia GPU devices are supported by
> libcuda, which GCC doesn't influence.
>
> (A related topic also is PR96005 "Add possibility to use newer ptx isa",
> etc., but that's for a separate discussion thread.)

Well, if we do not update the PTX ISA version, we are stuck to PTX ISA 3.1
features and the PTX for sm_30 and sm_35. – Those will work with newer
hardware but prevent supporting the newer hardware better.

Thus, I regard the newer ptx isa question as tightly coupled with the
questions raised in this thread – and do not think that discussing it
in another thread makes sense.

> I can confirm that as of present, things are still working fine with
> many-years-old CUDA 6.5/libcuda 346.46 (2014), with Nvidia Tesla K40c
> GPU, for example.
>
> Until recently (but not anymore), I've been running testing with a Dell
> Precision M4700 laptop's Nvidia Quadro K1000M GPU (GK107 core).  If I
> remember correctly, that one supports Compute Capability 3.0, which Tom
> recently removed (default) support for via PR97348 "[nvptx] Make
> -misa=sm_35 the default".  For avoidance of doubt: that's OK as far as
> I'm concerned; not using this laptop anymore for GCC testing, and, as has
> been noted, 3.5 now is the minimum version with CUDA 11 anyway.
(I wonder whether that's 11.0 or 11.1.)

> (CUDA 11 still supports now, but does deprecate compute capabilities 3.5,
> 3.7, 5.0, by the way.)
> [...](CUDA 8.0 has been released in 2016.)
> [...](2014 CUDA 6.5)

Tobias

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter


More information about the Gcc mailing list