# 简单编译命令
nvcc my_program.cu -o my_program

.cu源文件
    ↓
预处理阶段
    ↓
代码分离
    ├── 主机代码 → 主机编译器 → .o文件
    └── 设备代码 → NVCC编译器 → PTX代码 → SASS代码 → .cubin文件
    ↓
设备代码链接
    ↓
主机链接
    ↓
可执行文件

# 预处理CUDA文件
nvcc -E my_program.cu -o my_program.i

// sample.cu
#include <iostream>
#include <cuda_runtime.h>

// 主机函数
void hostFunction() {
    std::cout << "Host function" << std::endl;
}

// 设备函数
__device__ float deviceFunction(float x) {
    return x * 2.0f;
}

// 全局函数
__global__ void kernel(float* data) {
    int idx = threadIdx.x;
    data[idx] = deviceFunction(data[idx]);
}

int main() {
    hostFunction();
    // 主机代码继续...
}

# 生成PTX中间代码
nvcc -ptx my_program.cu -o my_program.ptx

# 使用ptxas优化PTX代码
ptxas my_program.ptx -o my_program.sass

# 生成CUDA二进制文件
nvcc -cubin my_program.cu -o my_program.cubin

# 查看主机编译命令
nvcc --dry-run my_program.cu -o my_program

# 链接阶段
nvcc my_program.o -o my_program -lcudart

// 示例PTX代码
.version 7.8
.target sm_70
.address_size 64

.visible .entry addKernel(
    .param .u64 addKernel_param_0,
    .param .u64 addKernel_param_1,
    .param .u64 addKernel_param_2
)
{
    .reg .b64   %rd<3>;
    .reg .f32   %f<4>;
    .reg .b32   %r<6>;

    ld.param.u64    %rd1, [addKernel_param_0];
    ld.param.u64    %rd2, [addKernel_param_1];
    cvta.to.global.u64  %rd1, %rd1;
    cvta.to.global.u64  %rd2, %rd2;

    // 计算线程ID
    mov.u32     %r1, %tid.x;
    mov.u32     %r2, %ctaid.x;
    mov.u32     %r3, %ntid.x;
    mad.lo.s32  %r4, %r2, %r3, %r1;
    cvt.s64.s32 %rd2, %r4;

    // 内存访问和计算
    mul.wide.s32  %rd3, %r4, 4;
    add.s64    %rd1, %rd1, %rd3;
    add.s64    %rd2, %rd2, %rd3;
    ld.global.f32  %f1, [%rd1];
    ld.global.f32  %f2, [%rd2];
    add.f32    %f3, %f1, %f2;

    // 存储结果
    ld.param.u64    %rd1, [addKernel_param_2];
    cvta.to.global.u64  %rd1, %rd1;
    add.s64    %rd2, %rd1, %rd3;
    st.global.f32  [%rd2], %f3;

    ret;
}

# 查看SASS代码
nvdisasm my_program.cubin > my_program.sass

# 编译但不链接
nvcc -c my_program.cu -o my_program.o

# 生成可执行文件
nvcc my_program.cu -o my_program

# 生成静态库
nvcc -lib my_program.cu -o libmyprogram.a

# 指定虚拟架构（生成PTX）
nvcc -arch=compute_70 my_program.cu

# 指定真实架构（生成SASS）
nvcc -arch=sm_70 my_program.cu

# 同时生成多种架构代码
nvcc -gencode arch=compute_60,code=sm_60 \
     -gencode arch=compute_70,code=sm_70 \
     -gencode arch=compute_80,code=sm_80 \
     my_program.cu

# 向前兼容（为未来架构生成PTX）
nvcc -gencode arch=compute_70,code=compute_70 my_program.cu

# 主机代码优化级别
nvcc -O0 my_program.cu    # 无优化
nvcc -O1 my_program.cu    # 基本优化
nvcc -O2 my_program.cu    # 标准优化（默认）
nvcc -O3 my_program.cu    # 高级优化
nvcc -Ofast my_program.cu # 最快优化（可能不严格遵循标准）

# 设备代码优化
nvcc -dopt=on my_program.cu   # 启用设备代码优化
nvcc -dopt=off my_program.cu  # 禁用设备代码优化

# 特定优化类型
nvcc --use_fast_math my_program.cu  # 使用快速数学函数
nvcc --ftz=true my_program.cu       # 清零非正规数
nvcc --prec-div=false my_program.cu # 非精确除法

# 生成调试信息
nvcc -g -G my_program.cu  # -g为主机代码，-G为设备代码

