Can You Direct a CUDA Kernel to a Specific SM?

1. CUDA Architecture and SM Basics

Streaming Multiprocessors (SMs) are the primary compute units in NVIDIA GPUs, each containing its own register file, shared memory, caches, and many cores. A grid of thread blocks is distributed across multiple SMs, and each SM can run several thread blocks concurrently depending on resource usage such as registers and shared memory.

2. SM Scheduling Mechanism

Warp Scheduler

A warp consists of 32 threads that must execute the same instruction. The warp scheduler selects the next warp to execute and sends it to the dispatch unit. Because each SM has limited warp concurrency—resource constraints require registers and shared memory per thread block—the scheduler’s dynamic policy maximizes resource utilization and avoids idle cycles.

Dispatch Unit

The dispatch unit forwards the selected warp’s instructions to the appropriate functional units within the SM. Each SM contains two dispatch units, allowing two warps to be dispatched simultaneously.

3. CUDA Programming and SM Scheduling Relationship

In CUDA, kernels are launched with

cuLaunchKernel(cuFunction, blockWidth, blockHeight, blockDepth, gridWidth, gridHeight, gridDepth, 0, 0, 0, extra);

. The grid maps to multiple SMs, while each block resides on a single SM and can access that SM’s shared memory. Threads within a block access registers or local memory.

Grid level decides which SMs receive work.

Within an SM, the warp scheduler decides which threads execute at any moment.

4. Can We Specify Scheduling?

The answer is no: both the grid‑to‑SM assignment and the warp‑to‑SM scheduling are performed entirely inside the hardware, and there is no external API to manually bind a kernel to a particular SM. The motivation for such control would be higher utilization, but the hardware already handles placement.

One indirect approach is described in the NanoFlow project, which introduces an intra‑device parallelism engine. It splits a global batch into smaller “nano‑batches” so that operations with different resource bottlenecks (e.g., GEMV vs. GEMM) can overlap on the GPU. A custom execution‑unit scheduling mechanism limits the number of SMs used by each kernel, reducing interference and improving overall resource usage.

For deeper CUDA learning, the article recommends consulting dedicated books and references the original NanoFlow paper (https://zhuanlan.zhihu.com/p/19187838229).