32

Can anyone give me a good explanation as to the nature of CUDA C and C++? As I understand it, CUDA is supposed to be C with NVIDIA's GPU libraries. As of right now CUDA C supports some C++ features but not others.

What is NVIDIA's plan? Are they going to build upon C and add their own libraries (e.g. Thrust vs. STL) that parallel those of C++? Are they eventually going to support all of C++? Is it bad to use C++ headers in a .cu file?

2
  • 1
    I wouldn't guess they will put any effort into supporting C++ features with poor performance on GPU hardware, such as exceptions, since CUDA is designed around high-performance numerical computation.
    – Ben Voigt
    Mar 23, 2012 at 21:54
  • Can we rephrase the title to: "What is the difference between CUDA C/C++ and ISO C/C++"? This is what has been actually answered, and what NVIDIA will or not do can only be speculated upon. Apr 6, 2016 at 19:53

5 Answers 5

28

CUDA C is a programming language with C syntax. Conceptually it is quite different from C.

The problem it is trying to solve is coding multiple (similar) instruction streams for multiple processors.

CUDA offers more than Single Instruction Multiple Data (SIMD) vector processing, but data streams >> instruction streams, or there is much less benefit.

CUDA gives some mechanisms to do that, and hides some of the complexity.

CUDA is not optimised for multiple diverse instruction streams like a multi-core x86. CUDA is not limited to a single instruction stream like x86 vector instructions, or limited to specific data types like x86 vector instructions.

CUDA supports 'loops' which can be executed in parallel. This is its most critical feature. The CUDA system will partition the execution of 'loops', and run the 'loop' body simultaneously across an array of identical processors, while providing some of the illusion of a normal sequential loop (specifically CUDA manages the loop "index"). The developer needs to be aware of the GPU machine structure to write 'loops' effectively, but almost all of the management is handled by the CUDA run-time. The effect is hundreds (or even thousands) of 'loops' complete in the same time as one 'loop'.

CUDA supports what looks like if branches. Only processors running code which match the if test can be active, so a subset of processors will be active for each 'branch' of the if test. As an example this if... else if ... else ..., has three branches. Each processor will execute only one branch, and be 're-synched' ready to move on with the rest of the processors when the if is complete. It may be that some of the branch conditions are not matched by any processor. So there is no need to execute that branch (for that example, three branches is the worst case). Then only one or two branches are executed sequentially, completing the whole if more quickly.

There is no 'magic'. The programmer must be aware that the code will be run on a CUDA device, and write code consciously for it.

CUDA does not take old C/C++ code and auto-magically run the computation across an array of processors. CUDA can compile and run ordinary C and much of C++ sequentially, but there is very little (nothing?) to be gained by that because it will run sequentially, and more slowly than a modern CPU. This means the code in some libraries is not (yet) a good match with CUDA capabilities. A CUDA program could operate on multi-kByte bit-vectors simultaneously. CUDA isn't able to auto-magically convert existing sequential C/C++ library code into something which would do that.

CUDA does provides a relatively straightforward way to write code, using familiar C/C++ syntax, adds a few extra concepts, and generates code which will run across an array of processors. It has the potential to give much more than 10x speedup vs e.g. multi-core x86.

Edit - Plans: I do not work for NVIDIA

For the very best performance CUDA wants information at compile time.

So template mechanisms are the most useful because it gives the developer a way to say things at compile time, which the CUDA compiler could use. As a simple example, if a matrix is defined (instantiated) at compile time to be 2D and 4 x 8, then the CUDA compiler can work with that to organise the program across the processors. If that size is dynamic, and changes while the program is running, it is much harder for the compiler or run-time system to do a very efficient job.

EDIT: CUDA has class and function templates. I apologise if people read this as saying CUDA does not. I agree I was not clear.

I believe the CUDA GPU-side implementation of templates is not complete w.r.t. C++.

User harrism has commented that my answer is misleading. harrism works for NVIDIA, so I will wait for advice. Hopefully this is already clearer.

