std::atomic<\/code> or a mutex it is correct. This version compiled and ran with g++ using C++17:<\/p>\n#include <atomic>\n#include <iostream>\n#include <thread>\n#include <vector>\n\nint main() {\n const int threads = 4, per = 100000;\n std::atomic<int> counter{0};\n std::vector<std::thread> ts;\n for (int t = 0; t < threads; ++t)\n ts.emplace_back([&] {\n for (int i = 0; i < per; ++i) counter.fetch_add(1);\n });\n for (auto& t : ts) t.join();\n std::cout << counter << \" expected \" << threads * per << \"\\n\";\n}\n\/\/ prints: 400000 expected 400000<\/code><\/pre>\nBe ready for the follow-up: an atomic increment from many threads still makes them contend for one cache line, so a faster design gives each thread a local counter and sums at the end.<\/p>\n
GPU and CUDA basics<\/h2>\n
For GPU roles, learn how the hardware shapes the programming model:<\/p>\n
\n- CPU versus GPU:<\/strong> a CPU has a few cores optimized for low latency on one thread; a GPU has thousands of simpler cores optimized for throughput on many threads running the same program on different data.<\/li>\n
- Thread hierarchy:<\/strong> threads are grouped into blocks, and blocks into a grid. Threads in a block can share fast on-chip memory and synchronize; blocks run independently.<\/li>\n
- Warps:<\/strong> threads execute in groups of 32. When threads in a warp take different branches, the paths run one after another (divergence), which costs time.<\/li>\n
- Memory:<\/strong> global memory is large and slow; shared memory is small and fast; registers are fastest. Neighboring threads reading neighboring addresses lets the hardware combine loads (coalescing). Strided or scattered access wastes bandwidth.<\/li>\n
- Host and device:<\/strong> the CPU launches kernels asynchronously, and data must be on the device (or in unified memory) to be used there.<\/li>\n<\/ul>\n
The standard first kernel adds two arrays. NVIDIA’s official tutorial, “An Even Easier Introduction to CUDA” on the NVIDIA developer blog, builds it step by step with unified memory and the grid-stride loop. A kernel in that style looks like this (it needs the CUDA toolkit and nvcc<\/code> to build):<\/p>\n__global__ void add(int n, float *x, float *y) {\n int index = blockIdx.x * blockDim.x + threadIdx.x;\n int stride = blockDim.x * gridDim.x;\n for (int i = index; i < n; i += stride)\n y[i] = x[i] + y[i];\n}\n\/\/ launch: add<<<numBlocks, 256>>>(n, x, y); then cudaDeviceSynchronize();\n\/\/ numBlocks = (n + 255) \/ 256 covers all elements one per thread.<\/code><\/pre>\nThe index math is the part to practice. Each thread starts at its global index and jumps by the total number of threads in the grid, so every element is handled exactly once whatever the launch size. You can check that claim with a plain Python simulation of the same loop: for n = 1000 with 4 blocks of 256 threads, or with 4 blocks of 64 threads, and for n = 10 with 4 blocks of 256 threads, every element is visited exactly once. For n = 1,000,000 and 256 threads per block, (n + 255) \/ 256<\/code> gives 3,907 blocks.<\/p>\nCommon mistakes: forgetting the bounds check, forgetting to synchronize before the CPU reads results, copying data between host and device more often than needed, and ignoring error codes from CUDA calls.<\/p>\n