# 行号信息（用于profiling）
nvcc -lineinfo my_program.cu

# 禁用优化以获得更好的调试体验
nvcc -O0 -g -G my_program.cu

# 设备端调试
nvcc -device-debug my_program.cu
nvcc -expt-relaxed-constexpr my_program.cu

# 生成PTX代码
nvcc -ptx my_program.cu -o my_program.ptx

# 生成CUBIN文件
nvcc -cubin my_program.cu -o my_program.cubin

# 生成FATBIN（多种架构的二进制包）
nvcc -fatbin my_program.cu -o my_program.fatbin

# 保留中间文件
nvcc -keep my_program.cu  # 保留所有中间文件
nvcc -keep-dir temp my_program.cu  # 指定中间文件目录

# 添加头文件搜索路径
nvcc -I/usr/local/cuda/include my_program.cu

# 添加库文件搜索路径
nvcc -L/usr/local/cuda/lib64 my_program.cu

# 链接特定库
nvcc -lcudart -lcublas my_program.cu

# 预定义宏
nvcc -DMY_MACRO=1 my_program.cu

# 并行编译（需要支持的编译器）
nvcc -t 4 my_program.cu  # 使用4个线程编译

# 预编译头文件
nvcc -use_fast_math -Xcompiler -fopenmp my_program.cu

# 交叉编译到不同平台
nvcc -m64 my_program.cu     # 64位目标
nvcc -m32 my_program.cu     # 32位目标

// 运行时查询计算能力
int device;
cudaGetDevice(&device);
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, device);
printf("Compute Capability: %d.%d\n", prop.major, prop.minor);

# 生成PTX，运行时JIT编译
nvcc -arch=compute_70 -code=compute_70 my_program.cu

# 支持多种当前架构和未来架构
nvcc -gencode arch=compute_60,code=sm_60 \
     -gencode arch=compute_70,code=sm_70 \
     -gencode arch=compute_75,code=sm_75 \
     -gencode arch=compute_80,code=compute_80 \
     -gencode arch=compute_80,code=compute_80 \
     my_program.cu

# CMakeLists.txt示例
set(CMAKE_CUDA_ARCHITECTURES 60 70 75 80)
set(CMAKE_CUDA_ARCHITECTURE_ALL ON)
set(CMAKE_CUDA_ARCHITECTURE_VALUES OFF)

# 查询设备属性
nvcc --device-c my_program.cu  # 仅编译设备代码

# 运行时检测最佳架构
deviceQuery  # NVIDIA提供的示例程序