The hardest stuff to do efficiently across multiple processors is dynamic branching down many alternate paths because that effectively serialises the code; in the worst case only one processor can execute at a time, which wastes the benefit of a GPU. So virtual functions seem to be very hard to do well.

There are some very smart whole-program-analysis tools which can deduce much more type information than the developer might understand. Existing tools might deduce enough to eliminate virtual functions, and hence move analysis of branching to compile time. There are also techniques for instrumenting program execution which feeds directly back into recompilation of programs which might reach better branching decisions.

AFAIK (modulo feedback) the CUDA compiler is not yet state-of-the-art in these areas.

(IMHO it is worth a few days for anyone interested, with a CUDA or OpenCL-capable system, to investigate them, and do some experiments. I also think, for people interested in these areas, it is well worth the effort to experiment with Haskell, and have a look at Data Parallel Haskell)

6
  • 1
    Thanks to everyone for the answers. I'm just starting to learn CUDA and haven't quite grasped the finer points of GPU programming. Mar 23, 2012 at 22:46
  • @xiongtx - I am pleased to help. It took me some time before it became clear. I was a bit saddened that they didn't do a good job of explaining it from the start. Instead they got obsessed about how clever it is.
    – gbulmer
    Mar 23, 2012 at 22:59
  • 3
    This answer is misinformed and incorrect. As examples of where the answer is wrong, CUDA does not automatically parallelize for loops, and it already supports templates (in addition to inheritance, virtual functions, operator overloading, references, new/delete on the device, and many other C++ features).
    – harrism
    Mar 26, 2012 at 2:21
  • 1
    @harrism - I apologise if my poor description of 'templates' is misleading. I was trying to stay at a high level. I'll gladly try to improve it. I did not say CUDA did not have templates, but I can see that I was not clear. I had believed that templates are still restricted (vs C++), for example partial template specialization which would provide ways to cope with the geneneral case, but improve for specific cases. If that is not true, would you please point me at examples? What do you mean "does not automatically parallelize" I didn't write that. CUDA can't do it, so I'd like to fix my words.
    – gbulmer
    Mar 26, 2012 at 14:52
  • 1
    I again apologise, I wrote 'loop' not loops deliberatly to get across the idea of taking a for loop, and transforming it. I thought the word kernel wouldn't help, but I can try again. I'd rather have a clear answer. The OP asked "Can anyone give me a good explanation as to the nature of CUDA C and C++" and I specifically labelled the first part C in order to separate the idea of kernels and control flow which can be (is) done using 'vanilla' C in CUDA (except the <<<...>>> of course). As for CUDA template support vs C++, is all of C++ (not C++11) template support in CUDA?
    – gbulmer
    Mar 27, 2012 at 8:15
28

CUDA is a platform (architecture, programming model, assembly virtual machine, compilation tools, etc.), not just a single programming language. CUDA C is just one of a number of language systems built on this platform (CUDA C, C++, CUDA Fortran, PyCUDA, are others.)

CUDA C++

Currently CUDA C++ supports the subset of C++ described in Appendix D ("C/C++ Language Support") of the CUDA C Programming Guide.

To name a few:

  • Classes
  • __device__ member functions (including constructors and destructors)
  • Inheritance / derived classes
  • virtual functions
  • class and function templates
  • operators and overloading
  • functor classes

Edit: As of CUDA 7.0, CUDA C++ includes support for most language features of the C++11 standard in __device__ code (code that runs on the GPU), including auto, lambda expressions, range-based for loops, initializer lists, static assert, and more.

Examples and specific limitations are also detailed in the same appendix linked above. As a very mature example of C++ usage with CUDA, I recommend checking out Thrust.

Future Plans

(Disclosure: I work for NVIDIA.)

