Compiling with nvcc

nvcc myfile.cu -o myapp               # basic compile + link
nvcc myfile.cu -o myapp -O2           # with optimization
nvcc myfile.cu -o myapp -arch=sm_89   # target Ada Lovelace (RTX 4090)

nvcc myfile.cu -o myapp \
  -gencode arch=compute_80,code=sm_80 \    # Ampere
  -gencode arch=compute_89,code=sm_89 \    # Ada
  -gencode arch=compute_90,code=sm_90 \    # Hopper
  -gencode arch=compute_100,code=sm_100 \  # Blackwell data center (CUDA ≥ 12.8)
  -gencode arch=compute_120,code=sm_120 \  # Blackwell RTX 50 (CUDA ≥ 12.8)
  -gencode arch=compute_120,code=compute_120  # PTX (JIT for future GPUs)

-gencode arch=compute_80,code=sm_80 -gencode arch=compute_80,code=compute_80

# Compile each file to a relocatable device object
nvcc -rdc=true -c fileA.cu -o fileA.o
nvcc -rdc=true -c fileB.cu -o fileB.o

# Link (device link step + host link)
nvcc -rdc=true fileA.o fileB.o -o myapp

# Or explicitly device-link then host-link:
nvcc -rdc=true -dlink fileA.o fileB.o -o device_link.o
g++ fileA.o fileB.o device_link.o -lcudart -o myapp

cmake_minimum_required(VERSION 3.18)
project(MyCUDAProject LANGUAGES CXX CUDA)

set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_ARCHITECTURES "80;89;90;120")  # or "native" for local GPU

add_executable(myapp main.cu kernels.cu)

target_compile_options(myapp PRIVATE
    $<$<COMPILE_LANGUAGE:CUDA>:--use_fast_math -lineinfo>
)

# Separate compilation (for dynamic parallelism)
set_target_properties(myapp PROPERTIES CUDA_SEPARABLE_COMPILATION ON)

nvcc -ptx myfile.cu -o myfile.ptx    # emit PTX
nvcc --ptxas-options=-v myfile.cu    # show register/smem usage

__global__ void inlinePtxDemo(unsigned* out) {
    unsigned tid;
    asm("mov.u32 %0, %%tid.x;" : "=r"(tid));    // read threadIdx.x via PTX
    *out = tid;

    // Atomic add with relaxed ordering
    asm volatile("atom.relaxed.global.add.u32 [%0], 1;" :: "l"(out) : "memory");
}

nvcc myfile.cu -o myapp -lcublas -lcufft -lcurand

CUDA_VISIBLE_DEVICES=0,1 ./myapp    # restrict to GPUs 0 and 1
CUDA_LAUNCH_BLOCKING=1 ./myapp      # synchronize after every kernel (debugging)

#include <nvrtc.h>
#include <cuda.h>

const char* src = R"(
    extern "C" __global__ void myKernel(float* data, int n) {
        int i = blockIdx.x * blockDim.x + threadIdx.x;
        if (i < n) data[i] *= 2.0f;
    }
)";

nvrtcProgram prog;
nvrtcCreateProgram(&prog, src, "myKernel.cu", 0, nullptr, nullptr);

const char* opts[] = {"--gpu-architecture=compute_89"};
nvrtcCompileProgram(prog, 1, opts);

size_t ptxSize;
nvrtcGetPTXSize(prog, &ptxSize);
char* ptx = new char[ptxSize];
nvrtcGetPTX(prog, ptx);

// Load and launch via Driver API
CUmodule module; CUfunction kernel;
cuModuleLoadDataEx(&module, ptx, 0, nullptr, nullptr);
cuModuleGetFunction(&kernel, module, "myKernel");

void* args[] = { &d_data, &n };
cuLaunchKernel(kernel, gridDim, 1, 1, blockDim, 1, 1, 0, 0, args, nullptr);

nvrtcDestroyProgram(&prog);
delete[] ptx;
cuModuleUnload(module);

# Show PTX for a single function
nvcc -ptx -arch=compute_89 myfile.cu && grep -A 50 "myKernel" myfile.ptx

# Show SASS (native assembly)
cuobjdump -sass myapp | grep -A 30 "myKernel"

# Show embedded cubins / PTX in a binary
cuobjdump -all myapp
cuobjdump -ptx  myapp    # extract PTX sections
cuobjdump -elf  myapp    # extract ELF sections

# Register / shared mem usage per kernel
nvcc --ptxas-options=-v myfile.cu 2>&1 | grep "Used"
# Output: ptxas info: Used 32 registers, 4096 bytes smem, 352 bytes cmem[0]