// 告诉编译器指针不重叠
__global__ void vectorAdd(const float* __restrict__ a,
                         const float* __restrict__ b,
                         float* __restrict__ c,
                         int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

// 使用快速内置函数
__global__ void fastMath(float* data) {
    int idx = threadIdx.x;
    // 使用更快的数学函数
    data[idx] = __fmul_rn(data[idx], data[idx]);  // 快速乘法
    data[idx] = __fdiv_rn(data[idx], 2.0f);      // 快速除法
}

// 告诉编译器循环不依赖迭代
#pragma unroll
for (int i = 0; i < 4; i++) {
    sum += data[i];
}

// 控制循环展开
#pragma unroll 8
for (int i = 0; i < 16; i++) {
    // 编译器会展开前8次迭代
}

# 使用快速数学函数（可能牺牲精度）
nvcc --use_fast_math my_program.cu

# 等效于以下选项组合
nvcc --ftz=true --prec-div=false \
     --prec-sqrt=false --fmad=true \
     my_program.cu

// 强制内联
__device__ __forceinline__ float fastFunc(float x) {
    return x * x;
}

// 不内联
__device__ __noinline__ float slowFunc(float x) {
    // 复杂计算
}

// 确保内存对齐
struct __align__(16) Float4 {
    float x, y, z, w;
};

// 使用矢量化类型
__global__ void vectorizedLoad(float4* input) {
    float4 data = input[threadIdx.x];
    // 处理4个float
}

// 将只读数据放入常量内存
__constant__ float constData[1024];

// 在设备代码中使用
__global__ void kernel() {
    float value = constData[threadIdx.x];  // 快速访问
}

# 编译设备代码为对象文件
nvcc -dc device_code1.cu -o device_code1.o
nvcc -dc device_code2.cu -o device_code2.o

# 设备代码链接
nvcc -dlink device_code1.o device_code2.o -o device_link.o

# 主机链接
nvcc device_code1.o device_code2.o device_link.o host_code.o -o final_program

// 使用C++14/17特性
template<typename T>
__global__ void modernKernel(T* data) {
    // C++14: 泛型lambda
    auto process = [](auto x) { return x * 2; };

    // C++17: 结构化绑定
    auto [idx, value] = std::make_tuple(threadIdx.x, data[threadIdx.x]);
    data[idx] = process(value);
}

# 链接CUB库
nvcc my_program.cu -lcub

# 包含Thrust头文件
nvcc -I/usr/local/cuda/include my_program.cu

# 生成详细的编译信息
nvcc -v my_program.cu

# 生成PTX和SASS用于分析
nvcc -ptx -keep my_program.cu

# 生成二进制信息
nvcc -Xptxas=-v my_program.cu

# 生成profiling信息
nvcc -lineinfo my_program.cu

# 启用NVTX标记
nvcc -Xcompiler -DNVTX_PROFILING my_program.cu

# 使用compute sanitizer
nvcc -g -G my_program.cu
compute-sanitizer ./my_program

# Makefile示例
CUDA_DIR = /usr/local/cuda
NVCC = $(CUDA_DIR)/bin/nvcc

# 编译选项
ARCH = -gencode arch=compute_70,code=sm_70
OPT = -O3 -use_fast_math
DEBUG = -g -G
INC = -I$(CUDA_DIR)/include

# 目标文件
CU_SRCS = kernel1.cu kernel2.cu
CU_OBJS = $(CU_SRCS:.cu=.o)

all: my_program

%.o: %.cu
	$(NVCC) $(ARCH) $(OPT) $(INC) -c $< -o $@

my_program: $(CU_OBJS) host.o
	$(NVCC) $(ARCH) $^ -o $@

clean:
	rm -f *.o my_program

# CMakeLists.txt
cmake_minimum_required(VERSION 3.18)
project(MyCudaProject LANGUAGES CXX CUDA)

# 启用CUDA
enable_language(CUDA)

# 设置CUDA架构
set(CMAKE_CUDA_ARCHITECTURES 70 75 80)

# 查找CUDA
find_package(CUDA REQUIRED)

# 添加可执行文件
add_executable(my_program main.cu kernel.cu)

# 设置编译选项
target_compile_options(my_program PRIVATE
    $<$<COMPILE_LANGUAGE:CUDA>:-O3 -use_fast_math>
)

# 链接库
target_link_libraries(my_program PRIVATE
    ${CUDA_CUDART_LIBRARY}
    ${CUDA_CUBLAS_LIBRARY}
)

# setup.py
from pybind11.setup_helpers import Pybind11Extension, build_ext
from pybind11 import get_cmake_dir
import pybind11

ext_modules = [
    Pybind11Extension(
        "cuda_module",
        [
            "bindings/main.cpp",
            "src/kernel.cu",
        ],
        include_dirs=[
            pybind11.get_include(),
            "/usr/local/cuda/include",
        ],
        libraries=["cudart"],
        extra_compile_args=["-O3", "-use_fast_math"],
        cxx_std=17,
    ),
]

setup(
    ext_modules=ext_modules,
    cmdclass={"build_ext": build_ext},
)

# 快速编译，便于调试
nvcc -O0 -g -G -lineinfo my_program.cu -o my_program_dev

# 标准优化，保留调试信息
nvcc -O2 -g -lineinfo my_program.cu -o my_program_test

# 最高优化，去除调试信息
nvcc -O3 -use_fast_math my_program.cu -o my_program_release

# 检查编译器返回值
if ! nvcc my_program.cu -o my_program; then
    echo "Compilation failed"
    exit 1
fi

# 使用set -e在脚本中
set -e  # 遇到错误立即退出
nvcc my_program.cu -o my_program

# 记录CUDA版本
nvcc --version > cuda_version.txt

# 在代码中检查CUDA版本
#if CUDA_VERSION >= 11000
    // 使用CUDA 11.0+特性
#endif

# 解决方案：明确指定头文件路径
nvcc -I/usr/local/cuda/include my_program.cu

# 解决方案：指定库路径和库名
nvcc -L/usr/local/cuda/lib64 -lcudart my_program.cu

# 检查GPU架构
nvidia-smi

# 使用正确的架构编译
nvcc -arch=sm_70 my_program.cu

# 查看预处理后的代码
nvcc -E my_program.cu > my_program.i

# 生成PTX
nvcc -ptx my_program.cu

# 生成SASS
nvcc -cubin my_program.cu
nvdisasm my_program.cubin

# 查看完整的编译命令
nvcc --dry-run my_program.cu