I can't be explicit about future releases and timing, but I can illustrate the trend that almost every release of CUDA has added additional language features to get CUDA C++ support to its current (In my opinion very useful) state. We plan to continue this trend in improving support for C++, but naturally we prioritize features that are useful and performant on a massively parallel computational architecture (GPU).

3
  • 1
    Hello @harrism, would it be possible to update this answer to reflect the current state of C++ in CUDA and future plans after 2015?
    – Bar
    Sep 26, 2015 at 12:49
  • Are CUDA Fortran and PyCUDA actually kernel versions of Fortran and Python that compile to run on the GPU? Or just host APIs? Apr 6, 2016 at 19:56
  • CUDA Fortran is a Fortran compiler with CUDA extensions, along with a host API. PyCUDA is more of a host API and convenience utilities, but kernels still have to be written in CUDA C++. "CUDA Python", part of Numba, is a compiler for CUDA-annotated Python to run on GPUs.
    – harrism
    Apr 16, 2016 at 5:55
10

Not realized by many, CUDA is actually two new programming languages, both derived from C++. One is for writing code that runs on GPUs and is a subset of C++. Its function is similar to HLSL (DirectX) or Cg (OpenGL) but with more features and compatibility with C++. Various GPGPU/SIMT/performance-related concerns apply to it that I need not mention. The other is the so-called "Runtime API," which is hardly an "API" in the traditional sense. The Runtime API is used to write code that runs on the host CPU. It is a superset of C++ and makes it much easier to link to and launch GPU code. It requires the NVCC pre-compiler which then calls the platform's C++ compiler. By contrast, the Driver API (and OpenCL) is a pure, standard C library, and is much more verbose to use (while offering few additional features).

Creating a new host-side programming language was a bold move on NVIDIA's part. It makes getting started with CUDA easier and writing code more elegant. However, truly brilliant was not marketing it as a new language.

2

Sometimes you hear that CUDA would be C and C++, but I don't think it is, for the simple reason that this impossible. To cite from their programming guide:

For the host code, nvcc supports whatever part of the C++ ISO/IEC 14882:2003 specification the host c++ compiler supports.

For the device code, nvcc supports the features illustrated in Section D.1 with some restrictions described in Section D.2; it does not support run time type information (RTTI), exception handling, and the C++ Standard Library.

As I can see, it only refers to C++, and only supports C where this happens to be in the intersection of C and C++. So better think of it as C++ with extensions for the device part rather than C. That avoids you a lot of headaches if you are used to C.

0
1

What is NVIDIA's plan?

I believe the general trend is that CUDA and OpenCL are regarded as too low level techniques for many applications. Right now, Nvidia is investing heavily into OpenACC which could roughly be described as OpenMP for GPUs. It follows a declarative approach and tackles the problem of GPU parallelization at a much higher level. So that is my totally subjective impression of what Nvidia's plan is.

11
  • It would be cool to have a functional interface to GPU programming, with an intelligent library behind it farming off the computations to the GPU.
    – quant_dev
    Mar 23, 2012 at 21:40
  • @quant_dev: It would have to be very intelligent, as it would need to be able to identify sufficiently large volumes of contiguous compute to amortize away the overhead, and to hide the latency. Mar 23, 2012 at 22:22
  • Why? The task is not that much different from writing a BLAS implementation.
    – quant_dev
    Mar 23, 2012 at 22:27
  • @quant_dev: With an accelerated BLAS implementation, you (the programmer) typically has to explicitly invoke the accelerated routine (or indirectly via something like Thrust). The runtime doesn't (AFAIK) make decisions about whether it would be more beneficial to execute each workitem on the CPU or the GPU. Mar 23, 2012 at 22:43
  • 2
    @quant_dev I work at AccelerEyes and we do exactly that. Check out ArrayFire (www.accelereyes.com/arrayfire) The base functions are free to use. You can contact me through my email address (on my profile) if you need more information. Mar 23, 2012 at 22:47

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge that you have read and understand our privacy policy and code of conduct.

Not the answer you're looking for? Browse other questions tagged or ask your own question.