[{"id":"ch01","contentType":"text","title":"Introduction to CUDA Programming","content":"

Introduction to CUDA Programming

Since its first release in 2007, Compute Unified Device Architecture (CUDA) has grown to become the de facto standard when it comes to using Graphic Computing Units (GPUs) for general-purpose computation, that is, non-graphics applications. So, what exactly is CUDA? Someone might ask the following:

Is it a programming language?
Is it a compiler?
Is it a new computing paradigm?

In this chapter, we will demystify some of the myths around GPU and CUDA. This chapter lays the foundation for heterogeneous computing by providing a simplified view of High-Performance Computing (HPC) history and substantiating it with laws such as Moore's Law and Dennard Scaling, which were—and still are—driving the semiconductor industry and hence the processor architecture itself. You will also be introduced to the CUDA programming model and get to know the fundamental difference between CPU and GPU architecture. By the end of this chapter, you will be able to write and understand Hello World! programs using CUDA programming constructs in the C language.

While this chapter primarily uses C to demonstrate CUDA constructs, we will be covering other programming languages such as Python, Fortran, and OpenACC in other chapters.

The following topics will be covered in this chapter:

The history of high-performance computing
Hello World from CUDA
Vector addition using CUDA
Error reporting with CUDA
Data type support in CUDA

","url":"https://subscription.packtpub.com/book/programming/9781788996242/1","metaDescription":"","titleSlug":"introduction-to-cuda-programming"},{"id":"ch01lvl1sec01","contentType":"text","title":"The history of high-performance computing","content":"

The history of high-performance computing

HPC has always pushed the limits in order to deliver scientific discoveries. The fundamental shift in processor architecture and design has helped to cross FLOP barriers, starting from Mega-Floating Point Operations (MFLOPs) to now being able to do PetaFLOP calculation in a second.

Floating-Point Operations (FLOPs) per second is the fundamental unit for measuring the theoretical peak of any compute processor. MegaFLOP stands for 10 to the 6^th power of FLOPS. PetaFLOP stands for 10 to the 15^th power of FLOPS.
Instruction-Level Parallelism (ILP) is a concept wherein code-independent instructions can execute at the same time. For the instructions to execute in parallel, they need to be independent of each other. All modern CPU architecture (even GPU architecture) provides five to 15+ stages to allow for faster clock rates:

Instr 1: add = inp1 + inp2
Instr 2: mult = inp1 * inp2
Instr 3: final_result = mult / add

Operations for calculating the mult and add variables do not depend on each other, so they can be calculated simultaneously while calculating final_result, which depends on the results of the Instr 1 and Instr 2 operations. Therefore, it cannot be calculated until add and mult have been calculated.

When we look at the history of HPC in terms of technology changes, which resulted in a fundamental shift in designing new processors and its impact on the scientific community, there are three primary ones that stand out and can be referred to as epochs:

Epoch 1: The history of the supercomputer goes back to CRAY-1, which was basically a single vector CPU architecture providing peak 160 MegaFLOP/MFLOP compute power.
Epoch 2: The MegaFLOP barrier was crossed by moving from single-core design to multi-core design in CRAY-2, which was a 4 Core Vector CPU that gave 2 GigaFLOPs of peak performance.
Epoch 3: Crossing GigaFLOP compute performance was a fundamental shift and required compute nodes to work with each other and communicate by a network to deliver higher performance. Cray T3D was one of the first machines that delivered 1 TeraFLOP of compute performance. The network was 3D Torus and provided a bandwidth of 300 MB/s. It was the first significant implementation of a rich shell around a standard microprocessor.

After this, for almost 20 years, there were no fundamental innovations. Technological innovations were primarily focused on three architectural innovations:

Moving from an 8-bit to a 16-bit to a 32-bit and now a 64-bit instruction set
Increasing ILP
Increasing the number of cores

This was supported by increasing the clock rate, which currently stands at 4 GHz. It was possible to deliver this because of the fundamental laws that drove the semiconductor industry.

Moore's Law: This law observes the number of transistors in a dense integrated circuit double every two years.

Moore's prediction proved accurate for several decades and still does. Moore's Law is an observation and projection of a historical trend.

Dennard scaling: This a scaling law that keeps Moore's Law alive. Dennard made an observation with respect to the relationship between transistor size and power density and summarized it in the following formula:

P = QfCV² + V I_leakage

In this equation, Q is the number of transistors, f is the operating frequency, C is the capacitance, V is the operating voltage, and I_leakage is the leakage current.

Dennard scaling and Moore's Law are related to each other as it's inferred that reducing the size of transistors can lead to more and more transistors per chip in terms of cost-effectiveness.

With Dennard scaling rules, the total chip power for a given size stayed the same for many processor generations. Transistor count doubled while size kept shrinking (1/S rate) and increased in frequency by 40% every two years. This stopped after the feature size reached below 65 nm as these rules could no longer be sustained due to the leakage current growing exponentially. To reduce the effect of leakage current, new innovations were enforced on the switching process. However, these breakthroughs still were not sufficient to revive how voltage was scaled. The voltage remained constant at 1 V for many processor designs. It was no longer possible to keep the power envelope constant. This is also popularly known as Powerwall.

Dennard scaling held its own from 1977 until 1997 and then began to fade. Due to this, from 2007 to 2017, processors went from 45 nm to 16 nm but resulted in a threefold increase in energy/chip size.

At the same time, the pipeline stages went from five stages to 15+ in the latest architecture. To keep the instruction pipeline full, advance techniques such as speculation were used. The speculation unit involves predicting the program's behavior, such as predicting branches and memory addresses. If a prediction is accurate, it can proceed; otherwise, it undoes the work it did and restarts. Deep pipeline stages and the way legacy software is written resulted in unused transistors and wasted clock cycles, which means that there was no improvement in terms of performance for the application.

Then came GPU, which was primarily used for graphics processing. A researcher named Mark Harris made use of GPU for non-graphics tasks for the first time, and the new term General Purpose Computation using GPU (GPGPU) was coined. GPU was proven to be efficient when it came to certain tasks that fell into the category of data parallelism. Unsurprisingly, most of the compute-intensive tasks in many HPC applications are data-parallel in nature. They were mostly matrix to matrix multiplications, which is a routine in the Basic Linear Algebra Specification (BLAS) and used extensively.

The only problem for users when it came to adapting and using GPU was that they had to understand the graphics pipeline to make use of GPU. The only interface that was provided for any computation work on GPU centered around shader execution. There was a need to provide a more general interface that was known to developers who were working in the HPC community. This was solved by the introduction of CUDA in 2007.

While the GPU architecture is also bound by the same laws (Moore's Law and Dennard scaling), the design of processors takes a different approach and dedicates transistors for different usage and achieves higher performance than traditional homogeneous architectures.

The following diagram shows the evolution of computer architecture from sequential processing to distributed memory and its impact on programming models:

With GPU being added to existing servers, there are two types of processors (CPU and GPU) on which the application runs, which brings in a notion of heterogeneity. This is what we will introduce in the next section.

Heterogeneous computing

The common misconception around GPU is that it is an alternative to CPU. GPUs are used to accelerate the parts of the code that are parallel in nature. Accelerator is a common term that's used for GPUs because they accelerate an application by running the parallel part of the code faster, while CPUs run the other part of the code, which is latency bound. Hence, a highly efficient CPU coupled with a high throughput GPU results in improved performance for the application.

The following diagram represents an application running on multiple processor types:

This concept can be very well defined with the help of Amdahl's law. Amdahl's law is used to define the maximum speedup that can be achieved when only a fraction of the application is parallelized. To demonstrate this, the preceding diagram shows two parts of the code. One part is latency bound, while the other is throughput bound. We will cover what these two terms mean in the next section, which differentiates between the CPU and GPU architecture.

The key point is that CPU is good for a certain fraction of code that is latency bound, while GPU is good at running the Single Instruction Multiple Data (SIMD) part of the code in parallel. If only one of them, that is, CPU code or GPU code, runs faster after optimization, this won't necessarily result in good speedup for the overall application. It is required that both of the processors, when used optimally, give maximum benefit in terms of performance. This approach of essentially offloading certain types of operations from the processor onto a GPU is called heterogeneous computing.

The following diagram depicts the two types of sections that all applications have, that is, latency bound and throughput bound:

Here, the importance of improving both sections is demonstrated using Amdahl's law.

Programming paradigm

The classification of computer architecture was done using Flynn's taxonomy, which describes four classes of architecture. One of Flynn's classification SIMDs is used to describe GPU architecture. However, there is a subtle difference between the two. SIMD is used to describe an architecture where the same instruction is applied in parallel to multiple data points. This description is suitable for processors that have the capability of doing vectorization. In contrast, in Single Instruction Multiple Threads (SIMTs), rather than a single thread issuing the instructions, multiple threads issue the same instruction to different data. The GPU architecture is more suitable in terms of the SIMT category compared to SIMD.

Let's look at an example of adding two arrays and storing data in a third array. The dataset for this operation consists of the arrays A, B, and C. The same operations that are used for addition are used on each element of the array:

Cx = Ax + Bx

It is obvious that each task is independent of each other, but the same operation is being applied by all of the threads.

The following screenshot shows vector addition, depicting an example of this paradigm:

Low latency versus higher throughput

As we mentioned in the previous section, CPU architecture is optimized for low latency access while GPU architecture is optimized for data parallel throughput computation. As shown in the following screenshot, the CPU architecture has a large amount of cache compared to GPU and has many types. The higher we go, that is, L3 to L1, the lower the amount of cache is present, but less latency. The CPU architecture is designed for low latency access to cached datasets. A large number of transistors are used to implement the speculative execution and out of order execution. Since CPUs run at a very high clock speed, it becomes necessary to hide the latency of fetching the data by frequently storing used data in caches and predicting the next instruction to execute. Applications that can explore this temporal locality can optimally make use of a CPU cache. Also, applications where it is easy to fill the instruction pipeline, for example, an application with no if and else statements in its code, can benefit from this by hiding the latency of fetching the instruction. Hence, the CPU architecture is a latency reducing architecture.

The following screenshot shows how the CPU and GPU architecture dedicate the chip die area for different memory and compute units. While GPU uses a lot of transistors for computing ALUs, CPU uses it to reduce latency:

The GPU architecture, on the other hand, is called a latency reducing or high throughput architecture. The GPU architecture hides latency with computations from other threads. When one thread is waiting for the data to be available for computation, the other threads can start execution and hence not waste any clock cycles. If you are familiar with CUDA, then you might know about the concept of warps. We will cover the concept of warps in the upcoming chapters. (In CUDA, the execution unit is a warp and not a thread. Due to this, context switching happens between warps and not threads).

Some of you might be already wondering why we can't create these threads in the CPU and do the same thing to hide latency. The reason for this is that GPUs have lots of registers, and all of the thread context switching information is already present in them. This is the fastest memory that's available. However, in CPU, there are limited sets of registers and hence thread-related information is usually stored in a lower memory hierarchy such as a cache. For example, Volta contains 20 MB of register storage. Due to this, the context switching time between threads in CPU, compared to GPU, is much higher.

Now, let's take a look at the different approaches when it comes to programming on GPU.

Programming approaches to GPU

Let's go back to our original question, that is, what is CUDA? CUDA is a parallel computing platform and programming model architecture developed by NVIDIA that exposes general-purpose computations on GPU as first-class capabilities. Like any other processor, the GPU architecture can be coded using various methods. The easiest method, which provides drop-in acceleration, is making use of existing libraries. Alternatively, developers can choose to make use of OpenACC directives for quick acceleration results and portability. Another option is to choose to dive into CUDA by making use of language constructs in C, C++, Fortran, Python, and more for the highest performance and flexibility. We will be covering all of these methods in detail in the subsequent chapters.

The following screenshot represents the various ways we can perform GPU programming:

In this section, we provided you with a perspective of how processors and high-performance computing have evolved over time. We provided you with an overview of why the heterogeneous programming model is key to getting the best performance from an application, followed by approaches to GPU programming. In the next section, we will start writing a Hello World program on a GPU.

","url":"https://subscription.packtpub.com/book/programming/9781788996242/1/ch01lvl1sec01/the-history-of-high-performance-computing","metaDescription":"","titleSlug":"the-history-of-high-performance-computing"},{"id":"ch01lvl1sec02","contentType":"text","title":"Technical requirements ","content":"

Technical requirements

A Linux/Windows PC with a modern NVIDIA GPU (Pascal architecture onwards) is required for this chapter, along with all of the necessary GPU drivers and the CUDA Toolkit (10.0 onward) installed. If you're unsure of your GPU's architecture, please visit NVIDIA's GPU site (https://developer.nvidia.com/cuda-gpus) and confirm your GPU's architecture. This chapter's code is also available on GitHub at https://github.com/PacktPublishing/Learn-CUDA-Programming.

The code examples in this chapter have been developed and tested with version 10.1 of CUDA Toolkit, but it is recommended to use the latest CUDA version, if possible.

","url":"https://subscription.packtpub.com/book/programming/9781788996242/1/ch01lvl1sec02/technical-requirements","metaDescription":"","titleSlug":"technical-requirements"},{"id":"ch01lvl1sec03","contentType":"text","title":"Hello World from CUDA","content":"

Hello World from CUDA

CUDA is a heterogeneous programming model that includes provisions for both CPU and GPU. The CUDA C/C++ programming interface consists of C language extensions so that you can target portions of source code for parallel execution on the device (GPU). It is based on industry-standard C/C++ and provides a library of C functions that can be executed on the host (CPU) so that it can interact with the device.

In CUDA, there are two processors that work with each other. The host is usually referred to as the CPU, while the device is usually referred to as the GPU. The host is responsible for calling the device functions. As we've already mentioned, part of the code that runs on the GPU is called device code, while the serial code that runs on the CPU is called host code.

Let's start by writing our first CUDA code in C. The intention is to take a systematic step-wise approach, start with some sequential code, and convert it into CUDA-aware code by adding some additional keywords. As we mentioned earlier, there is no necessity to learn a new language—all we need to do is add some keywords to the existing language so that we can run it in a heterogeneous environment with CPU and GPU.

Let's take a look at our first piece of code. All this code does is print Hello World! from both the host and device:

#include<stdio.h>
#include<stdlib.h>

__global__ void print_from_gpu(void) {
    printf("Hello World! from thread [%d,%d] \
        From device\n", threadIdx.x,blockIdx.x);
}

int main(void) {
    printf("Hello World from host!\n");
    print_from_gpu<<<1,1>>>();
    cudaDeviceSynchronize();
    return 0;
}

Let's try to compile and run the preceding snippet:

Compile the code: Place the preceding code into a file called hello_world.cu and compile it using the NVIDIA C Compiler (nvcc). Note that the extension of the file is .cu, which tells the compiler that this file has GPU code inside it:

$ nvcc -o hello_world hello_world.cu

Execute the GPU code: We should receive the following output after executing the GPU code:

By now, you might have already observed that the CUDA C code isn't used very differently and only requires that we learn some additional constructs to tell the compiler which function is GPU code and how to call a GPU function. It isn't like we need to learn a new language altogether.

In the preceding code, we added a few constructs and keywords, as follows:

__global__: This keyword, when added before the function, tells the compiler that this is a function that will run on the device and not on the host. However, note that it is called by the host. Another important thing to note here is that the return type of the device function is always "void". Data-parallel portions of an algorithm are executed on the device as kernels.
<<<,>>>: This keyword tells the compiler that this is a call to the device function and not the host function. Additionally, the 1,1 parameter basically dictates the number of threads to launch in the kernel. We will cover the parameters inside angle brackets later. For now, the 1,1 parameter basically means we are launching the kernel with only one thread, that is, sequential code with a thread since we are not doing anything important in the code apart from printing.
threadIdx.x, blockIdx.x: This is a unique ID that's given to all threads. We will cover this topic more in the next section.
cudaDeviceSynchronize(): All of the kernel calls in CUDA are asynchronous in nature. The host becomes free after calling the kernel and starts executing the next instruction afterward. This should come as no big surprise since this is a heterogeneous environment and hence both the host and device can run in parallel to make use of the types of processors that are available. In case the host needs to wait for the device to finish, APIs have been provided as part of CUDA programming that make the host code wait for the device function to finish. One such API is cudaDeviceSynchronize, which waits until all of the previous calls to the device have finished.

Try removing the cudaDeviceSynchronize() call and see whether the device output is visible or not. Alternatively, try putting this call before printing it on the host code.

Thread hierarchy

Now, let's start playing around with the two parameters, that is, threadIdx.x and blockIdx.x.

Experiment 1: First, change the parameter from <<<1,1>>> to <<<2,1>> and view the output. The output of running multiple thread-single blocks of Hello World code should be as follows:

As we can see, instead of one thread, we now have two threads printing the value. Note that their unique IDs are different.

Experiment 2: Now, instead of changing the first parameter, let's change the second, that is, change <<<1,1>>> to <<<1,2>>> and observe the output of running multiple single-thread blocks of Hello World code, as follows:

As you can see, the total number of threads that were launched into the kernel is two, just like before—the only difference is that their IDs are different. So, what are these thread and block concepts? To combat this, let's dive into the GPU architecture some more.

GPU architecture

One of the key reasons why CUDA became so popular is because the hardware and software have been designed and tightly bound to get the best performance out of the application. Due to this, it becomes necessary to show the relationship between the software CUDA programming concepts and the hardware design itself.

The following screenshot shows the two sides of CUDA:

We can see that the CUDA software has been mapped to the GPU hardware.

The following table, in accordance with the preceding screenshot, explains software and hardware mapping in terms of the CUDA programming model:

Software	Executes on/as	Hardware
CUDA thread		CUDA Core/SIMD code
CUDA block		Streaming multiprocessor
GRID/kernel		GPU device

Let's take a look at the preceding table's components in detail:

CUDA Threads: CUDA threads execute on a CUDA core. CUDA threads are different from CPU threads. CUDA threads are extremely lightweight and provide fast context switching. The reason for fast context switching is due to the availability of a large register size in a GPU and hardware-based scheduler. The thread context is present in registers compared to CPU, where the thread handle resides in a lower memory hierarchy such as a cache. Hence, when one thread is idle/waiting, another thread that is ready can start executing with almost no delay. Each CUDA thread must execute the same kernel and work independently on different data (SIMT).
CUDA blocks: CUDA threads are grouped together into a logical entity called a CUDA block. CUDA blocks execute on a single Streaming Multiprocessor (SM). One block runs on a single SM, that is, all of the threads within one block can only execute on cores in one SM and do not execute on the cores of other SMs. Each GPU may have one or more SM and hence to effectively make use of the whole GPU; the user needs to divide the parallel computation into blocks and threads.
GRID/kernel: CUDA blocks are grouped together into a logical entity called a CUDA GRID. A CUDA GRID is then executed on the device.

This may sound somewhat complicated at first glance. In this next section, we'll take a look at an example of vector addition to explain this. Hopefully, things will become much clearer.

","url":"https://subscription.packtpub.com/book/programming/9781788996242/1/ch01lvl1sec03/hello-world-from-cuda","metaDescription":"","titleSlug":"hello-world-from-cuda"},{"id":"ch01lvl1sec04","contentType":"text","title":"Vector addition using CUDA","content":"

Vector addition using CUDA

The problem that we are trying to solve is vector addition. As we are aware, vector addition is a data parallel operation. Our dataset consists of three arrays: A, B, and C. The same operation is performed on each element:

Cx = Ax + Bx

Each addition is independent of each other, but the same operation is applied by all CUDA threads. To get started, configure your environment according to the following steps:

Prepare your GPU application. This code will be placed in 01_cuda_introduction/01_vector_addition.
Compile your application with the nvcc compiler with the following command:

$nvcc -o vector_addition vector_addition.cu

The preceding code is sequential code. We will convert this code so that it can run on a GPU using a step-by-step approach, as follows:

#include<stdio.h>
#include<stdlib.h>

#define N 512

void host_add(int *a, int *b, int *c) {
    for(int idx=0;idx<N;idx++)
        c[idx] = a[idx] + b[idx];
}

//basically just fills the array with index.
void fill_array(int *data) {
    for(int idx=0;idx<N;idx++)
        data[idx] = idx;
}

void print_output(int *a, int *b, int*c) {
    for(int idx=0;idx<N;idx++)
        printf("\n %d + %d = %d", a[idx] , b[idx], c[idx]);
}

int main(void) {
    int *a, *b, *c;
    int size = N * sizeof(int);
   // Alloc space for host copies of a, b, c and setup input values
    a = (int *)malloc(size); fill_array(a);
    b = (int *)malloc(size); fill_array(b);
    c = (int *)malloc(size);
    host_add(a,b,c);
    print_output(a,b,c);
    free(a); free(b); free(c);
    return 0;
}

Before converting the sequential code, let's take a look at the fundamental changes or steps that are taken between the CUDA and sequential code:

Sequential code		CUDA code
Step 1	Allocate memory on the CPU, that is, `malloc new`.	Step 1	Allocate memory on the CPU, that is, `malloc new`.
Step 2	Populate/initialize the CPU data.	Step 2	Allocate memory on the GPU, that is, `cudaMalloc`.
Step 3	Call the CPU function that has the crunching of data. The actual algorithm is vector addition in this case.	Step 3	Populate/initialize the CPU data.
Step 4	Consume the crunched data, which is printed in this case.	Step 4	Transfer the data from the host to the device with `cudaMemcpy`.
		Step 5	Call the GPU function with `<<<,>>>` brackets.
		Step 6	Synchronize the device and host with `cudaDeviceSynchronize`.
		Step 7	Transfer data from the device to the host with `cudaMemcpy`.
		Step 8	Consume the crunched data, which is printed in this case.

This book is not a replacement for the CUDA API guide and does not cover all CUDA APIs. For extensive use of the API, please refer to the CUDA API guide.

As we can see, the CUDA processing flow has some additional steps that need to be added to the sequential code. These are as follows:

Memory allocation on GPU: CPU memory and GPU memory are physically separate memory. malloc allocates memory on the CPU's RAM. The GPU kernel/device function can only access memory that's allocated/pointing to the device memory. To allocate memory on the GPU, we need to use the cudaMalloc API. Unlike the malloc command, cudaMalloc does not return a pointer to allocated memory; instead, it takes a pointer reference as a parameter and updates the same with the allocated memory.

Transfer data from host memory to device memory: The host data is then copied to the device's memory, which was allocated using the cudaMalloc command used in the previous step. The API that's used to copy the data between the host and device and vice versa is cudaMemcpy. Like other memcopy commands, this API requires the destination pointer, source pointer, and size. One additional parameter it takes is the direction of copy, that is, whether we are copying from the host to the device or from the device to the host. In the latest version of CUDA, this is optional since the driver is capable of understanding whether the pointer points to the host memory or device memory. Note that there is an asynchronous alternative to cudaMemcpy. This will be covered in more detail in other chapters.
Call and execute a CUDA function: As shown in the Hello World CUDA program, we call a kernel by using <<<,>>> brackets, which provide parameters for the block and thread size, respectively. We will cover this in more detail after all of the steps are complete.
Synchronize: As we mentioned in the Hello World program, kernel calls are asynchronous in nature. In order for the host to make sure that kernel execution has finished, the host calls the cudaDeviceSynchronize function. This makes sure that all of the previously launched device calls have finished.
Transfer data from host memory to device memory: Use the same cudaMemcpy API to copy the data back from the device to the host for post-processing or validation duties such as printing. The only change here, compared to the first step, is that we reverse the direction of the copy, that is, the destination pointer points to the host while the source pointer points to the device allocated in memory.
Free the allocated GPU memory: Finally, free the allocated GPU memory using the cudaFree API.

Change the sequential vector addition code's main function to reflect these new steps. The main function will look like this:

int main(void) {
    int *a, *b, *c;
    int *d_a, *d_b, *d_c; // device copies of a, b, c
    int size = N * sizeof(int);
 
    // Alloc space for host copies of a, b, c and setup input values
    a = (int *)malloc(size); fill_array(a);
    b = (int *)malloc(size); fill_array(b);
    c = (int *)malloc(size);
 
    // Alloc space for device copies of vector (a, b, c)
    cudaMalloc((void *)&d_a, N * sizeof(int));
    cudaMalloc((void *)&d_b, N *sizeof(int));
    cudaMalloc((void *)&d_c, N * sizeof(int));
 
    // Copy from host to device
    cudaMemcpy(d_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, N* sizeof(int), cudaMemcpyHostToDevice);
 
    device_add<<<1,1>>>(d_a,d_b,d_c);
 
    // Copy result back to host
    cudaMemcpy(c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);
 
    print_output(a,b,c);
    free(a); free(b); free(c);
 
    //free gpu memory
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
     
    return 0;
}

Now, let's look at how kernel code is written and manage the thread and block sizes. For this, we will be conducting multiple experiments.

Experiment 1 – creating multiple blocks

In this section, we will make use of CUDA blocks to run the vector addition code in parallel on the GPU. Additional keywords will be exposed that are related to how we can index CUDA blocks. Change the call to the device_add function, as follows:

//changing from device_add<<<1,1>>> to
device_add<<<N,1>>>

This will execute the device_add function N times in parallel instead of once. Each parallel invocation of the device_add function is referred to as a block. Now, let's add a __global__ device function, as follows:

__global__ void device_add(int *a, int *b, int *c) {
 c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

By using blockIdx.x to index the array, each block handles a different element of the array. On the device, each block can execute in parallel. Let's take a look at the following screenshot:

The preceding screenshot represents the vector addition GPU code in which every block shows indexing for multiple single-thread blocks.

Experiment 2 – creating multiple threads

In this section, we will make use of CUDA threads to run the vector addition code in parallel on GPU. Additional keywords will be exposed that are related to how we can index CUDA threads.

A block can be split into multiple threads. Change the call to the device_add function, as follows:

//changing from device_add<<<1,1>>> to
device_add<<<1,N>>>

This will execute the device_add function N times in parallel instead of once. Each parallel invocation of the device_add function is referred to as a thread. Change the device routine to reflect the kernel, as follows:

__global__ void device_add(int *a, int *b, int *c) {
     c[threadIdx.x] = a[threadIdx.x] + b[threadIdx.x];
}

One notable difference is that, instead of blockIdx.x, we make use of threadIdx.x, as shown in the following screenshot:

The preceding screenshot represents vector addition GPU code in which every block shows indexing for a single block-multiple threads.

Experiment 3 – combining blocks and threads

So far, we've looked at parallel vector addition through the use of several blocks with one thread in the Experiment 1 – creating multiple blocks section and one block with several threads in the Experiment 2 – creating multiple threads section. In this experiment, we'll use multiple blocks as well as separate blocks containing multiple threads. This becomes more challenging in terms of how to find the index because we need to combine both threadIdx and blockIdx to generate a unique ID.

Let's take a look at two scenarios that depict different combinations that the developer can choose from:

Scenario 1: Let's consider that the total number of vector elements is 32. Each block contains eight threads and a total of four blocks.
Scenario 2: Let's consider that the total number of vector elements is 32. Each block contains four threads and a total of eight blocks.

In both scenarios, the number of parallel executions is 32, where all 32 elements get populated in parallel. The developer makes the choice between the threads within a block and the number of blocks based on the problem's size and restriction by each piece of hardware. We will be covering details about the right choice of sizing based on the architecture in another chapter.

The following screenshot shows the vector addition GPU indexing code for different block and thread configurations:

Now, let's look at how the kernel code can be changed to combine both threads and blocks to calculate a global index:

__global__ void device_add(int *a, int *b, int *c) {
     int index = threadIdx.x + blockIdx.x * blockDim.x;
     c[index] = a[index] + b[index];
}

While calling the kernel from the main() function, the developer chooses the block and thread configuration, as depicted in the following code, for the two scenarios we mentioned previously:

Scenario 1: Following is the code for vector addition GPU grid and block size calculation for eight threads per block:

threads_per_block = 8;
no_of_blocks = N/threads_per_block;
device_add<<<no_of_blocks,threads_per_block>>>(d_a,d_b,d_c);

Scenario 2: Following is the code for vector addition GPU grid and block size calculation for four threads per block:

threads_per_block = 4;
no_of_blocks = N/threads_per_block;
device_add<<<no_of_blocks,threads_per_block>>>(d_a,d_b,d_c);

With a combination of threads and blocks, the unique ID of a thread can be calculated. As shown in the preceding code, another variable is given to all threads. This is called blockDim. This variable consists of the block's dimensions, that is, the number of threads per block. Let's take a look at the following screenshot:

Here, we can see a vector addition GPU indexing calculation for scenario 1.

Why bother with threads and blocks?

It might not be obvious why we need this additional hierarchy of threads and blocks. They add a level of complexity where the developer needs to find out the right block and grid size. Also, global indexing becomes a challenge. The reason for this is because of the restrictions that the CUDA programming model put it place.

Unlike parallel blocks, threads have mechanisms to communicate and synchronize efficiently. Real-world applications require threads to communicate with each other and may want to wait for certain data to be interchanged before proceeding further. This kind of operation requires threads to communicate, and the CUDA programming model allows this communication for threads within the same block. Threads belonging to different blocks cannot communicate/synchronize with each other during the execution of the kernel. This restriction allows the scheduler to schedule the blocks on the SM independently of each other. The result of this is that, if new hardware is released with more SMs and if the code has enough parallelism, the code can be scaled linearly. In other words, this allows the hardware to scale the number of blocks running in parallel based on the GPU's capability.

The threads communicate with each other using a special memory known as shared memory. We will cover shared memory extensively in Chapter 2, CUDA Memory Management, where we will expose other memory hierarchies in the GPU and their optimal usage. The following screenshot demonstrates scaling blocks across different GPUs consisting of different amounts of SMs:

Now, let's find out more about launching kernels in multiple dimensions.

Launching kernels in multiple dimensions

So far, we have been launching threads and blocks in one dimension. This means we have only using indexes for one dimension; for example, we've been using threadIdx.x, where x represents that we are using only an x dimension thread index. Similarly, we've been using blockIdx.x, where x represents that we are using only an x dimension block index. We can launch threads and blocks in one, two, or three dimensions. One example of launching threads and blocks in two dimensions is when we use parallel operations on an image, for example, to blur the image using a filter. The developer has the choice of launching threads and blocks in two dimensions, which is a more natural choice given that images are two-dimensional in nature.

It is important to understand that every GPU architecture also puts a restriction on the dimensions of threads and blocks. For example, the NVIDIA Pascal card allows a maximum of 1,024 threads per thread block in the x and y dimensions, while in the z dimension, you can only launch 64 threads. Similarly, the maximum blocks in a grid are restricted to 65,535 in the y and z dimensions in the Pascal architecture and 2^31 -1 in the x dimension. If the developer launches a kernel with an unsupported dimension, the application throws a runtime error.

So far, we have been assuming that the code we have written is error-free. But in the real world, every programmer writes code that has bugs in it, and it is necessary to catch those errors. In this next section, we'll take a look at how error reporting in CUDA works.

","url":"https://subscription.packtpub.com/book/programming/9781788996242/1/ch01lvl1sec04/vector-addition-using-cuda","metaDescription":"","titleSlug":"vector-addition-using-cuda"},{"id":"ch01lvl1sec05","contentType":"text","title":"Error reporting in CUDA","content":"

Error reporting in CUDA

In CUDA, the host code manages errors. Most CUDA functions call cudaError_t, which is basically an enumeration type. cudaSuccess (value 0) indicates a 0 error. The user can also make use of the cudaGetErrorString() function, which returns a string describing the error condition, as follows:

 cudaError_t e;
 e = cudaMemcpy(...);
 if(e)
     printf("Error: %sn", cudaGetErrorString(err));

Kernel launches have no return value. We can make use of a function such as cudaGetLastError() here, which returns the error code for the last CUDA function (including kernel launches). In the case of multiple errors, only the last one is reported:

MyKernel<<< ... >>> (...);
cudaDeviceSynchronize();
e = cudaGetLastError();

When it comes to production code, it is advised to make use of error checking code at logical checkpoints as the CPU code will continue with normal execution even if the GPU kernel has crashed, resulting in incorrect results.

In the next section, we will introduce you to the data types that are supported in the CUDA programming model.

","url":"https://subscription.packtpub.com/book/programming/9781788996242/1/ch01lvl1sec05/error-reporting-in-cuda","metaDescription":"","titleSlug":"error-reporting-in-cuda"},{"id":"ch01lvl1sec06","contentType":"text","title":"Data type support in CUDA","content":"

Data type support in CUDA

Like any processor architecture, a GPU also has different types of memories, each meant for a different purpose. We will cover them in more detail in Chapter 2, CUDA Memory Management. However, it is important to understand the different data types that are supported and their implications on performance and accuracy. CUDA programming supports all of the standard data types that developers are familiar with in terms of their respective languages. Along with standard data types with different sizes (char is 1 byte, float is 4 bytes, double is 8 bytes, and so on), it also supports vector types such as float2 and float4.

It is recommended that the data types are naturally aligned since aligned data access for data types that are 1, 2, 4, 8, or 16 bytes in size ensure that the GPU calls a single memory instruction. If they are not aligned, the compiler generates multiple instructions, which are interleaved, resulting in inefficient utilization of the memory and instruction bus. Due to this, the recommendation is to use types that are naturally aligned for data residing in GPU memory. The alignment requirement is automatically fulfilled for the built-in types of char, short, int, long, long long, float, and double such as float2 and float4.

Also, CUDA programming supports complex data structures such as structures and classes (in the context of C and C++). For complex data structures, the developer can make use of alignment specifiers to the compiler to enforce the alignment requirements, as shown in the following code:

struct __align__(16) {
    float r;
    float g;
    float b;
};

Every GPU has a limited set of cores and so the FLOPS are different. For example, a Tesla V100 card with the Volta architecture has 2,560 FP64 cores (double precision), while it has double the number of 32-bit single precision cores. It is quite evident that using the right data types based on the precision requirements of the algorithm is essential. Mixed precision algorithms are now being developed to make use of different types of cores where some part of the algorithm runs with higher precision while some parts run with lower precision. We will cover more on this topic in the upcoming chapters as well. For now, it is important to understand that the GPU memory hierarchy is different and, hence, using the right data type matters.

While this was a general introduction to the data types that are supported in GPU, more details about all of the supported data types can be found at https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#built-in-vector-types.

","url":"https://subscription.packtpub.com/book/programming/9781788996242/1/ch01lvl1sec06/data-type-support-in-cuda","metaDescription":"","titleSlug":"data-type-support-in-cuda"},{"id":"ch01lvl1sec07","contentType":"text","title":"Summary","content":"

Summary

In this chapter, we provided you with a perspective on heterogeneous computing with the help of history and high-performance computing. We went into detail about how the two processors, that is, CPU and GPU, are different. We also wrote a Hello World and vector addition CUDA program on a GPU. Finally, we looked at how to detect errors in CUDA since the majority of calls that are made to a CUDA API are asynchronous in nature.

In the next chapter, we will look at the different types of GPU memory that are available and how to utilize them optimally.

","url":"https://subscription.packtpub.com/book/programming/9781788996242/1/ch01lvl1sec07/summary","metaDescription":"","titleSlug":"summary"}]

Learn CUDA Programming

By : Jaegeun Han, Bharatkumar Sharma

Learn CUDA Programming

By: Jaegeun Han, Bharatkumar Sharma

Overview of this book

Related Content you might be interested in

Current Title:

Learn CUDA Programming

Hands-On GPU Programming with Python and CUDA

Hands-On GPU-Accelerated Computer Vision with OpenCV and CUDA

Hands-On GPU Computing with Python