CUDA C na bancada: shared memory, medição honesta e o caminho até o Jetson

CUDA C na bancada: shared memory, medição honesta e o caminho até o Jetson

Quem chega em CUDA pelo tutorial de “hello GPU” costuma sair com a sensação de que o problema acabou: o kernel rodou, o printf apareceu e o cudaMemcpy devolveu o vetor. Em projeto real — especialmente em visão embarcada, fusão de sensores ou pré-processamento em SBC NVIDIA — o gargalo quase nunca é “saber lançar um kernel”. É saber o que a memória está fazendo, medir sem se enganar e escolher se o ganho cabe no envelope térmico e energético da placa.

O repositório c_cuda foi montado exatamente nessa lógica de bancada: oito exemplos progressivos (cuda_1 … cuda_8), Makefile que detecta a arquitetura da GPU e tabelas de tempo medidas em duas máquinas diferentes. Este artigo não reescreve o README. Usa os números publicados lá como âncora e discute as decisões de projeto por trás de shared memory, reduction e streams — e por que isso importa quando o destino não é só um desktop com GTX 1660, mas um caminho de edge que termina no Jetson.

Por que começar “vazio” e com erro explícito

cuda_1 (kernel vazio + printf) e cuda_2 (soma de dois números com error checking) parecem triviais. Em bancada, eles resolvem três falhas clássicas de setup:

1. Toolchain e runtime: nvcc, driver e capability batem? O Makefile do repo força essa conversa cedo: make tenta descobrir a arch; se não, você força make ARCH=sm_50 (MX130) ou make ARCH=sm_75 (GTX 1660). Compilar para a capability errada é uma forma silenciosa de “funciona no meu notebook e falha no outro”.
2. Modelo host/device: até o primeiro <<<grid, block>>> e o primeiro cudaDeviceSynchronize, muita gente ainda pensa em CUDA como “uma função C mais rápida”. O kernel vazio força a distinção entre código que roda na CPU e código que o driver despacha para a SMD.
3. Error checking como hábito: cuda_2 e cuda_3 insistirem em checar retorno não é pedantismo. Em kernels de matriz e reduction, um cudaMemcpy falho ou um launch com grid inválido se manifesta como resultado “quase certo” — o pior tipo de bug em validação de algoritmo.

cuda_3 (soma de vetores) fecha o ciclo de indexação 1D: threadIdx, blockIdx, stride e limites. É o ponto em que o engenheiro deixa de tratar a GPU como calculadora mágica e passa a tratar como arranjo de workers com memória própria.

Armadilhas de setup que o README não dramatiza

Na prática de bancada, três sintomas voltam com frequência:

Sem essa higiene, shared memory e streams viram teatro: você mede ruído de setup, não ganho de algoritmo.

Matriz em global memory: o baseline honesto

cuda_4 multiplica matrizes só com global memory e já introduz timing com cudaEvent. Esse exemplo é o baseline de decisão, não o “código final”. Em multiplicação de matrizes densa, cada elemento de saída reutiliza linhas e colunas; se cada thread busca tudo direto da DRAM da GPU, você paga latência e largura de banda demais. Ainda assim, você precisa desse número: sem baseline, “shared memory ficou 2× mais rápido” vira marketing interno.

A lição de projeto é simples: otimize contra uma medição, não contra a intuição. Em visão embarcada, o equivalente é medir o pré-processamento (resize, conversão de cor, convolução) no caminho real de câmera antes de reescrever em TensorRT ou CUDA custom.

Como medir sem se enganar

cudaEventRecord / cudaEventElapsedTime medem o device com menos ruído do que clock() no host, mas ainda dá para mentir para si mesmo:

• Esquecer cudaDeviceSynchronize (ou o sync implícito do evento) e cronometrar só o enqueue.
• Incluir o primeiro launch “frio” (JIT, alocação, warmup de cache) na média de produção.
• Misturar H2D + kernel + D2H num único número e chamar isso de “tempo do algoritmo”.
• Comparar MX130 em WSL com 1660 nativa e concluir “shared memory não vale a pena”.

Protocolo mínimo alinhado ao espírito do c_cuda: warmup descartado, N corridas, GPU e ambiente nomeados, e — quando houver claim de overlap — decomposição transferência versus compute.

Shared memory tiled: o salto conceitual

cuda_5 traz a versão tiled com shared memory; cuda_6 compara lado a lado Global vs Shared. Aqui mora o ângulo do repositório.

Shared memory é memória on-chip, por bloco, com latência muito menor que a global, mas com tamanho limitado e regras de bank conflict. O padrão tiled carrega um tile de A e um tile de B para __shared__, sincroniza o bloco (__syncthreads), computa a contribuição parcial e avança o tile. O ganho vem de reuso de dados: o mesmo elemento lido da global alimenta várias contas dentro do bloco.

