Device Query & Compute Capability

The deviceQuery sample (cuda-samples on GitHub) prints all properties; nvidia-smi -q shows driver-level info.

Compute capability (CC) = architecture version (major.minor), determining available features:

P100, GTX 10xx A100, RTX 30xx, Jetson Orin L4, L40, RTX 40xx B200, GB200, RTX 50xx

int count;
cudaGetDeviceCount(&count);
for (int d = 0; d < count; ++d) {
    cudaDeviceProp p;
    cudaGetDeviceProperties(&p, d);
    printf("Device %d: %s\n", d, p.name);
    printf("  Compute capability: %d.%d\n", p.major, p.minor);
    printf("  SMs: %d, warp size: %d\n", p.multiProcessorCount, p.warpSize);
    printf("  Global memory: %.1f GB\n", p.totalGlobalMem / 1e9);
    printf("  Shared memory per block: %zu KB (opt-in max %zu KB)\n",
           p.sharedMemPerBlock / 1024, p.sharedMemPerBlockOptin / 1024);
    printf("  Registers per SM: %d, max threads per block: %d, per SM: %d\n",
           p.regsPerMultiprocessor, p.maxThreadsPerBlock, p.maxThreadsPerMultiProcessor);
    printf("  L2 cache: %d MB, memory bus: %d-bit\n", p.l2CacheSize >> 20, p.memoryBusWidth);
}
cudaSetDevice(0);
int value; cudaDeviceGetAttribute(&value, cudaDevAttrMultiProcessorCount, 0);   // faster single attribute

🧪 Part I Exercises

Run the hello kernel with different grid/block shapes (e.g., <<<3, 5>>>, dim3(2,2)) and explain the print order.

Write a kernel saxpy(n, a, x, y) computing y = a*x + y; verify against a CPU loop.

Convert vecAdd to a grid-stride loop and launch it with 4 × SM count blocks.

Write a 2-D kernel that converts an RGB image (interleaved bytes) to grayscale.

Add CUDA_CHECK to all calls and deliberately trigger three different errors (too many threads, bad pointer, wrong size).

Write a device query program printing name, CC, SM count, memory, and shared memory per block.

Time vecAdd on CPU vs GPU (including and excluding transfers) for sizes 1K to 100M; plot the results.

Write a __host__ __device__ function clamp and use it in both a CPU loop and a kernel.

Implement element-wise matrix multiplication (Hadamard product) with 2-D indexing.

Write a kernel that initializes an array with its global index and verify on the host.