学习目标学完本节你将能够理解常量内存的物理特性与广播机制掌握常量内存的最佳使用场景及限制了解纹理内存的缓存优化与硬件插值能力使用纹理对象进行基本的图像采样对比常量内存、纹理内存与全局内存在只读场景下的性能差异在卷积、图像处理等场景中合理选择只读缓存优化手段1. 常量内存Constant Memory1.1 物理特性常量内存是位于设备内存上的一块只读区域但拥有独立的片上常量缓存Constant Cache每 SM 通常有 64 KB 的常量缓存。总容量64 KB所有 SM 共享访问延迟同一 Warp 内所有线程读取相同地址时广播机制使访问速度接近寄存器。如果同一 Warp 内线程读取不同地址访问会串行化性能急剧下降。生命周期从 Host 端通过cudaMemcpyToSymbol写入后在整个程序运行期间有效但 Kernel 内只能读不能写。1.2 广播机制当同一 Warp 内的 32 个线程都读取常量内存的同一个地址时硬件会将数据广播给所有线程相当于一次读操作服务整个 Warp这是常量内存最高效的使用方式。适用场景算法参数如缩放系数、阈值滤波器系数固定查找表如果访问模式统一不适用场景每个线程读取不同的常量数组元素例如const_arr[tid % 64]这会导致同一 Warp 内访问多个地址冲突严重。数据量超过 64 KB 的只读数组。1.3 使用方式// 声明常量内存 __constant__ float scaleFactor; __constant__ float filterCoeffs[16]; // Host 端初始化 float scale 2.5f; cudaMemcpyToSymbol(scaleFactor, scale, sizeof(float)); float coeffs[16] {...}; cudaMemcpyToSymbol(filterCoeffs, coeffs, 16 * sizeof(float)); // Kernel 内使用 __global__ void applyFilter(float *in, float *out, int N) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx N) { float val in[idx] * scaleFactor; // 使用 filterCoeffs 进行计算... out[idx] val; } }2. 常量内存 vs 全局内存性能对比下面通过一个简单示例对比常量内存和全局内存在所有线程读取相同参数时的性能差异。#include cstdio #include cuda_runtime.h __constant__ float c_scale; // 常量内存 // 使用常量内存 __global__ void useConstant(float *in, float *out, int N) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx N) { out[idx] in[idx] * c_scale; } } // 使用全局内存参数作为 Kernel 参数传入 __global__ void useGlobalParam(float *in, float *out, float scale, int N) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx N) { out[idx] in[idx] * scale; } } int main() { const int N 1 22; size_t bytes N * sizeof(float); float *d_in, *d_out; cudaMalloc(d_in, bytes); cudaMalloc(d_out, bytes); // 初始化省略... float scale 2.5f; cudaMemcpyToSymbol(c_scale, scale, sizeof(float)); cudaEvent_t start, end; cudaEventCreate(start); cudaEventCreate(end); float ms; // 常量内存版本 cudaEventRecord(start); useConstant(N255)/256, 256(d_in, d_out, N); cudaEventRecord(end); cudaEventSynchronize(end); cudaEventElapsedTime(ms, start, end); printf(Constant memory: %f ms\n, ms); // 全局参数版本 cudaEventRecord(start); useGlobalParam(N255)/256, 256(d_in, d_out, scale, N); cudaEventRecord(end); cudaEventSynchronize(end); cudaEventElapsedTime(ms, start, end); printf(Global parameter: %f ms\n, ms); cudaFree(d_in); cudaFree(d_out); return 0; }**预期结果**两者性能非常接近因为编译器通常会将 Kernel 参数保存在常量缓存中。但关键区别在于容量和可更新性常量内存可以存储较大的只读参数数组如 64 KB 的系数表而 Kernel 参数数量有限。3. 纹理内存Texture Memory3.1 什么是纹理内存纹理内存是 GPU 中专门优化的只读数据访问路径通过纹理缓存Texture Cache提供二维空间局部性优化。底层数据仍存储在全局内存但通过纹理缓存访问。支持硬件自动处理边界clamp, wrap, mirror支持硬件双线性插值、归一化坐标等缓存针对二维局部性优化适合图像处理、卷积等场景3.2 纹理内存与 L1/L2 缓存的区别常规全局内存访问经过 L1/L2 缓存对一维连续访问优化较好。纹理缓存针对二维空间邻近进行优化当线程访问的像素在图像上邻近时缓存命中率很高。纹理缓存是只读的因此不能用于写入。3.3 纹理对象Texture Object的基本用法现代 CUDA 推荐使用纹理对象它在 Kernel 中以cudaTextureObject_t类型传递。#include cuda_runtime.h // Kernel 使用纹理对象 __global__ void textureSampleKernel(cudaTextureObject_t texObj, float *out, int width, int height) { int x blockIdx.x * blockDim.x threadIdx.x; int y blockIdx.y * blockDim.y threadIdx.y; if (x width y height) { out[y * width x] tex2Dfloat(texObj, x, y); } } int main() { int width 1024, height 1024; size_t bytes width * height * sizeof(float); float *d_data; cudaMalloc(d_data, bytes); // 创建纹理资源描述 cudaResourceDesc resDesc {}; resDesc.resType cudaResourceTypeLinear; resDesc.res.linear.devPtr d_data; resDesc.res.linear.desc.f cudaChannelFormatKindFloat; resDesc.res.linear.desc.x 32; // 32‑bit float resDesc.res.linear.sizeInBytes bytes; // 创建纹理描述 cudaTextureDesc texDesc {}; texDesc.addressMode[0] cudaAddressModeClamp; // 边界处理 texDesc.addressMode[1] cudaAddressModeClamp; texDesc.filterMode cudaFilterModePoint; // 点采样不插值 texDesc.readMode cudaReadModeElementType; // 返回原始类型 cudaTextureObject_t texObj 0; cudaCreateTextureObject(texObj, resDesc, texDesc, NULL); // 启动 Kernel dim3 block(32, 32); dim3 grid((width 31) / 32, (height 31) / 32); textureSampleKernelgrid, block(texObj, d_data, width, height); // 销毁纹理对象 cudaDestroyTextureObject(texObj); cudaFree(d_data); return 0; }3.4 纹理内存的硬件插值纹理对象支持在采样时自动进行双线性插值只需将filterMode设置为cudaFilterModeLinear并使用tex2Dfloat的归一化坐标形式。这可以在图像缩放、变形等场景中直接使用硬件加速。4. 常量内存与纹理内存的选用指南场景推荐原因所有线程读取同一个参数常量内存广播机制极快只读数组访问模式统一常量内存容量 64 KB 以内图像处理、二维邻近访问纹理内存二维缓存优化 硬件插值不规则只读访问如查找表纹理内存纹理缓存对随机只读有优化大数据量只读64 KB纹理内存或全局内存 __ldg常量内存容量不够需要写入全局内存常量和纹理均为只读5. 代码演示常量内存与全局内存的广播读取对比为了更明显地展示常量内存的广播优势我们构造一个场景所有线程在循环中反复读取同一个常量值与反复读取全局内存中的同一个值进行对比。#include cstdio #include cuda_runtime.h __constant__ float c_val; // 从常量内存广播读取 __global__ void constantBroadcast(float *out, int N, int iterations) { int idx blockIdx.x * blockDim.x threadIdx.x; float sum 0.0f; for (int i 0; i iterations; i) { sum c_val; // 广播读取 } if (idx N) out[idx] sum; } // 从全局内存读取同一个地址 __global__ void globalBroadcast(float *in, float *out, int N, int iterations) { int idx blockIdx.x * blockDim.x threadIdx.x; float sum 0.0f; for (int i 0; i iterations; i) { sum in[0]; // 所有线程读同一地址可能被 L1 缓存优化 } if (idx N) out[idx] sum; } int main() { const int N 1 20; const int iterations 1000; float h_val 1.5f; cudaMemcpyToSymbol(c_val, h_val, sizeof(float)); float *d_in, *d_out; cudaMalloc(d_in, sizeof(float)); cudaMalloc(d_out, N * sizeof(float)); cudaMemcpy(d_in, h_val, sizeof(float), cudaMemcpyHostToDevice); cudaEvent_t start, end; cudaEventCreate(start); cudaEventCreate(end); float ms; cudaEventRecord(start); constantBroadcast(N255)/256, 256(d_out, N, iterations); cudaEventRecord(end); cudaEventSynchronize(end); cudaEventElapsedTime(ms, start, end); printf(Constant broadcast: %f ms\n, ms); cudaEventRecord(start); globalBroadcast(N255)/256, 256(d_in, d_out, N, iterations); cudaEventRecord(end); cudaEventSynchronize(end); cudaEventElapsedTime(ms, start, end); printf(Global broadcast: %f ms\n, ms); cudaFree(d_in); cudaFree(d_out); return 0; }**预期结果**常量内存广播版本通常比全局内存广播版本快一些因为常量缓存专用于广播而全局内存广播依赖于 L1 缓存策略可能不如常量缓存高效。6. 课后练习练习1常量内存容量限制编写一个程序尝试声明一个超过 64 KB 的__constant__数组编译运行观察错误信息。然后改用纹理内存或全局内存实现相同功能。练习2常量内存广播 vs 不同地址读取修改常量内存示例使每个线程读取c_arr[threadIdx.x % 8]不同地址与读取c_arr[0]广播对比性能验证广播机制的优越性。练习3纹理对象基本使用编写一个 Kernel使用纹理对象对一张二维图像进行水平翻转通过调整采样坐标展示纹理对象的坐标处理和边界处理能力。练习4纹理内存硬件插值创建一个纹理对象设置filterMode cudaFilterModeLinear使用tex2Dfloat采样图像对比线性插值与点采样的结果差异。可输入一个渐变图像观察效果。练习5卷积中的常量内存优化在卷积算子中将卷积核系数存储在常量内存中所有线程读取相同的核系数。对比使用全局内存传递核系数的性能分析原因。7. 下一步下一节将深入 局部内存与寄存器分配策略你将学习寄存器分配规则与编译器行为寄存器溢出spill的触发条件与性能影响使用-maxrregcount和__launch_bounds__控制寄存器使用局部内存的访问模式与缓存优化局部内存使用的方法避免寄存器溢出
第2板块·第4节：常量内存与纹理内存的优化使用
学习目标学完本节你将能够理解常量内存的物理特性与广播机制掌握常量内存的最佳使用场景及限制了解纹理内存的缓存优化与硬件插值能力使用纹理对象进行基本的图像采样对比常量内存、纹理内存与全局内存在只读场景下的性能差异在卷积、图像处理等场景中合理选择只读缓存优化手段1. 常量内存Constant Memory1.1 物理特性常量内存是位于设备内存上的一块只读区域但拥有独立的片上常量缓存Constant Cache每 SM 通常有 64 KB 的常量缓存。总容量64 KB所有 SM 共享访问延迟同一 Warp 内所有线程读取相同地址时广播机制使访问速度接近寄存器。如果同一 Warp 内线程读取不同地址访问会串行化性能急剧下降。生命周期从 Host 端通过cudaMemcpyToSymbol写入后在整个程序运行期间有效但 Kernel 内只能读不能写。1.2 广播机制当同一 Warp 内的 32 个线程都读取常量内存的同一个地址时硬件会将数据广播给所有线程相当于一次读操作服务整个 Warp这是常量内存最高效的使用方式。适用场景算法参数如缩放系数、阈值滤波器系数固定查找表如果访问模式统一不适用场景每个线程读取不同的常量数组元素例如const_arr[tid % 64]这会导致同一 Warp 内访问多个地址冲突严重。数据量超过 64 KB 的只读数组。1.3 使用方式// 声明常量内存 __constant__ float scaleFactor; __constant__ float filterCoeffs[16]; // Host 端初始化 float scale 2.5f; cudaMemcpyToSymbol(scaleFactor, scale, sizeof(float)); float coeffs[16] {...}; cudaMemcpyToSymbol(filterCoeffs, coeffs, 16 * sizeof(float)); // Kernel 内使用 __global__ void applyFilter(float *in, float *out, int N) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx N) { float val in[idx] * scaleFactor; // 使用 filterCoeffs 进行计算... out[idx] val; } }2. 常量内存 vs 全局内存性能对比下面通过一个简单示例对比常量内存和全局内存在所有线程读取相同参数时的性能差异。#include cstdio #include cuda_runtime.h __constant__ float c_scale; // 常量内存 // 使用常量内存 __global__ void useConstant(float *in, float *out, int N) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx N) { out[idx] in[idx] * c_scale; } } // 使用全局内存参数作为 Kernel 参数传入 __global__ void useGlobalParam(float *in, float *out, float scale, int N) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx N) { out[idx] in[idx] * scale; } } int main() { const int N 1 22; size_t bytes N * sizeof(float); float *d_in, *d_out; cudaMalloc(d_in, bytes); cudaMalloc(d_out, bytes); // 初始化省略... float scale 2.5f; cudaMemcpyToSymbol(c_scale, scale, sizeof(float)); cudaEvent_t start, end; cudaEventCreate(start); cudaEventCreate(end); float ms; // 常量内存版本 cudaEventRecord(start); useConstant(N255)/256, 256(d_in, d_out, N); cudaEventRecord(end); cudaEventSynchronize(end); cudaEventElapsedTime(ms, start, end); printf(Constant memory: %f ms\n, ms); // 全局参数版本 cudaEventRecord(start); useGlobalParam(N255)/256, 256(d_in, d_out, scale, N); cudaEventRecord(end); cudaEventSynchronize(end); cudaEventElapsedTime(ms, start, end); printf(Global parameter: %f ms\n, ms); cudaFree(d_in); cudaFree(d_out); return 0; }**预期结果**两者性能非常接近因为编译器通常会将 Kernel 参数保存在常量缓存中。但关键区别在于容量和可更新性常量内存可以存储较大的只读参数数组如 64 KB 的系数表而 Kernel 参数数量有限。3. 纹理内存Texture Memory3.1 什么是纹理内存纹理内存是 GPU 中专门优化的只读数据访问路径通过纹理缓存Texture Cache提供二维空间局部性优化。底层数据仍存储在全局内存但通过纹理缓存访问。支持硬件自动处理边界clamp, wrap, mirror支持硬件双线性插值、归一化坐标等缓存针对二维局部性优化适合图像处理、卷积等场景3.2 纹理内存与 L1/L2 缓存的区别常规全局内存访问经过 L1/L2 缓存对一维连续访问优化较好。纹理缓存针对二维空间邻近进行优化当线程访问的像素在图像上邻近时缓存命中率很高。纹理缓存是只读的因此不能用于写入。3.3 纹理对象Texture Object的基本用法现代 CUDA 推荐使用纹理对象它在 Kernel 中以cudaTextureObject_t类型传递。#include cuda_runtime.h // Kernel 使用纹理对象 __global__ void textureSampleKernel(cudaTextureObject_t texObj, float *out, int width, int height) { int x blockIdx.x * blockDim.x threadIdx.x; int y blockIdx.y * blockDim.y threadIdx.y; if (x width y height) { out[y * width x] tex2Dfloat(texObj, x, y); } } int main() { int width 1024, height 1024; size_t bytes width * height * sizeof(float); float *d_data; cudaMalloc(d_data, bytes); // 创建纹理资源描述 cudaResourceDesc resDesc {}; resDesc.resType cudaResourceTypeLinear; resDesc.res.linear.devPtr d_data; resDesc.res.linear.desc.f cudaChannelFormatKindFloat; resDesc.res.linear.desc.x 32; // 32‑bit float resDesc.res.linear.sizeInBytes bytes; // 创建纹理描述 cudaTextureDesc texDesc {}; texDesc.addressMode[0] cudaAddressModeClamp; // 边界处理 texDesc.addressMode[1] cudaAddressModeClamp; texDesc.filterMode cudaFilterModePoint; // 点采样不插值 texDesc.readMode cudaReadModeElementType; // 返回原始类型 cudaTextureObject_t texObj 0; cudaCreateTextureObject(texObj, resDesc, texDesc, NULL); // 启动 Kernel dim3 block(32, 32); dim3 grid((width 31) / 32, (height 31) / 32); textureSampleKernelgrid, block(texObj, d_data, width, height); // 销毁纹理对象 cudaDestroyTextureObject(texObj); cudaFree(d_data); return 0; }3.4 纹理内存的硬件插值纹理对象支持在采样时自动进行双线性插值只需将filterMode设置为cudaFilterModeLinear并使用tex2Dfloat的归一化坐标形式。这可以在图像缩放、变形等场景中直接使用硬件加速。4. 常量内存与纹理内存的选用指南场景推荐原因所有线程读取同一个参数常量内存广播机制极快只读数组访问模式统一常量内存容量 64 KB 以内图像处理、二维邻近访问纹理内存二维缓存优化 硬件插值不规则只读访问如查找表纹理内存纹理缓存对随机只读有优化大数据量只读64 KB纹理内存或全局内存 __ldg常量内存容量不够需要写入全局内存常量和纹理均为只读5. 代码演示常量内存与全局内存的广播读取对比为了更明显地展示常量内存的广播优势我们构造一个场景所有线程在循环中反复读取同一个常量值与反复读取全局内存中的同一个值进行对比。#include cstdio #include cuda_runtime.h __constant__ float c_val; // 从常量内存广播读取 __global__ void constantBroadcast(float *out, int N, int iterations) { int idx blockIdx.x * blockDim.x threadIdx.x; float sum 0.0f; for (int i 0; i iterations; i) { sum c_val; // 广播读取 } if (idx N) out[idx] sum; } // 从全局内存读取同一个地址 __global__ void globalBroadcast(float *in, float *out, int N, int iterations) { int idx blockIdx.x * blockDim.x threadIdx.x; float sum 0.0f; for (int i 0; i iterations; i) { sum in[0]; // 所有线程读同一地址可能被 L1 缓存优化 } if (idx N) out[idx] sum; } int main() { const int N 1 20; const int iterations 1000; float h_val 1.5f; cudaMemcpyToSymbol(c_val, h_val, sizeof(float)); float *d_in, *d_out; cudaMalloc(d_in, sizeof(float)); cudaMalloc(d_out, N * sizeof(float)); cudaMemcpy(d_in, h_val, sizeof(float), cudaMemcpyHostToDevice); cudaEvent_t start, end; cudaEventCreate(start); cudaEventCreate(end); float ms; cudaEventRecord(start); constantBroadcast(N255)/256, 256(d_out, N, iterations); cudaEventRecord(end); cudaEventSynchronize(end); cudaEventElapsedTime(ms, start, end); printf(Constant broadcast: %f ms\n, ms); cudaEventRecord(start); globalBroadcast(N255)/256, 256(d_in, d_out, N, iterations); cudaEventRecord(end); cudaEventSynchronize(end); cudaEventElapsedTime(ms, start, end); printf(Global broadcast: %f ms\n, ms); cudaFree(d_in); cudaFree(d_out); return 0; }**预期结果**常量内存广播版本通常比全局内存广播版本快一些因为常量缓存专用于广播而全局内存广播依赖于 L1 缓存策略可能不如常量缓存高效。6. 课后练习练习1常量内存容量限制编写一个程序尝试声明一个超过 64 KB 的__constant__数组编译运行观察错误信息。然后改用纹理内存或全局内存实现相同功能。练习2常量内存广播 vs 不同地址读取修改常量内存示例使每个线程读取c_arr[threadIdx.x % 8]不同地址与读取c_arr[0]广播对比性能验证广播机制的优越性。练习3纹理对象基本使用编写一个 Kernel使用纹理对象对一张二维图像进行水平翻转通过调整采样坐标展示纹理对象的坐标处理和边界处理能力。练习4纹理内存硬件插值创建一个纹理对象设置filterMode cudaFilterModeLinear使用tex2Dfloat采样图像对比线性插值与点采样的结果差异。可输入一个渐变图像观察效果。练习5卷积中的常量内存优化在卷积算子中将卷积核系数存储在常量内存中所有线程读取相同的核系数。对比使用全局内存传递核系数的性能分析原因。7. 下一步下一节将深入 局部内存与寄存器分配策略你将学习寄存器分配规则与编译器行为寄存器溢出spill的触发条件与性能影响使用-maxrregcount和__launch_bounds__控制寄存器使用局部内存的访问模式与缓存优化局部内存使用的方法避免寄存器溢出