Understanding the basics of CUDA thread hierarchies

Elemar Júnior

In this post, I would like to explain a basic but confusing concept of CUDA programming: Thread Hierarchies. It will not be an exhaustive reference. We will not cover all aspects, but it could be a nice first step.

If you are starting with CUDA and want to know how to setup your environment, using VS2017, I recommend you to read this post.

From CPU to GPU

To get started, let’s write something straightforward to run on the CPU.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <cstdio>

void printHelloCPU()
{
	printf("Hello World from the CPU");
}

int main()
{
	printHelloCPU();
	getchar();
	return 0;
}

Now, let’s change this code to run on the GPU.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <cstdio>

__global__ void printHelloGPU()
{
	printf("Hello World from the GPUn");
}

int main()
{
	printHelloGPU<<<1, 1>>>();
	cudaDeviceSynchronize();
	
	getchar();
	return 0;
}

The cudaDeviceSyncronize function determines that all the processing on the GPU must be done before continuing.

Let’s remember some concepts we learned in a previous post:

  • The __global__ keyword indicates that the following function will run on the GPU.
  • The code executed on the CPU is referred to as host code, and code executed on the GPU is referred to as device code.
  • It is required that functions defined with the __global__ keyword return type void.
  • When calling a function to run on the GPU, we call this function a kernel (In the example, printHelloGPU is the kernel).
  • When launching a kernel, we must provide an execution configuration, which is done by using the <<< ... >>> syntax.

The Execution Configuration

At a high level, the execution configuration allows programmers to specify the thread hierarchy for a kernel launch, which defines the number of thread blocks, as well as
how many threads to execute in each block.

Notice, in the previous example, the kernel is launching with 1 block of threads (the first execution configuration argument) which contains 1 thread (the second configuration argument).

The execution configuration allows programmers to specify details about launching the kernel to run in parallel on multiple GPU threads. The syntax for this is:

<<< NUMBER_OF_BLOCKS, NUMBER_OF_THREADS_PER_BLOCK>>>

A kernel is executed once for every thread in every thread block configured when the kernel is launched.

Thus, under the assumption that a kernel called printHelloGPU has been defined, the following are true:

  • printHelloGPU<<<1, 1>>>() is configured to run in a single thread block which has a single thread and will, therefore, run only once.
  • printHelloGPU<<<1, 5>>>() is configured to run in a single thread block which has 5 threads and will, therefore, run 5 times.
  • printHelloGPU<<<5, 1>>>() is configured to run in 5 thread blocks which each have a single thread and will, therefore, run five times.
  • printHelloGPU<<<5, 5>>>() is configured to run in 5 thread blocks which each have five threads and will, therefore, run 25 times.

Let me try to explain this graphically:

In the drawing, each blue rectangle represents a thread. Each gray rectangle represents a block.The green rectangle represents the grid.

Thread Hierarchy Variables

In the kernel’s code, we can access variables provided by CUDA. These variables describe the thread, thread block, and grid.

gridDim.x is the number of the blocks in the grids.

blockIdx.x is the index of the current block within the grid.

blockDim.x is the number of threads in the block. All blocks in a grid contain the same number of threads.

threadIdx.x is index of the thread within a block (starting at 0).

X… Y, and Z

As you noted, we have been using the suffix .x for all variables. But, we could use, .y and .z as well.

The CUDA threads hierarchy can be 3-dimensional.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <cstdio>


__global__ void printHelloGPU()
{
	printf("Hello x: #%d  y: #%dn", threadIdx.x, threadIdx.y);
}

int main()
{
	dim3 threads(3, 3);
	printHelloGPU<<<1, threads>>>();
	cudaDeviceSynchronize();
	
	getchar();
	return 0;
}

We can use the dim3 structure to specify dimensions for blocks and threads. In the example, we specified we are creating a 2-dimensional structure (3x3x1).

Deciding what execution configuration to use

