Hello CUDA Revisited

The configurable battalion — feel block × thread in your fingertips

The earlier <<<1, 4>>> launch is a single squad of four ants. Realistic launches have hundreds of blocks each holding hundreds of threads. The expanded version below takes block count and thread count from the command line so you can dial any formation and watch what happens:

Code

ant_battalion.cu — variable launch geometry·cuda

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

__global__ void ant_battalion_report() {
    int ant_id = blockIdx.x * blockDim.x + threadIdx.x;
    printf("Ant %d reporting from block %d, slot %d.\n",
           ant_id, blockIdx.x, threadIdx.x);
}

int main(int argc, char **argv) {
    if (argc != 3) {
        fprintf(stderr, "Usage: %s <num_blocks> <threads_per_block>\n", argv[0]);
        return 1;
    }
    int blocks  = atoi(argv[1]);
    int threads = atoi(argv[2]);
    if (threads % 32 != 0) {
        fprintf(stderr, "Tip: threads-per-block should be a multiple of 32 (warp size).\n");
    }

    ant_battalion_report<<<blocks, threads>>>();
    cudaDeviceSynchronize();
    return 0;
}

Build + try a few formations·bash

nvcc -arch=sm_89 ant_battalion.cu -o ant_battalion

./ant_battalion 1   4    # 1 squad of 4 — same as the original hello
./ant_battalion 2   8    # 2 squads, 16 ants total — order interleaves
./ant_battalion 32  32   # 1024 ants — first hint of scheduler chaos
./ant_battalion 128 4    # 512 ants split into many tiny squads

# Notice:
#  1) Ants don't print in numerical order — scheduler picks blocks freely.
#  2) WITHIN a block, threadIdx values are usually grouped.
#  3) printf throughput drops sharply past ~10K threads — buffer fills.

Exercise

Compile ant_battalion.cu and run it with ./ant_battalion 4 32, ./ant_battalion 32 32, and ./ant_battalion 256 256. For each, count how many lines you get (should equal blocks × threads) and notice how 'random' the order looks at higher counts. Then redirect to file: ./ant_battalion 256 256 > out.txt and grep for one specific ant_id — confirm the GPU is accurate even when output ordering is chaotic.

The configurable battalion — feel block × thread in your fingertips

Code

External links

Exercise

Progress

Comments 0