Nos benchmarks publicados no README do c_cuda (WSL):

Três leituras de engenharia:

• O speedup relativo não é constante entre GPUs. Na MX130 o ganho relativo foi maior; na 1660 o absoluto é muito menor (8.4 ms → 4.1 ms). Em decisão de produto, “2,5×” na GPU fraca pode ser a diferença entre caber ou não no ciclo de 30 fps; na GPU forte, o mesmo algoritmo pode já estar folgado e o próximo gargalo ser I/O de câmera.
• A GTX 1660 ficou ~4–6× mais rápida que a MX130 nos kernels testados no repo. Isso calibra expectativa: portar o mesmo .cu para um notebook Maxwell de entrada não é “só recompilar e esquecer”.
• Medição honesta inclui ambiente. O README também compara WSL vs Windows nativo na mesma GTX 1660: global memory 8.43 ms (WSL) vs 7.90 ms (nativo); shared 4.10 ms nos dois. A conclusão do autor — kernels de compute parecidos; nativo um pouco melhor em transferências e kernels curtos — é exatamente o tipo de nuance que some em posts que citam um único número “milagroso”.

Para quem vem de automação: pense em shared memory como um buffer local de bloco análogo a um cache de ciclo rápido no CLP — útil quando há reuso, inútil (ou prejudicial) se o padrão de acesso não justifica a sincronização e o footprint.

Quando shared memory não paga

Nem todo kernel mere tiled. Shared memory deixa de ser vitória fácil quando:

• O padrão de acesso tem pouco reuso (streaming puro de um passe).
• O tile força __syncthreads demais e o bloco fica ocioso esperando o straggler.
• Bank conflicts transformam “memória rápida” em fila serializada.
• O footprint __shared__ reduz occupancy a ponto de esconder o ganho de latência.

A decisão correta no cuda_6 não é “sempre Shared”. É medir Global, medir Shared, escolher com a tabela na mão. Em Jetson Nano, onde a margem de milissegundos é curta, esse hábito vale mais do que memorizar a API.

Parallel reduction: de “somar vetor” a padrão de biblioteca

cuda_7 faz parallel reduction com shared memory (1M elementos no benchmark do README: ~5–6 ms na MX130, 1.4 ms na GTX 1660 em WSL). Reduction é o esqueleto de soma, máximo, norma, histogramas parciais e várias etapas de redes neurais. O ponto de ensino não é o algoritmo em si — é reconhecer padrões de comunicação entre threads: árvore de redução, warps, e o fato de que a primeira implementação correta quase nunca é a mais rápida (bank conflicts, divergência, últimas etapas sequenciais).

O README lista próximos objetivos (bank conflicts, occupancy, __shfl, scan, histogramas, Nsight). Isso importa na narrativa do site: o repositório se posiciona como laboratório contínuo, não como checklist concluído. Em conteúdo técnico monetizável, honestidade de roadmap gera mais confiança do que fingir cobertura completa de CUDA.

Redução na bancada: o que quebra primeiro

Reduction bem medida ensina humildade: o algoritmo “certo” no quadro branco ainda pode ser lento no silício.

Streams: quando o kernel não é o único ator

cuda_8 introduz CUDA Streams e overlap CPU/GPU. Números do README (WSL):

Streams brilham quando há transferência + compute para pipelinar. Em desktop com PCIe razoável e workload balanceado, o overlap aparece. Em Jetson, a memória é unificada de forma diferente (e Unified Memory / zero-copy entram na conversa); o princípio continua: esconder latência de movimento de dados. O erro clássico é otimizar só o kernel e deixar o pipeline de câmera → CPU → H2D → kernel → D2H serializado.

Comparação WSL vs nativo no README ainda mostra nuance: na 1660, speedup de streams foi 1.53× no WSL e 1.10× no Windows nativo para aquele ensaio — outro lembrete de que “ganho de stream” depende do balanceamento do pipeline no ambiente medido.

Tabela de decisão: streams ou não?

Streams não são “modo turbo”. São ferramenta de escalonamento de latência. Sem perfil de transferência, o cudaStreamCreate só adiciona complexidade.

Makefile com auto GPU arch: detalhe chato, decisão certa

Detectar arquitetura no make (Windows + Linux/WSL) parece infra. Em time misto notebook/desktop/CI, é o que evita binários incompatíveis e discussões intermináveis de “no meu PC compila”. Forçar ARCH=sm_50 ou sm_75 documenta a intenção. Para Jetson, o equivalente mental é fixar a compute capability da geração (e o JetPack) no script de build — não deixar o default do host “vazar” para o target.

