速查表

// 一维int i = blockIdx.x * blockDim.x + threadIdx.x;if (i < n) { ... } // 二维（x 对应列，即内存连续方向）int col = blockIdx.x * blockDim.x + threadIdx.x;int row = blockIdx.y * blockDim.y + threadIdx.y;if (row < rows && col < cols) m[row * cols + col] = ...; // grid-stride loop：grid 尺寸与数据规模解耦for (int i = blockIdx.x * blockDim.x + threadIdx.x;     i < n; i += gridDim.x * blockDim.x) { ... } // block 内线性 tid（warp 就是按它每 32 个切一组）int tid = threadIdx.x        + threadIdx.y * blockDim.x        + threadIdx.z * blockDim.x * blockDim.y; // 向上取整的 grid 尺寸int blocks = (n + threads - 1) / threads;

cudaMalloc(&d_p, bytes);                 // 注意是 &d_pcudaMallocManaged(&p, bytes);            // 统一内存cudaHostAlloc(&h_p, bytes, cudaHostAllocDefault);   // pinned，异步拷贝的前提cudaMemcpy(dst, src, bytes, cudaMemcpyHostToDevice);cudaMemcpyAsync(dst, src, bytes, kind, stream);cudaMemset(d_p, 0, bytes);cudaFree(d_p);  cudaFreeHost(h_p); cudaStreamCreate(&s);cudaStreamSynchronize(s);                // 只等这条流cudaStreamWaitEvent(s2, evt, 0);         // 跨流依赖，CPU 不阻塞cudaDeviceSynchronize();                 // 等一切，仅用于调试和退出前 cudaEventCreate(&e);  cudaEventRecord(e, s);cudaEventSynchronize(e);cudaEventElapsedTime(&ms, start, stop);  // 唯一正确的 kernel 计时方式 cudaGetLastError();                      // launch 时的同步错误cudaGetErrorString(err);cudaOccupancyMaxPotentialBlockSize(&minGrid, &block, kernel, 0, 0);cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 96*1024);

__inline__ __device__ float warpReduceSum(float v) {    for (int off = 16; off > 0; off >>= 1)        v += __shfl_down_sync(0xffffffff, v, off);    return v;}

# 第一步：看整体时间线，找出真正的瓶颈 kernelnsys profile -o report --stats=true ./app # 第二步：深挖单个 kernelncu --set full -o profile ./appncu --kernel-name myKernel --launch-count 3 ./app # 只看关键指标（快很多）ncu --metrics \  sm__throughput.avg.pct_of_peak_sustained_elapsed,\  gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed,\  sm__warps_active.avg.pct_of_peak_sustained_active \  ./app # 检查访存合并：sectors/requests 理想为 4，接近 32 说明完全没合并ncu --metrics \  l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\  l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum ./app # 检查 bank conflictncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum ./app # 正确性compute-sanitizer --tool memcheck ./appcompute-sanitizer --tool racecheck ./app # 看编译产物cuobjdump -sass ./app | lessnvcc -ptx -arch=sm_80 -o k.ptx k.cu