Hello, CUDA!
Configuring Ubuntu 26.04 LTS for GPU computing: From fresh OS install to running your first nvcc benchmark
I can’t believe my stroke of luck, I actually managed to nab a Nvidia 5090 Founders Edition at retail price! I’ve been pretty curious about the limits of throughput performance that is causing such a shortage of every type of computer hardware possible today (RAM, SSDs, Flash, CPUs, and of course GPUs). Who would have thought the price of computer parts would increase by such a factor! And all because these AI companies complain about not enough compute, even though they’re furiously buying up all the compute.
I had to dive in. But I didn’t really want to rent a GPU online. So I waited my turn, until I landed a consumer-grade one that wasn’t scalped (b/c I refuse to support scalpers). Plus, I wanted one for gaming.
After building the computer, setting up the environment, and running some CUDA code, I found the initial benchmark results to be promising! But, there’s a wrong way to use CUDA that severely handicaps performance, so I’ll demonstrate and explain.
Notes to self:
* this post is about the basics of performance benchmarking.
* will need to introduce some basic terms, like SIMD and SIMT
* will need really good diagrams for blocks/threads/warps
* vectorization
* performance
* gpu: naive vs efficient, speedup 100x
* cpu is better than naive gpu
* is it possible for cpu to be worse than naive gpu when there are multiple operations?
* this post assumes knowledge of programming and of c++.
* makefile. makes it easier to compile. like a cheatsheet so you don't have to remember all the argument options.Code
#include <iostream>
#include <cuda_runtime.h>
// -----------------------------------------------------------------------------
// KERNEL 1: Naive Fill
// Uses a sub-optimal 1D layout (1 thread per block) which underutilizes
// GPU warps and Streaming Multiprocessors (SMs).
// -----------------------------------------------------------------------------
__global__ void naiveFillKernel(float* data, int N) {
int i = blockIdx.x; // only using block index
if (i < N) {
data[i] = 1.0f;
}
}
// -----------------------------------------------------------------------------
// KERNEL 2: Efficient Fill
// Uses proper 2D/1D grid-block stride indexing (256 threads per block)
// to fully saturate GPU compute pipelines and memory bandwidth.
// -----------------------------------------------------------------------------
__global__ void efficientFillKernel(float* data, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
data[i] = 1.0f;
}
}
float timeKernel(void (*kernel)(float*, int), dim3 grid, dim3 block, float* d_data, int N) {
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
kernel<<<grid, block>>>(d_data, N);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms = 0;
cudaEventElapsedTime(&ms, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return ms;
}
int main() {
const int N = 10000000; // 10 million
size_t size = N * sizeof(float);
std::cout << "Running CUDA Hello World with " << N << " elements..." << std::endl;
// =========================================================================
// APPROACH 1: NAIVE WAY
// Slow Host-to-Device transfer + terrible thread block geometry (1 thread/block)
// =========================================================================
float* h_data_naive = (float*)malloc(size);
float* d_data_naive;
cudaMalloc((void**)&d_data_naive, size);
cudaMemcpy(d_data_naive, h_data_naive, size, cudaMemcpyHostToDevice);
dim3 naiveBlock(1);
dim3 naiveGrid(N);
float naiveTime = timeKernel(naiveFillKernel, naiveGrid, naiveBlock, d_data_naive, N);
cudaMemcpy(h_data_naive, d_data_naive, size, cudaMemcpyDeviceToHost);
cudaFree(d_data_naive);
free(h_data_naive);
// =========================================================================
// APPROACH 2: EFFICIENT WAY
// Uses Managed Memory + optimal 256 threads per block (32-thread warp aligned)
// =========================================================================
float* h_data_eff = (float*)malloc(size);
float* d_data_eff;
cudaMalloc((void**)&d_data_eff, size);
cudaMemcpy(d_data_eff, h_data_eff, size, cudaMemcpyHostToDevice);
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
dim3 effBlock(threadsPerBlock);
dim3 effGrid(blocksPerGrid);
float effTime = timeKernel(efficientFillKernel, effGrid, effBlock, d_data_eff, N);
cudaMemcpy(h_data_eff, d_data_eff, size, cudaMemcpyDeviceToHost);
cudaFree(d_data_eff);
free(h_data_eff);
// =========================================================================
// PRINT RESULTS
// =========================================================================
std::cout << "----------------------------------------------------" << std::endl;
std::cout << "Naive Execution Time : " << naiveTime << " ms" << std::endl;
std::cout << "Efficient Execution Time: " << effTime << " ms" << std::endl;
std::cout << "Speedup Factor : " << naiveTime / effTime << "x" << std::endl;
std::cout << "----------------------------------------------------" << std::endl;
return 0;
}Results:
Running CUDA Hello World with 10000000 elements...
----------------------------------------------------
Naive Execution Time : 4.79638 ms
Efficient Execution Time: 0.037664 ms
Speedup Factor : 127.347x
----------------------------------------------------Appendix A: Building the PC
Add in pictures…
And list the parts…
a
a
a
a
a
a
a
a
a
a
Appendix B: Environment Setup
Ubuntu 26.04 LTS.
So, I built the computer, installed Linux Ubuntu 26.04LTS,
install nvidia drivers, like nvcc.
a
a
a
a
a
a
a
a
a
a
prereqs:
References
https://developer.nvidia.com/blog/even-easier-introduction-cuda/
description options
Demonstrating >100x performance leap from CPU to GPU through blocks, threads, and parallel processing
GPU setup and initial benchmark
From fresh OS install to running your first nvcc benchmark
Configuring Ubuntu 26.04 LTS for GPU computing: From fresh OS install to running your first nvcc benchmark