Armadilhas de build multi-GPU

• Compilar no host sm_75 e copiar o binário para um ambiente que só tem Maxwell.
• Confiar no default do nvcc da máquina de CI sem pin de JetPack no target.
• Misturar fatbin multi-arch “por via das dúvidas” sem medir tamanho e tempo de load no edge.
• Tratar warning de arch como ruído cosmético — às vezes é incompatibilidade adiada.

O Makefile do c_cuda não é cerimônia: é o contrato de reproducibilidade entre MX130, 1660 e, depois, Jetson.

O edge path até o Jetson

Por que este estudo de CUDA C no desktop alimenta a linha Jetson do site?

1. Os mesmos conceitos de memória (global, shared, registers) aparecem em kernels custom, em plugins TensorRT e em pré-processamento CUDA de câmera CSI.
2. A disciplina de medição (cudaEvent, tabelas por GPU, WSL vs nativo) é o antídoto contra demos que “rodam 60 fps” sem dizer resolução, batch, precision nem temperatura.
3. A curva MX130 → 1660 ensina a ler hardware fraco vs forte. O Jetson Nano (Maxwell 128 CUDA cores, 4 GB) está mais próximo do regime “cada milissegundo conta” do que do regime desktop Turing folgado. Shared memory e redução bem feitas importam mais quando a margem é pequena.
4. Streams e overlap antecipam pipelines reais: captura, pré-processamento, inferência e pós-processamento em estágios — o mesmo desenho mental do protótipo 4x4 / veículo de teste com Jetson na Wilson Tecnologia.

Ponte industrial: do .cu ao chão de fábrica / AMR

Em automação e robótica móvel, CUDA raramente é o produto. É o acelerador de um estágio. A ponte honesta:

O engenheiro de painel que já aprendeu a não confiar em “melhorou” sem trend de ciclo reconhece o método. CUDA muda o silício; não muda a ética de medição.

Decisões de projeto que o código sugere

• Comece pelo baseline global (cuda_4 / lado Global de cuda_6) antes de shared. Sem isso, você não sabe se o tiling pagou.
• Publique GPU, ambiente e métrica juntos. “51.1 ms” só faz sentido com MX130 + WSL + Global Memory no cuda_6.
• Trate WSL como ambiente de desenvolvimento válido, mas valide números críticos em target nativo quando houver I/O intenso — o próprio README já mostra diferenças em reduction e em kernels curtos.
• Não confunda speedup algorítmico com capacidade absoluta. 2,5× na MX130 ainda pode ser lento demais para o ciclo; 2× na 1660 pode já ser irrelevante frente ao sensor.
• Planeje o port para Jetson como mudança de envelope, não como scp do binário. JetPack, power mode (nvpmodel), clocks (jetson_clocks / jtop) e termal mudam o resultado tanto quanto o kernel.

O que deliberadamente não fazer com este material

• Não inventar timings. Os únicos números de CUDA deste artigo são os do README do c_cuda (MX130 / GTX 1660, WSL e nativo onde citado). Sem novo benchmark, sem novo “X ms”.
• Não vender TensorRT / DeepStream como se estivessem no repo. O caminho conceitual existe; o código progressivo é cuda_1…cuda_8.
• Não tratar speedup de streams do WSL como garantia no Jetson. Memória unificada e ISP/CSI mudam o balanço.
• Não usar o estudo como tutorial de SaaS, cloud GPU rental ou “como montar startup de visão”. É bancada e decisão de envelope.
• Não omitir temperatura e power mode em claims de FPS no Nano — o irmão editorial jetson_nano existe exatamente para isso.

Checklist de bancada antes de declarar “shared memory ganhou”

1. Arch correta (sm_50 / sm_75 / target Jetson) e mesmo tipo numérico nos dois lados.
2. Warmup + média; eventos CUDA — não só relógio de parede no host.
3. Tabela Global vs Shared na mesma GPU e ambiente.
4. Separar tempo de memcpy do tempo de kernel quando o claim for overlap.
5. Anotar se a medição foi WSL ou nativo.
6. Só então discutir port para Jetson (JetPack, nvpmodel, jtop).

Conclusão

O valor do c_cuda para a editorial técnica da Wilson não é “mais um tutorial de soma de vetores”. É um método: progressão host→device→global→shared→medição→reduction→streams, com Makefile consciente de arch e números reais em duas GPUs. Shared memory deixa de ser buzzword e vira decisão mensurável. Streams deixam de ser feature de slides e viram overlap de pipeline. E o caminho até o Jetson fica claro: quem aprendeu a medir memória e latência no desktop está preparado para decidir o que cabe — e o que não cabe — em percepção embarcada sob orçamento de watts.