CUDA tutorial: Difference between revisions

= A simple CUDA C program= <!--T:10-->

The following example shows how to add two numbers on the GPU using CUDA. Note that this is just an exercise, it's very simple, so don't expect to see any actual acceleration.

<syntaxhighlight lang="cpp" line highlight="1,5">

___global__ void add (int *a, int *b, int *c){

}

</syntaxhighlight>

<!--T:17-->

That code does not look parallel!

Solution: Let's look at what's inside the triple brackets in the kernel call and make some changes :

<syntaxhighlight lang="cpp" line highlight="1,5">

add <<< N, 1 >>> (dev_a, dev_b, dev_c);

</syntaxhighlight>

Here we replaced 1 by N, so that N different CUDA blocks will be executed at the same time. However, in order to achieve parallelism we need to make some changes to the kernel as well:

<syntaxhighlight lang="cpp" line highlight="1,5">

__global__  void add (int *a, int *b, int *c){

   c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];

</syntaxhighlight>

where blockIdx.x is the unique number identifying a CUDA block. This way each CUDA block adds a value from a[ ] to b[ ].

[[File:Cuda-blocks-parallel.png|thumbnail|CUDA blocks-based parallelism. ]]

<!--T:19-->

Can we again make some modifications in those triple brackets ?  

<syntaxhighlight lang="cpp" line highlight="1,5">

add <<< 1, '''N''' >>> (dev_a, dev_b, dev_c);

</syntaxhighlight>

Now instead of blocks, the job is distributed across parallel threads. What is the advantage of having parallel threads ? Unlike blocks, threads can communicate between each other: in other words, we parallelize across multiple threads in the block when heavy communication is involved. The chunks of code that can run independently, i.e. with little or no communication, are distributed across parallel blocks.

= Advantages of shared memory= <!--T:20-->

So far all the memory transfers in the kernel have been done via the regular GPU (global) memory which is relatively slow. Often we have so many communications between the threads that the performance decreases significantly.  In order to address this issue there exists another type of memory called '''shared memory''' which can be used to speed-up the memory operations between the threads. However the trick is that only the threads within a block can communicate.  In order to demonstrate the usage of such shared memory we consider the dot product example where two vectors are multiplied together element by element and then summed. Below is the kernel:

__global__  void dot(int *a, int *b, int *c){

   int temp = a[threadIdx.x]*b[threadIdx.x];  

}

</syntaxhighlight>

After each thread computes its portion, we need to add everything together: each thread has to share its data.  However, the problem is that each copy of thread's temp variable is private. This can be resolved by the use of shared memory. Below is the kernel with the modifications to use shared memory:

<syntaxhighlight lang="cpp" line highlight="1,4">

#define N 512

}

</syntaxhighlight>

= Basic performance considerations = <!--T:26-->