Cheat sheet

// 1Dint i = blockIdx.x * blockDim.x + threadIdx.x;if (i < n) { ... } // 2D (x is the column, the contiguous memory direction)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 size is decoupled from data sizefor (int i = blockIdx.x * blockDim.x + threadIdx.x;     i < n; i += gridDim.x * blockDim.x) { ... } // Linear tid in the block (warps slice this into groups of 32)int tid = threadIdx.x        + threadIdx.y * blockDim.x        + threadIdx.z * blockDim.x * blockDim.y; // Ceiling division for the grid sizeint blocks = (n + threads - 1) / threads;

cudaMalloc(&d_p, bytes);                 // note the &d_pcudaMallocManaged(&p, bytes);            // unified memorycudaHostAlloc(&h_p, bytes, cudaHostAllocDefault);   // pinned; required for async copiescudaMemcpy(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);                // wait on this stream onlycudaStreamWaitEvent(s2, evt, 0);         // cross-stream dependency; CPU does not blockcudaDeviceSynchronize();                 // wait for everything; debug and shutdown only cudaEventCreate(&e);  cudaEventRecord(e, s);cudaEventSynchronize(e);cudaEventElapsedTime(&ms, start, stop);  // the only correct way to time a kernel cudaGetLastError();                      // synchronous error from the launchcudaGetErrorString(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;}

# Step 1: inspect the overall timeline and find the real bottleneck kernelnsys profile -o report --stats=true ./app # Step 2: drill into a single kernelncu --set full -o profile ./appncu --kernel-name myKernel --launch-count 3 ./app # Key metrics only (much faster)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 # Check coalescing: sectors/requests ideal is 4; near 32 means no coalescing at allncu --metrics \  l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\  l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum ./app # Check bank conflictncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum ./app # Correctnesscompute-sanitizer --tool memcheck ./appcompute-sanitizer --tool racecheck ./app # Inspect the compiler outputcuobjdump -sass ./app | lessnvcc -ptx -arch=sm_80 -o k.ptx k.cu