cudaMallocPitch 分配的内存必须用 pitch 计算行首地址再访问，不可直接按宽度索引；绑定纹理后应通过 tex2D 等采样函数读取，禁止指针写入；更新需 cudaMemcpy2D 配合重建纹理对象或使用可写纹理。
cudaMallocPitch 分配的纹理内存怎么用指针访问
不能直接用普通指针解引用访问 cudaMallocPitch 分配的内存——它的行对齐（pitch）不等于图像宽度 × 元素大小，硬算偏移会越界或读错行。必须用 pitch 值参与地址计算。
- 假设分配了 1024×768 的
float纹理：float *d_data; size_t pitch; cudaMallocPitch(&d_data, &pitch, 1024 * sizeof(float), 768);
- 在 kernel 中访问第
y行、第x列元素，正确写法是：float *row = (float*)((char*)d_data + y * pitch); float val = row[x];
- 错误写法（忽略 pitch）：
d_data[y * 1024 + x]—— 在大多数 GPU 上会跨行错位，尤其当 1024×sizeof(float) 不是 128 字节对齐时
绑定到纹理对象后还能不能用指针读写
可以读，但不能安全写；且必须通过 tex2D / tex3D 等采样函数，不能用原始指针解引用。
- 纹理内存绑定后，GPU 会启用缓存和插值逻辑，底层物理地址可能被重映射，
cudaMemcpy或 kernel 直接写d_data地址会绕过纹理缓存一致性机制，导致读取结果不可预测 - 若需更新纹理内容：先用
cudaMemcpy2D把新数据拷回 pitch 分配的内存，再调用cudaTexObject_t对应的更新接口（如重新cudaCreateTextureObject），或使用 CUDA 11.0+ 的cudaTexObject_t可写纹理（需设备支持 compute capability ≥ 7.0 且显存为 unified memory） - 常见误操作：
tex2D读取正常，但 kernel 里同时用d_data[y * pitch + x]写入 —— 这不是原子同步行为，读写结果不一致
__restrict__ 和纹理指针要不要加
对指向 cudaMallocPitch 内存的指针加 __restrict__ 有用；但对纹理采样器对象（cudaTextureObject_t）本身不适用。
C++ Code Review Master
下载
组合式C++代码评审方案，融合静态分析、AI推理、多轮迭代评审和C++专项检查，适用于PR审查、增量代码审查、全项目评审和代码质量评分，触发词包括review cpp、cpp代码评审、C++review、代码审查。
-
__restrict__告诉编译器该指针不与其他指针别名，能提升 kernel 中基于 pitch 手动寻址的访存优化程度（比如循环展开、向量化） - 示例：
__global__ void copy_kernel(float * __restrict__ dst, const float * __restrict__ src, size_t src_pitch, size_t dst_pitch, int w, int h) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < w && y < h) { float *src_row = (float*)((char*)src + y * src_pitch); float *dst_row = (float*)((char*)dst + y * dst_pitch); dst_row[x] = src_row[x]; } } - 不加
__restrict__时，编译器可能因担心指针重叠而禁用某些优化，实测在密集图像处理中性能可差 15%–30%
cudaMallocArray 和 cudaMallocPitch 哪个更适合纹理
优先选 cudaMallocArray；除非你要做非纹理用途的 pitch 对齐内存（比如双缓冲帧拷贝）。
立即学习“C++免费学习笔记（深入）”；
-
cudaMallocArray分配的是“纹理原生格式”，绑定时无需额外配置 format、numChannels，且 GPU 硬件能直接理解其布局，采样延迟更低 -
cudaMallocPitch是通用对齐内存，绑定纹理前必须用cudaChannelFormatDesc显式描述数据类型（如cudaCreateChannelDesc<float>()</float>），否则cudaBindTexture2D会失败并返回cudaErrorInvalidValue - 兼容性注意：老驱动（cudaMallocArray 的 32-bit 整数纹理支持不稳定；若必须用
int4格式，建议 fallback 到cudaMallocPitch+ 手动tex2D<int4></int4>
实际项目里最容易被忽略的是 pitch 对齐带来的跨平台差异：同一段 kernel 在 A100 上可能因默认对齐到 512 字节而侥幸跑通，在 GTX 1060 上因只对齐到 256 字节就出现每行末尾几个像素全零——别依赖“看起来正常”，务必用 cudaMemcpy2D 验证拷入/拷出的数据完整性。
