PTX, SASS & NVRTC

Embedding PTX (-gencode arch=compute_90,code=compute_90) lets future GPUs JIT-compile your code.

34.1 PTX and SASS

PTX — NVIDIA's virtual assembly (forward-compatible, JIT-compiled by the driver).

SASS — actual machine code for a specific architecture.

Embedding PTX (-gencode arch=compute_90,code=compute_90) lets future GPUs JIT-compile your code.

nvcc -ptx kernel.cu -o kernel.ptx
cuobjdump -sass app               # disassemble SASS from a binary
nvdisasm kernel.cubin
cuobjdump -ptx app

34.2 Inline PTX

__device__ unsigned laneId() {
    unsigned id;
    asm volatile("mov.u32 %0, %%laneid;" : "=r"(id));
    return id;
}
__device__ float fastAdd(float a, float b) {
    float r;
    asm("add.f32 %0, %1, %2;" : "=f"(r) : "f"(a), "f"(b));
    return r;
}

34.3 Runtime compilation with NVRTC

Used for specializing kernels at runtime (constants baked in), code generation frameworks, and Python JITs. nvJitLink links LTO-IR at runtime.

#include <nvrtc.h>
#include <cuda.h>
const char* src = R"(
extern "C" __global__ void axpy(float a, float* x, float* y, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
})";
nvrtcProgram prog;
nvrtcCreateProgram(&prog, src, "axpy.cu", 0, nullptr, nullptr);
const char* opts[] = {"--gpu-architecture=compute_80"};
nvrtcCompileProgram(prog, 1, opts);
size_t ptxSize; nvrtcGetPTXSize(prog, &ptxSize);
std::string ptx(ptxSize, '\0'); nvrtcGetPTX(prog, ptx.data());
CUmodule mod; CUfunction fn;
cuModuleLoadData(&mod, ptx.c_str());
cuModuleGetFunction(&fn, mod, "axpy");
void* args[] = {&a, &d_x, &d_y, &n};
cuLaunchKernel(fn, blocks, 1, 1, threads, 1, 1, 0, nullptr, args, nullptr);