March 17, 2009 at 16:00CUDA: How the GPU Works
The nVidia GPU internal model is the key to understanding GPGPUs using CUDA. This time I will try to tell in more detail about the software device GPUs. I will talk about the key points of the CUDA compiler, the CUDA runtime API, and, in conclusion, I will give an example of using CUDA for simple mathematical calculations.
Let's get started.
Consider the GPU computational model in more detail.
Fig. 2. GPU unit device.
When using the GPU, you can use the grid of the required size and configure the blocks to the needs of your task.
The CUDA technology itself (the nvcc.exe compiler) introduces a number of additional extensions for the C language, which are necessary for writing code for the GPU:
As mentioned, function specifiers determine how and where functions will be called. In total, CUDA has 3 such qualifiers:
Kernel start qualifiers are used to describe the number of blocks, threads and memory that you want to allocate when calculating on a GPU. The kernel startup syntax is as follows:
myKernelFunc <<>> (float * param1, float * param2), where
And of course, myKernelFunc itself is a kernel function (__global__ specifier). Some variables can be omitted when invoking the kernel, for example sharedMemSize and cudaStream.
It is also worth mentioning the built-in variables:
By the way, gridDim and blockDim are the very variables that we pass when the GPU kernel starts up, however, they can be read only in the kernel.
Additional types of variables and their qualifiers will be considered directly in the examples of working with memory.
Before you start using CUDA directly for computing, you need to familiarize yourself with the so-called CUDA host API, which is the link between the CPU and GPU. The CUDA host API, in turn, can be divided into a low-level API called the CUDA driver API, which provides access to the CUDA user-mode driver, and a high-level API, the CUDA runtime API. In my examples, I will use the CUDA runtime API.
The CUDA runtime API includes the following groups of functions:
As said, a thread is a direct performer of calculations. How, then, is the parallelization of computations between threads? Consider the operation of a single block.
Task. It is required to calculate the sum of two vectors of dimension N elements.
We know the maximum sizes of our block: 512 * 512 * 64 threads. Since our vector is one-dimensional, so far we will restrict ourselves to using the x-dimension of our block, that is, we will use only one strip of threads from the block (Fig. 3).
Fig. 3. Our strip of threads from the used block.
Note that the x-dimension of block 512, that is, we can add up vectors at a time, the length of which is N <= 512 elements. In other cases, with more massive calculations, you can use a larger number of blocks and multidimensional arrays. I also noticed one interesting feature, perhaps some of you thought that you can use 512 * 512 * 64 = 16777216 threads in one block, of course this is not so, in general, this product cannot exceed 512 (at least my graphics card).
In the program itself, you must perform the following steps:
We proceed directly to writing the code:
First, we write the kernel function, which will carry out the addition of vectors:
Thus, parallelization will be performed automatically when the kernel starts. This function also uses the built-in variable threadIdx and its field x, which allows you to set the correspondence between the calculation of a vector element and the thread in the block. We do the calculation of each element of the vector in a separate thread.
We write the code that is responsible for points 1 and 2 in the program:
To allocate memory on the video card, the cudaMalloc function is used , which has the following prototype:
cudaError_t cudaMalloc (void ** devPtr, size_t count), where
Returns:
To copy data to the video card's memory, cudaMemcpy is used, which has the following prototype:
cudaError_t cudaMemcpy (void * dst, const void * src, size_t count, enum cudaMemcpyKind kind), where
Returns:
Now we turn to the direct call of the kernel for calculation on the GPU.
In our case, it is not necessary to determine the size of the grid and block, since we use only one block and one measurement in the block, so the code above can be written:
Now we just have to copy the calculation result from the video memory to the host memory. But the kernel functions have a peculiarity - asynchronous execution, that is, if the next section of code starts working after calling the kernel, this does not mean that the GPU has performed the calculations. To complete a given kernel function, you must use synchronization tools, such as events. Therefore, before copying the results to the host, we synchronize the GPU threads through event.
Code after calling the kernel:
Let's take a closer look at the functions from the Event Managment API.
Event is created using the cudaEventCreate function , the prototype of which is of the form:
cudaError_t cudaEventCreate (cudaEvent_t * event), where
Returns:
Event recording is performed using the cudaEventRecord function , the prototype of which is:
cudaError_t cudaEventRecord (cudaEvent_t event, CUstream stream), where
Returns:
Event synchronization is performed by the cudaEventSynchronize function. This function expects the completion of all GPU threads and the passage of the specified event, and only then gives control to the calling program. The prototype of the function is:
cudaError_t cudaEventSynchronize (cudaEvent_t event), where
Returns:
Understand how cudaEventSynchronize, it is possible from the following diagram:
Fig. 4. Synchronization of the work of the main and GPU programs.
In Figure 4, the block “Waiting for the Event to pass” is the cudaEventSynchronize function call.
Well, in conclusion, display the result on the screen and clean the allocated resources.
I think that it is not necessary to describe the functions of releasing resources. Unless, it can be recalled that they also return cudaError_t values, if there is a need to check their work.
Hope this material helps you understand how the GPU works. I described the most important points that you need to know to work with CUDA. Try to write the addition of two matrices yourself, but do not forget about the hardware limitations of the video card.
PS: It didn’t work very briefly. I hope not tired. If you need all the source code, I can send it by mail.
PSS: Ask questions.
Let's get started.
GPU Computing Model:
Consider the GPU computational model in more detail.
- The upper level of the GPU core consists of blocks that are grouped in a grid or grid of dimension N1 * N2 * N3. This can be represented as follows:
Fig. 1. Computing device GPU.
The dimension of the grid of blocks can be found using the cudaGetDeviceProperties function; in the resulting structure, the maxGridSize field is responsible for this. For example, on my GeForce 9600M GS the block grid dimension is 65535 * 65535 * 1, that is, the block grid I have is two-dimensional (the received data satisfies Compute Capability v.1.1). - Any block in turn consists of threads, which are the direct executors of the calculations. The threads in the block are formed in the form of a three-dimensional array (Fig. 2), the dimension of which can also be found using the cudaGetDeviceProperties function; the maxThreadsDim field is responsible for this.
Fig. 2. GPU unit device.
When using the GPU, you can use the grid of the required size and configure the blocks to the needs of your task.
CUDA and C language:
The CUDA technology itself (the nvcc.exe compiler) introduces a number of additional extensions for the C language, which are necessary for writing code for the GPU:
- Function specifiers that show how and where the functions will be executed from.
- Variable qualifiers that indicate the type of GPU memory used.
- GPU kernel launch qualifiers.
- Built-in variables for identifying threads, blocks, and other parameters when executing code in the GPU core.
- Additional types of variables.
As mentioned, function specifiers determine how and where functions will be called. In total, CUDA has 3 such qualifiers:
- __host__ - will be executed on the CPU, called from the CPU (in principle, it can be omitted).
- __global__ - runs on the GPU, called from the CPU.
- __device__ - runs on the GPU, called from the GPU.
Kernel start qualifiers are used to describe the number of blocks, threads and memory that you want to allocate when calculating on a GPU. The kernel startup syntax is as follows:
myKernelFunc <<
- gridSize - dimension of the grid of blocks (dim3) allocated for calculations,
- blockSize - block size (dim3) allocated for calculations,
- sharedMemSize - the size of the additional memory allocated at kernel startup,
- cudaStream - cudaStream_t variable that defines the stream in which the call will be made.
And of course, myKernelFunc itself is a kernel function (__global__ specifier). Some variables can be omitted when invoking the kernel, for example sharedMemSize and cudaStream.
It is also worth mentioning the built-in variables:
- gridDim - grid dimension, of type dim3. Lets you know the size of the grid allocated during the current kernel call.
- blockDim - block dimension, also of type dim3. Allows you to find out the size of the block allocated during the current kernel call.
- blockIdx - the index of the current block in the calculation on the GPU, is of type uint3.
- threadIdx - index of the current thread in the calculation on the GPU, is of type uint3.
- warpSize - the size of warp, it is of type int (I have not tried to use it myself).
By the way, gridDim and blockDim are the very variables that we pass when the GPU kernel starts up, however, they can be read only in the kernel.
Additional types of variables and their qualifiers will be considered directly in the examples of working with memory.
CUDA host API:
Before you start using CUDA directly for computing, you need to familiarize yourself with the so-called CUDA host API, which is the link between the CPU and GPU. The CUDA host API, in turn, can be divided into a low-level API called the CUDA driver API, which provides access to the CUDA user-mode driver, and a high-level API, the CUDA runtime API. In my examples, I will use the CUDA runtime API.
The CUDA runtime API includes the following groups of functions:
- Device Management - includes functions for general GPU control (receiving information about GPU capabilities, switching between GPUs when working in SLI mode, etc.).
- Thread Management - thread management.
- Stream Management - stream management.
- Event Management - the function of creating and managing events.
- Execution Control - CUDA kernel launch and execution functions.
- Memory Management - GPU memory management features.
- Texture Reference Manager - work with texture objects through CUDA.
- OpenGL Interoperability - features for interacting with the OpenGL API.
- Direct3D 9 Interoperability - features for interacting with the Direct3D 9 API.
- Direct3D 10 Interoperability - features for interacting with the Direct3D 10 API.
- Error Handling - error handling functions.
Understanding the work of the GPU:
As said, a thread is a direct performer of calculations. How, then, is the parallelization of computations between threads? Consider the operation of a single block.
Task. It is required to calculate the sum of two vectors of dimension N elements.
We know the maximum sizes of our block: 512 * 512 * 64 threads. Since our vector is one-dimensional, so far we will restrict ourselves to using the x-dimension of our block, that is, we will use only one strip of threads from the block (Fig. 3).
Fig. 3. Our strip of threads from the used block.
Note that the x-dimension of block 512, that is, we can add up vectors at a time, the length of which is N <= 512 elements. In other cases, with more massive calculations, you can use a larger number of blocks and multidimensional arrays. I also noticed one interesting feature, perhaps some of you thought that you can use 512 * 512 * 64 = 16777216 threads in one block, of course this is not so, in general, this product cannot exceed 512 (at least my graphics card).
In the program itself, you must perform the following steps:
- Get data for calculations.
- Copy this data to the GPU memory.
- Perform a calculation in the GPU through the kernel function.
- Copy calculated data from GPU memory to RAM.
- View results.
- Free up used resources.
We proceed directly to writing the code:
First, we write the kernel function, which will carry out the addition of vectors:
// Функция сложения двух векторов
__global__ void addVector(float* left, float* right, float* result)
{
//Получаем id текущей нити.
int idx = threadIdx.x;
//Расчитываем результат.
result[idx] = left[idx] + right[idx];
}
* This source code was highlighted with Source Code Highlighter.Thus, parallelization will be performed automatically when the kernel starts. This function also uses the built-in variable threadIdx and its field x, which allows you to set the correspondence between the calculation of a vector element and the thread in the block. We do the calculation of each element of the vector in a separate thread.
We write the code that is responsible for points 1 and 2 in the program:
#define SIZE 512
__host__ int main()
{
//Выделяем память под вектора
float* vec1 = new float[SIZE];
float* vec2 = new float[SIZE];
float* vec3 = new float[SIZE];
//Инициализируем значения векторов
for (int i = 0; i < SIZE; i++)
{
vec1[i] = i;
vec2[i] = i;
}
//Указатели на память видеокарте
float* devVec1;
float* devVec2;
float* devVec3;
//Выделяем память для векторов на видеокарте
cudaMalloc((void**)&devVec1, sizeof(float) * SIZE);
cudaMalloc((void**)&devVec2, sizeof(float) * SIZE);
cudaMalloc((void**)&devVec3, sizeof(float) * SIZE);
//Копируем данные в память видеокарты
cudaMemcpy(devVec1, vec1, sizeof(float) * SIZE, cudaMemcpyHostToDevice);
cudaMemcpy(devVec2, vec2, sizeof(float) * SIZE, cudaMemcpyHostToDevice);
…
}
* This source code was highlighted with Source Code Highlighter.To allocate memory on the video card, the cudaMalloc function is used , which has the following prototype:
cudaError_t cudaMalloc (void ** devPtr, size_t count), where
- devPtr - a pointer to the address of the allocated memory,
- count - the size of the allocated memory in bytes.
Returns:
- cudaSuccess - upon successful allocation of memory
- cudaErrorMemoryAllocation - on memory allocation error
To copy data to the video card's memory, cudaMemcpy is used, which has the following prototype:
cudaError_t cudaMemcpy (void * dst, const void * src, size_t count, enum cudaMemcpyKind kind), where
- dst - pointer containing the address of the copy destination,
- src - pointer containing the copy source address,
- count - the size of the copied resource in bytes,
- cudaMemcpyKind - an enumeration indicating the direction of copying (may be cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyHostToHost, cudaMemcpyDeviceToDevice).
Returns:
- cudaSuccess - upon successful copying
- cudaErrorInvalidValue - invalid argument parameters (for example, copy size is negative)
- cudaErrorInvalidDevicePointer - invalid memory pointer in the graphics card
- cudaErrorInvalidMemcpyDirection - wrong direction (for example, the source and destination of the copy are confused)
Now we turn to the direct call of the kernel for calculation on the GPU.
…
dim3 gridSize = dim3(1, 1, 1); //Размер используемого грида
dim3 blockSize = dim3(SIZE, 1, 1); //Размер используемого блока
//Выполняем вызов функции ядра
addVector<<>>(devVec1, devVec2, devVec3);
…
* This source code was highlighted with Source Code Highlighter. In our case, it is not necessary to determine the size of the grid and block, since we use only one block and one measurement in the block, so the code above can be written:
addVector<<<1, SIZE>>>(devVec1, devVec2, devVec3);
* This source code was highlighted with Source Code Highlighter.Now we just have to copy the calculation result from the video memory to the host memory. But the kernel functions have a peculiarity - asynchronous execution, that is, if the next section of code starts working after calling the kernel, this does not mean that the GPU has performed the calculations. To complete a given kernel function, you must use synchronization tools, such as events. Therefore, before copying the results to the host, we synchronize the GPU threads through event.
Code after calling the kernel:
//Выполняем вызов функции ядра
addVector<<>>(devVec1, devVec2, devVec3);
//Хендл event'а
cudaEvent_t syncEvent;
cudaEventCreate(&syncEvent); //Создаем event
cudaEventRecord(syncEvent, 0); //Записываем event
cudaEventSynchronize(syncEvent); //Синхронизируем event
//Только теперь получаем результат расчета
cudaMemcpy(vec3, devVec3, sizeof(float) * SIZE, cudaMemcpyDeviceToHost);
* This source code was highlighted with Source Code Highlighter. Let's take a closer look at the functions from the Event Managment API.
Event is created using the cudaEventCreate function , the prototype of which is of the form:
cudaError_t cudaEventCreate (cudaEvent_t * event), where
- * event - a pointer to write the event handle.
Returns:
- cudaSuccess - if successful
- cudaErrorInitializationError - initialization error
- cudaErrorPriorLaunchFailure - error during previous asynchronous function start
- cudaErrorInvalidValue - invalid value
- cudaErrorMemoryAllocation - memory allocation error
Event recording is performed using the cudaEventRecord function , the prototype of which is:
cudaError_t cudaEventRecord (cudaEvent_t event, CUstream stream), where
- event - the handle of the event
- stream - the number of the stream in which we write (in our case, this is the main zero stream).
Returns:
- cudaSuccess - if successful
- cudaErrorInvalidValue - invalid value
- cudaErrorInitializationError - initialization error
- cudaErrorPriorLaunchFailure - error during previous asynchronous function start
- cudaErrorInvalidResourceHandle - invalid event handle
Event synchronization is performed by the cudaEventSynchronize function. This function expects the completion of all GPU threads and the passage of the specified event, and only then gives control to the calling program. The prototype of the function is:
cudaError_t cudaEventSynchronize (cudaEvent_t event), where
- event - the event handle, the passage of which is expected.
Returns:
- cudaSuccess - if successful
- cudaErrorInitializationError - initialization error
- cudaErrorPriorLaunchFailure - error during previous asynchronous function start
- cudaErrorInvalidValue - invalid value
- cudaErrorInvalidResourceHandle - invalid event handle
Understand how cudaEventSynchronize, it is possible from the following diagram:
Fig. 4. Synchronization of the work of the main and GPU programs.
In Figure 4, the block “Waiting for the Event to pass” is the cudaEventSynchronize function call.
Well, in conclusion, display the result on the screen and clean the allocated resources.
//Результаты расчета
for (int i = 0; i < SIZE; i++)
{
printf("Element #%i: %.1f\n", i , vec3[i]);
}
//
// Высвобождаем ресурсы
//
cudaEventDestroy(syncEvent);
cudaFree(devVec1);
cudaFree(devVec2);
cudaFree(devVec3);
delete[] vec1; vec1 = 0;
delete[] vec2; vec2 = 0;
delete[] vec3; vec3 = 0;
* This source code was highlighted with Source Code Highlighter.I think that it is not necessary to describe the functions of releasing resources. Unless, it can be recalled that they also return cudaError_t values, if there is a need to check their work.
Conclusion
Hope this material helps you understand how the GPU works. I described the most important points that you need to know to work with CUDA. Try to write the addition of two matrices yourself, but do not forget about the hardware limitations of the video card.
PS: It didn’t work very briefly. I hope not tired. If you need all the source code, I can send it by mail.
PSS: Ask questions.
