code/day03-setup/devicequery.cuThis is the source used by the lesson and its recorded evidence. Compile commands and expected output live in the directory README.
// SPDX-License-Identifier: MIT
// Day 3: print your GPU's card, then check three numbers you worked out
// from it.
//
// What it does: reads device 0 with cudaGetDeviceProperties, launches one
// kernel that reports the architecture this binary was actually compiled for,
// prints both, then grades the three answers you fill in below.
//
// The two architecture numbers are the point. The device has one compute
// capability and the binary was built for another, and when those two do not
// meet you get "no kernel image is available for execution on the device" at
// launch time with no other clue.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o devicequery devicequery.cu
// Run: ./devicequery
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
// The one error macro. This program is standalone so a reader can copy it into
// Compiler Explorer or a Colab cell, so it carries its own verbatim copy
// instead of including course/check.cuh. `err_` carries a trailing underscore
// so it cannot collide with a variable at the call site.
#define CUDA_CHECK(call) \
do { \
cudaError_t err_ = (call); \
if (err_ != cudaSuccess) { \
std::fprintf(stderr, "CUDA error %s:%d: %s: %s\n", __FILE__, \
__LINE__, #call, cudaGetErrorString(err_)); \
std::exit(EXIT_FAILURE); \
} \
} while (0)
// ---------------------------------------------------------------------------
// Your three answers.
//
// Run the program once with these left at -1. It prints your card and reports
// all three as blank. Work each number out from the card, replace the -1,
// rebuild, run again. A wrong answer prints what you wrote next to what the
// card says and exits non-zero; a blank one does not, because not having done
// the exercise yet is not a failure.
// ---------------------------------------------------------------------------
constexpr int kUnanswered = -1;
// The number you write after sm_ in -arch=sm_XX for this GPU.
constexpr int kArch = kUnanswered;
// How many warps can be resident on one SM at once.
constexpr int kWarpsPerSm = kUnanswered;
// How many threads can be resident on the whole GPU at once.
constexpr int kMaxResidentThreads = kUnanswered;
// ---------------------------------------------------------------------------
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr size_t kSlots = 1; // one int for the kernel to write
// Writes the compute capability this binary was compiled for into out[i].
//
// One thread does the whole job. A warp's 32 addresses are not interesting
// here: only lane 0 of warp 0 passes the guard, and it writes four bytes.
// The launch is a full 256-thread block anyway, so the bounds check is the
// same one every other kernel in this course carries.
//
// Launch assumption: n >= 1, and at least one thread.
//
// __CUDA_ARCH__ exists only during the device compilation pass, and nvcc
// assigns it "a three-digit value string xy0" for compute_xy, so a binary
// built with -arch=sm_75 reports 750.
// snippet: report-arch
__global__ void reportCompiledArch(int* out, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
#if defined(__CUDA_ARCH__)
out[i] = __CUDA_ARCH__;
#else
out[i] = 0;
#endif
}
}
// end snippet
// Prints one graded line. Returns 1 if the answer was filled in and wrong,
// 0 if it was right or still blank.
static int checkAnswer(const char* label, int got, int want,
const char* where) {
if (got == kUnanswered) {
std::printf(" BLANK %-21s work it out from %s\n", label, where);
return 0;
}
if (got != want) {
std::printf(" FAIL %-21s you wrote %d, the card says %d (%s)\n",
label, got, want, where);
return 1;
}
std::printf(" PASS %-21s %d\n", label, got);
return 0;
}
// Formats the packed integer that cudaDriverGetVersion and
// cudaRuntimeGetVersion return. 12060 is 12.6.
static void printVersion(const char* label, int version) {
std::printf(" %-28s %d.%d\n", label, version / 1000,
(version % 1000) / 10);
}
int main() {
// The one CUDA call in this file that is not wrapped in CUDA_CHECK, and
// the reason is this specific lesson: half its readers are here because
// they do not have a GPU yet, and "no CUDA-capable device is detected"
// followed by nothing is the least useful thing we could print them.
int deviceCount = 0;
const cudaError_t countErr = cudaGetDeviceCount(&deviceCount);
if (countErr != cudaSuccess || deviceCount == 0) {
std::fprintf(stderr,
"No CUDA device on this machine (%s).\n"
"Colab, Kaggle T4 x2 and Compiler Explorer all run this "
"program for free.\n",
cudaGetErrorString(countErr));
return EXIT_FAILURE;
}
const int device = 0;
CUDA_CHECK(cudaSetDevice(device));
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
int driverVersion = 0;
int runtimeVersion = 0;
CUDA_CHECK(cudaDriverGetVersion(&driverVersion));
CUDA_CHECK(cudaRuntimeGetVersion(&runtimeVersion));
int* d_arch = nullptr;
CUDA_CHECK(cudaMalloc(&d_arch, kSlots * sizeof(int)));
const int blocks =
static_cast<int>((kSlots + kThreadsPerBlock - 1) / kThreadsPerBlock);
// Launch, then both checks in this order. cudaGetLastError() returns the
// launch configuration error, which on the wrong -arch is 209,
// cudaErrorNoKernelImageForDevice. cudaDeviceSynchronize() returns the
// execution error. A program that checks only after the sync sees nothing.
reportCompiledArch<<<blocks, kThreadsPerBlock>>>(d_arch, kSlots);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
int compiledArch = 0;
CUDA_CHECK(
cudaMemcpy(&compiledArch, d_arch, sizeof(int), cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaFree(d_arch));
const int deviceArch = prop.major * 10 + prop.minor;
const int warpsPerSm = prop.maxThreadsPerMultiProcessor / prop.warpSize;
const int maxResidentThreads =
prop.multiProcessorCount * prop.maxThreadsPerMultiProcessor;
std::printf("Your GPU's card\n");
std::printf(" %-28s %d\n", "CUDA devices visible", deviceCount);
std::printf(" %-28s %s\n", "device 0", prop.name);
std::printf(" %-28s %d.%d (-arch=sm_%d)\n", "compute capability",
prop.major, prop.minor, deviceArch);
std::printf(" %-28s %d (__CUDA_ARCH__)\n", "this binary was built for",
compiledArch);
printVersion("runtime version (nvcc)", runtimeVersion);
printVersion("max CUDA the driver takes", driverVersion);
std::printf(" %-28s %d\n", "SMs", prop.multiProcessorCount);
std::printf(" %-28s %d\n", "warp size", prop.warpSize);
std::printf(" %-28s %d\n", "max threads per block",
prop.maxThreadsPerBlock);
std::printf(" %-28s %d\n", "max threads per SM",
prop.maxThreadsPerMultiProcessor);
std::printf(" %-28s %zu MiB\n", "global memory",
prop.totalGlobalMem / (1024 * 1024));
std::printf(" %-28s %zu KiB\n", "shared mem per block",
prop.sharedMemPerBlock / 1024);
std::printf(" %-28s %zu KiB\n", "shared mem per block, opt-in",
prop.sharedMemPerBlockOptin / 1024);
std::printf(" %-28s %d KiB\n", "L2 cache", prop.l2CacheSize / 1024);
std::printf(" %-28s %d\n", "32-bit registers per block",
prop.regsPerBlock);
std::printf("\nYour answers\n");
int wrong = 0;
wrong += checkAnswer("arch", kArch, deviceArch,
"compute capability, written without the dot");
wrong += checkAnswer("warps per SM", kWarpsPerSm, warpsPerSm,
"max threads per SM and warp size");
wrong += checkAnswer("max resident threads", kMaxResidentThreads,
maxResidentThreads, "SMs and max threads per SM");
if (wrong != 0) {
std::fprintf(stderr, "\n%d of 3 answers wrong.\n", wrong);
return EXIT_FAILURE;
}
return EXIT_SUCCESS;
}