Consider:

  • printHelloGPU<<<1, 25>>>() is configured to run in a single thread block which has 25 threads and will, therefore, run 25 times.
  • printHelloGPU<<<1, dim3(5, 5)>>>() is configured to run in a single thread block which has 25 threads and will, therefore, run 25 times.
  • printHelloGPU<<<5, 5>>>() is configured to run in 5 thread blocks which each has 5 threads and will therefore run 25 times.

So, what configuration is right? Answer: All choices are valid. What should you use? It depends.

As you know, each thread will run the kernel once. If you are working on some data in memory, you should use the configuration that makes easier to address the data, using the thread hierarchy variables. Also, your graphics card has limitations that you need to consider.

Conclusions

If you are like me, you will need some time to understand Thread Hierarchies. In future posts, I will start to share some practical examples that can make it simpler.

For a while, feel free to comment this post.

Cover: Ilze Lucero

Compartilhe este insight:

Comentários

Participe deixando seu comentário sobre este artigo a seguir:

Subscribe
Notify of
guest
2 Comentários
Oldest
Newest Most Voted
Inline Feedbacks
View all comments
binhminhdanh2
binhminhdanh2
2 anos atrás

why dont meet unlikely it?
i can and i qualifications out
assess me[b][i][/i][/b]
LSA Intercession Inducement

binhminhdanh2
binhminhdanh2
2 anos atrás

why dont gain it?
i can and i elongate
assessment me[b][i][/i][/b]
review-product-launch

AUTOR

Elemar Júnior
Fundador e CEO da EximiaCo atua como tech trusted advisor ajudando empresas e profissionais a gerar mais resultados através da tecnologia.

NOVOS HORIZONTES PARA O SEU NEGÓCIO

Nosso time está preparado para superar junto com você grandes desafios tecnológicos.

Entre em contato e vamos juntos utilizar a tecnologia do jeito certo para gerar mais resultados.

Insights EximiaCo

Confira os conteúdos de negócios e tecnologia desenvolvidos pelos nossos consultores:

Arquitetura de Dados

Insights de um DBA na análise de um plano de execução

Especialista em performance de Bancos de Dados de larga escala
Arquitetura de Software

Estratégias para modernização do legado

Desenvolvedor .NET/NodeJs e especialista em Kafka com experiência em startups e grandes empresas
Infraestrutura e Nuvem

Migração para a nuvem, mais do que mudança tecnológica, implica em mudança da cultura organizacional

Engenheiro de nuvem, arquiteto de software e especialista em Containers e Devops

Acesse nossos canais

Simplificamos, potencializamos e aceleramos resultados usando a tecnologia do jeito certo

EximiaCo 2022 – Todos os direitos reservados

2
0
Queremos saber a sua opinião, deixe seu comentáriox
()
x

Understanding the basics of CUDA thread hierarchies

Para se candidatar nesta turma aberta, preencha o formulário a seguir:

Condição especial de pré-venda: R$ 14.000,00 - contratando a mentoria até até 31/01/2023 e R$ 15.000,00 - contratando a mentoria a partir de 01/02/2023, em até 12x com taxas.

Tenho interesse nessa capacitação

Para solicitar mais informações sobre essa capacitação para a sua empresa, preencha o formulário a seguir:

Tenho interesse em conversar

Se você está querendo gerar resultados através da tecnologia, preencha este formulário que um de nossos consultores entrará em contato com você:

O seu insight foi excluído com sucesso!

O seu insight foi excluído e não está mais disponível.

O seu insight foi salvo com sucesso!

Ele está na fila de espera, aguardando ser revisado para ter sua publicação programada.

Tenho interesse em conversar

Se você está querendo gerar resultados através da tecnologia, preencha este formulário que um de nossos consultores entrará em contato com você:

Tenho interesse nessa solução

Se você está procurando este tipo de solução para o seu negócio, preencha este formulário que um de nossos consultores entrará em contato com você:

Tenho interesse neste serviço

Se você está procurando este tipo de solução para o seu negócio, preencha este formulário que um de nossos consultores entrará em contato com você:

× Precisa de ajuda?