CUDA进阶学习与深入

什么需要错误处理？
CUDA API 调用可能失败，常见原因：

内存不足
设备不存在
内核启动失败
驱动程序错误
不检查错误会导致：

程序崩溃
结果错误
难以调试
CUDA 错误类型
typedef enum cudaError {
cudaSuccess = 0, // 成功
cudaErrorInvalidValue = 1, // 无效参数
cudaErrorMemoryAllocation = 2, // 内存分配失败
cudaErrorInvalidDevice = 10, // 无效设备
cudaErrorInvalidMemcpyDirection = 21, // 无效拷贝方向
// … 更多错误码
} cudaError;
错误检查函数
// 基本错误检查
cudaError_t err = cudaMalloc(&d_data, size);
if (err != cudaSuccess) {
printf(“CUDA 错误: %s\n”, cudaGetErrorString(err));
exit(1);
}
封装错误检查宏
// 定义错误检查宏
#define CUDA_CHECK(call)
do {
cudaError_t err = call;
if (err != cudaSuccess) {
fprintf(stderr, “CUDA 错误 at %s:%d: %s\n”,
FILE,LINE, cudaGetErrorString(err));
exit(1);
}
} while(0)

// 使用宏
CUDA_CHECK(cudaMalloc(&d_data, size));
CUDA_CHECK(cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice));
内核启动错误检查
globalvoid myKernel(int *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
data[idx] = idx * 2;
}
}

int main() {
// 启动内核
myKernel<<<grid, block>>>(d_data, n);

// 检查内核启动错误 cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { printf("内核启动失败: %s\n", cudaGetErrorString(err)); return -1; } // 等待内核完成并检查执行错误 err = cudaDeviceSynchronize(); if (err != cudaSuccess) { printf("内核执行失败: %s\n", cudaGetErrorString(err)); return -1; } return 0;

}
完整的错误处理模板
#include <stdio.h>
#include <stdlib.h>

#define CUDA_CHECK(call)
do {
cudaError_t err = call;
if (err != cudaSuccess) {
fprintf(stderr, “CUDA 错误 at %s:%d: %s\n”,
FILE,LINE, cudaGetErrorString(err));
exit(1);
}
} while(0)

#define CUDA_KERNEL_CHECK()
do {
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
fprintf(stderr, “内核启动错误 at %s:%d: %s\n”,
FILE,LINE, cudaGetErrorString(err));
exit(1);
}
err = cudaDeviceSynchronize();
if (err != cudaSuccess) {
fprintf(stderr, “内核执行错误 at %s:%d: %s\n”,
FILE,LINE, cudaGetErrorString(err));
exit(1);
}
} while(0)

int main() {
int n = 1000;
size_t size = n * sizeof(float);

float *d_data; CUDA_CHECK(cudaMalloc(&d_data, size)); myKernel<<<grid, block>>>(d_data, n); CUDA_KERNEL_CHECK(); CUDA_CHECK(cudaFree(d_data)); return 0;

}
练习题 9
CUDA 错误码 cudaSuccess 的值是什么？
cudaGetLastError() 和 cudaDeviceSynchronize() 分别检查什么错误？
为什么内核启动后需要调用 cudaDeviceSynchronize() 才能检测到执行错误？
第十课：原子操作知识点
什么是原子操作？
原子操作是不可分割的操作，在多线程环境下保证数据一致性。

问题场景：

// 非原子操作（危险！）
int count = 0;
globalvoid increment(int *count) {
(*count)++; // 多个线程同时执行，结果不确定
}
解决方案：使用原子操作

CUDA 原子函数
函数 操作 说明
atomicAdd() 加法 *addr += val
atomicSub() 减法 *addr -= val
atomicExch() 交换 *addr = val
atomicMin() 最小值 *addr = min(*addr, val)
atomicMax() 最大值 *addr = max(*addr, val)
atomicInc() 递增 *addr = (*addr >= val) ? 0 : *addr + 1
atomicDec() 递减 `addr = (addr == 0)
atomicCAS() 比较并交换 条件交换
atomicAnd() 与运算 *addr &= val
atomicOr() 或运算 *addr |= val
atomicXor() 异或运算 *addr ^= val
atomicAdd 示例
#include <stdio.h>

globalvoid atomicAddKernel(int *count, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
atomicAdd(count, 1); // 原子递增
}
}

int main() {
int n = 10000;
int h_count = 0;
int *d_count;

cudaMalloc(&d_count, sizeof(int)); cudaMemcpy(d_count, &h_count, sizeof(int), cudaMemcpyHostToDevice); int blockSize = 256; int gridSize = (n + blockSize - 1) / blockSize; atomicAddKernel<<<gridSize, blockSize>>>(d_count, n); cudaMemcpy(&h_count, d_count, sizeof(int), cudaMemcpyDeviceToHost); printf("计数结果: %d (预期: %d)\n", h_count, n); cudaFree(d_count); return 0;

}
atomicCAS（比较并交换）
// atomicCAS(int *addr, int compare, int val)
// 如果 *addr == compare，则 *addr = val
// 返回 *addr 的旧值

globalvoid casExample(int *data, int old_val, int new_val) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx == 0) {
int old = atomicCAS(data, old_val, new_val);
printf(“旧值: %d, 新值: %d\n”, old, new_val);
}
}
原子操作实现锁
struct Lock {
int *mutex;

Lock() { cudaMalloc(&mutex, sizeof(int)); cudaMemset(mutex, 0, sizeof(int)); } ~Lock() { cudaFree(mutex); } __device__ void lock() { while (atomicCAS(mutex, 0, 1) != 0) { // 等待锁释放 } } __device__ void unlock() { atomicExch(mutex, 0); }

};

globalvoid kernelWithLock(int *data, Lock lock) {
lock.lock();
// 临界区代码
(*data)++;
lock.unlock();
}
这段代码是 CUDA（GPU 编程）中非常经典的一种锁机制实现，叫做“自旋锁”（Spinlock）。

要理解这段代码，需要弄懂两个核心概念：atomicCAS 是什么，以及 while 循环在干什么。

1. 核心概念：atomicCAS
   atomicCAS 全称是 Atomic Compare-And-Swap（原子比较并交换）。
   在这个函数中：atomicCAS(mutex, 0, 1) 接收三个参数：

参数 1 (mutex)：你要操作的那个变量（锁的状态）。
参数 2 (0)：你期望此时锁的值是多少（0 表示锁当前是空闲的）。
参数 3 (1)：如果锁真的像你期望的一样是空闲的（为 0），你就把它改成新值（1 表示你占用了这个锁）。
⚠️ 最容易产生误解的地方（必须记住）：
atomicCAS 的返回值永远是 mutex 改变之前的“旧值”。它并不是返回一个 True 或 False！

“原子操作”意味着这个动作是瞬间完成的，绝对不可被打断。就算有 1000 个 GPU 线程同时执行这行代码，硬件也会保证它们一个一个排队执行这个判断和交换的过程。

2. 场景推演：它是怎么锁住的？
   我们假设有线程 A 和 线程 B 同时想要获取这个锁。初始状态下，锁是解开的，也就是 mutex = 0。

场景一：线程 A 先到达
线程 A 执行 atomicCAS(mutex, 0, 1)。
硬件一看，当前的 mutex 确实是 0（没人占用）。
于是硬件把 mutex 改成了 1（表示被线程 A 锁上了）。
返回值： 返回 mutex 被修改前的旧值，也就是 0。
来看 while 判断条件：while( 0 != 0 )。
这个条件是 假 (False)！所以线程 A 跳出 while 循环，成功拿到锁，去执行后面的代码了。
场景二：线程 B 紧接着到达（此时线程 A 还没释放锁）
此时 mutex 已经被线程 A 变成了 1。
线程 B 执行 atomicCAS(mutex, 0, 1)。
硬件一看，当前的 mutex 是 1，跟你期望的 0 不相等！
所以硬件什么都不做（不会把值改成 1）。
返回值： 依然返回 mutex 此时的旧值，也就是 1。
来看 while 判断条件：while( 1 != 0 )。
这个条件是 真 (True)！所以线程 B 被困在了 while 循环里，只能再次执行 atomicCAS 进行判断。
只要线程 A 不放开锁，线程 B 就会一直在 while 里面疯狂打转（这就是为什么叫“自旋锁”，它在原地自旋等待）。
3. 如何解锁？
虽然你没贴出解锁的代码，但配合起来看更容易理解。解锁的代码通常非常简单：

devicevoid unlock() {
atomicExch(mutex, 0); // 释放锁：直接把 mutex 变回 0
}
一旦线程 A 执行了 unlock()，把 mutex 改回了 0。
还在 while 里苦苦循环的线程 B 在下一次执行 atomicCAS(mutex, 0, 1) 时，就会发现旧值变成 0 了，于是它成功把 mutex 变成 1，返回值变为 0，跳出循环，成功接管这把锁。

💡 通俗的比喻
把 mutex 想象成公共厕所门上的那个指示牌（0 代表绿色的“无人”，1 代表红色的“有人”）。

atomicCAS 就是你跑过去看一眼并锁门的连贯动作。
你走过去，发现是绿色（0），你立刻进去并把牌子翻到红色（1）。因为你进来前是绿色（返回 0），所以你不用排队（跳出循环）。
另一个人走过来，发现是红色（1），他想翻牌子失败了。因为他看到的是红色（返回 1），所以他只能在门外一直转圈踱步（while 循环），死死盯着牌子，直到你出来把牌子翻回绿色（0）为止。
原子操作性能考虑
原子操作比普通操作慢
多个线程对同一地址原子操作会串行化
尽量减少原子操作的使用
考虑使用共享内存减少全局内存原子操作
练习题 10
为什么多线程环境下普通递增操作 (*count)++ 会产生错误结果？
atomicAdd(addr, val) 的作用是什么？返回值是什么？
如何使用原子操作实现一个简单的互斥锁？
第十一课：CUDA 流与异步执行知识点
什么是 CUDA 流？
CUDA 流是一系列按顺序执行的命令队列。不同流中的命令可以并发执行。

默认流（Stream 0）：
┌─────────────────────────────────────┐
│ 内核A → 拷贝1 → 内核B → 拷贝2 │ 串行执行
└─────────────────────────────────────┘

多流并发：
Stream 1: ┌─────────────────────────────┐
│ 内核A → 拷贝1 │
└─────────────────────────────┘
Stream 2: ┌─────────────────────────────┐
│ 内核B → 拷贝2 │ 并发执行
└─────────────────────────────┘
创建和使用流
cudaStream_t stream1, stream2;

// 创建流
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 在指定流中执行操作
cudaMemcpyAsync(d_a1, h_a1, size, cudaMemcpyHostToDevice, stream1);
kernel1<<<grid, block, 0, stream1>>>(d_a1, d_c1);
cudaMemcpyAsync(h_c1, d_c1, size, cudaMemcpyDeviceToHost, stream1);

cudaMemcpyAsync(d_a2, h_a2, size, cudaMemcpyHostToDevice, stream2);
kernel2<<<grid, block, 0, stream2>>>(d_a2, d_c2);
cudaMemcpyAsync(h_c2, d_c2, size, cudaMemcpyDeviceToHost, stream2);

// 同步流
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

// 销毁流
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);
异步内存拷贝
// 同步拷贝（阻塞）
cudaMemcpy(dst, src, size, cudaMemcpyHostToDevice);

// 异步拷贝（非阻塞）
cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream);
流同步
// 同步单个流
cudaStreamSynchronize(stream);

// 同步所有流
cudaDeviceSynchronize();

// 等待多个流
cudaStreamWaitEvent(stream, event);
流优先级
// 创建高优先级流
int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);

cudaStream_t stream_high;
cudaStreamCreateWithPriority(&stream_high, cudaStreamNonBlocking, priority_high);
完整示例：多流并发
#include <stdio.h>

#define N_STREAMS 4
#define N 1000000

globalvoid vectorAdd(float *a, float *b, float *c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}

int main() {
int n = N;
size_t size = n * sizeof(float);

// 分配主机内存（页锁定内存，用于异步传输） float *h_a[N_STREAMS], *h_b[N_STREAMS], *h_c[N_STREAMS]; for (int i = 0; i < N_STREAMS; i++) { cudaMallocHost(&h_a[i], size); cudaMallocHost(&h_b[i], size); cudaMallocHost(&h_c[i], size); } // 分配设备内存 float *d_a[N_STREAMS], *d_b[N_STREAMS], *d_c[N_STREAMS]; for (int i = 0; i < N_STREAMS; i++) { cudaMalloc(&d_a[i], size); cudaMalloc(&d_b[i], size); cudaMalloc(&d_c[i], size); } // 创建流 cudaStream_t streams[N_STREAMS]; for (int i = 0; i < N_STREAMS; i++) { cudaStreamCreate(&streams[i]); } // 并发执行 int blockSize = 256; int gridSize = (n + blockSize - 1) / blockSize; for (int i = 0; i < N_STREAMS; i++) { cudaMemcpyAsync(d_a[i], h_a[i], size, cudaMemcpyHostToDevice, streams[i]); cudaMemcpyAsync(d_b[i], h_b[i], size, cudaMemcpyHostToDevice, streams[i]); vectorAdd<<<gridSize, blockSize, 0, streams[i]>>>(d_a[i], d_b[i], d_c[i], n); cudaMemcpyAsync(h_c[i], d_c[i], size, cudaMemcpyDeviceToHost, streams[i]); } // 同步所有流 cudaDeviceSynchronize(); // 清理 for (int i = 0; i < N_STREAMS; i++) { cudaStreamDestroy(streams[i]); cudaFree(d_a[i]); cudaFree(d_b[i]); cudaFree(d_c[i]); cudaFreeHost(h_a[i]); cudaFreeHost(h_b[i]); cudaFreeHost(h_c[i]); } return 0;

}
页锁定内存（Pinned Memory）
// 普通主机内存（可分页）
floath_data = (float)malloc(size);

// 页锁定主机内存（不可分页，用于异步传输）
float *h_data_pinned;
cudaMallocHost(&h_data_pinned, size);

// 释放
cudaFreeHost(h_data_pinned);
优点：

支持异步传输
传输速度更快
DMA 直接访问
缺点：

占用物理内存
分配速度较慢
练习题 11
CUDA 流的作用是什么？
cudaMemcpy 和 cudaMemcpyAsync 的区别是什么？
为什么异步传输需要使用页锁定内存？
第十二课：CUDA 事件与性能计时知识点
什么是 CUDA 事件？
CUDA 事件是 GPU 上的时间标记，用于：

测量内核执行时间
流同步
性能分析
创建和使用事件
cudaEvent_t start, stop;

// 创建事件
cudaEventCreate(&start);
cudaEventCreate(&stop);

// 记录事件
cudaEventRecord(start);
myKernel<<<grid, block>>>(…);
cudaEventRecord(stop);

// 等待事件完成
cudaEventSynchronize(stop);

// 计算时间
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf(“执行时间: %.3f ms\n”, milliseconds);

// 销毁事件
cudaEventDestroy(start);
cudaEventDestroy(stop);
完整的性能测试示例
#include <stdio.h>

globalvoid vectorAdd(float *a, float *b, float *c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}

int main() {
int n = 10000000;
size_t size = n * sizeof(float);

// 分配内存 float *h_a = (float*)malloc(size); float *h_b = (float*)malloc(size); float *h_c = (float*)malloc(size); float *d_a, *d_b, *d_c; cudaMalloc(&d_a, size); cudaMalloc(&d_b, size); cudaMalloc(&d_c, size); // 初始化数据 for (int i = 0; i < n; i++) { h_a[i] = (float)i; h_b[i] = (float)(i * 2); } // 拷贝数据 cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice); // 创建事件 cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); // 测试不同 Block 大小 int blockSizes[] = {32, 64, 128, 256, 512, 1024}; int numTests = sizeof(blockSizes) / sizeof(int); for (int i = 0; i < numTests; i++) { int blockSize = blockSizes[i]; int gridSize = (n + blockSize - 1) / blockSize; // 预热 vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n); cudaDeviceSynchronize(); // 计时 cudaEventRecord(start); vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n); cudaEventRecord(stop); cudaEventSynchronize(stop); float ms; cudaEventElapsedTime(&ms, start, stop); printf("Block=%4d, Grid=%6d, 时间=%.3f ms\n", blockSize, gridSize, ms); } // 清理 cudaEventDestroy(start); cudaEventDestroy(stop); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); free(h_a); free(h_b); free(h_c); return 0;

}
事件同步
// 等待单个事件
cudaEventSynchronize(event);

// 流等待事件
cudaStreamWaitEvent(stream, event, 0);

// 事件完成检查
cudaError_t err = cudaEventQuery(event);
if (err == cudaSuccess) {
printf(“事件已完成\n”);
}
流间同步
cudaStream_t stream1, stream2;
cudaEvent_t event;

cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
cudaEventCreate(&event);

// Stream1 记录事件
kernel1<<<grid, block, 0, stream1>>>(…);
cudaEventRecord(event, stream1);

// Stream2 等待事件
cudaStreamWaitEvent(stream2, event, 0);
kernel2<<<grid, block, 0, stream2>>>(…); // 等待 kernel1 完成
练习题 12
CUDA 事件的主要用途是什么？
cudaEventRecord() 和 cudaEventSynchronize() 的区别？
如何使用事件实现两个流之间的同步？
第十三课：统一内存知识点
什么是统一内存？
统一内存（Unified Memory）创建一个在 CPU 和 GPU 之间共享的内存池，自动管理数据传输。

传统内存模型：
┌─────────┐ ┌─────────┐
│ CPU 内存 │ ←─────→ │ GPU 内存 │
└─────────┘ 手动传输 └─────────┘

统一内存模型：
┌─────────────────────────────┐
│ 统一内存池 │
│ CPU 和 GPU 共享访问 │
│ 自动管理数据迁移 │
└─────────────────────────────┘
创建统一内存
// 分配统一内存
float *data;
cudaMallocManaged(&data, size);

// CPU 和 GPU 都可以直接访问
data[0] = 1.0f; // CPU 写入

myKernel<<<grid, block>>>(data); // GPU 读取和修改

cudaDeviceSynchronize();

printf(“%f\n”, data[0]); // CPU 读取 GPU 修改后的值

// 释放
cudaFree(data);
完整示例
#include <stdio.h>

globalvoid add(float *a, float *b, float *c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}

int main() {
int n = 1000000;
size_t size = n * sizeof(float);

// 分配统一内存 float *a, *b, *c; cudaMallocManaged(&a, size); cudaMallocManaged(&b, size); cudaMallocManaged(&c, size); // CPU 初始化 for (int i = 0; i < n; i++) { a[i] = (float)i; b[i] = (float)(i * 2); } // GPU 计算 int blockSize = 256; int gridSize = (n + blockSize - 1) / blockSize; add<<<gridSize, blockSize>>>(a, b, c, n); // 等待 GPU 完成 cudaDeviceSynchronize(); // CPU 验证结果 bool success = true; for (int i = 0; i < n; i++) { if (c[i] != a[i] + b[i]) { success = false; break; } } printf("验证: %s\n", success ? "成功" : "失败"); // 释放 cudaFree(a); cudaFree(b); cudaFree(c); return 0;

}
内存迁移提示
// 提示内存将在设备上访问
cudaMemPrefetchAsync(data, size, deviceId, stream);

// 提示内存访问模式
cudaMemAdvise(data, size, cudaMemAdviseSetPreferredLocation, deviceId);
// 预取示例
int deviceId;
cudaGetDevice(&deviceId);

// 初始化数据（CPU）
for (int i = 0; i < n; i++) {
data[i] = i;
}

// 预取到 GPU
cudaMemPrefetchAsync(data, size, deviceId);

// GPU 计算
kernel<<<grid, block>>>(data, n);
统一内存的优势
优势 说明
简化编程 无需手动管理内存拷贝
减少代码 不需要 cudaMalloc/cudaMemcpy
自动迁移 数据按需在 CPU/GPU 间迁移
超额订阅 可使用超过 GPU 内存的数据
统一内存的限制
性能可能不如手动优化
需要调用 cudaDeviceSynchronize() 同步
频繁迁移会有开销
旧 GPU 不支持某些特性
练习题 13
统一内存与传统内存模型的主要区别是什么？
cudaMallocManaged() 分配的内存可以被谁访问？
为什么使用统一内存后还需要调用 cudaDeviceSynchronize()？
第十四课：常量内存知识点
什么是常量内存？
常量内存是只读内存，具有缓存优化，适合存储不会改变的数据。

常量内存特点：

• 大小限制：64KB
• 只读
• 有缓存
• 适合广播读取（所有线程读取相同地址）
  声明和使用常量内存
  // 声明常量内存（全局作用域）
  constantfloat constData[256];

// 从主机拷贝到常量内存
cudaMemcpyToSymbol(constData, h_data, size);

// 在内核中使用
globalvoid kernel(float *output) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
output[idx] = constData[idx % 256]; // 读取常量内存
}
完整示例：使用常量内存存储滤波器
#include <stdio.h>

#define FILTER_SIZE 5

// 声明常量内存
constantfloat filter[FILTER_SIZE];

globalvoid convolution(float *input, float *output, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;

if (idx >= FILTER_SIZE / 2 && idx < n - FILTER_SIZE / 2) { float sum = 0.0f; for (int i = 0; i < FILTER_SIZE; i++) { sum += input[idx - FILTER_SIZE / 2 + i] * filter[i]; } output[idx] = sum; }

}

int main() {
int n = 1000;
size_t size = n * sizeof(float);

// 准备滤波器数据 float h_filter[FILTER_SIZE] = {0.1f, 0.2f, 0.4f, 0.2f, 0.1f}; // 拷贝到常量内存 cudaMemcpyToSymbol(filter, h_filter, FILTER_SIZE * sizeof(float)); // 分配内存 float *h_input = (float*)malloc(size); float *h_output = (float*)malloc(size); float *d_input, *d_output; cudaMalloc(&d_input, size); cudaMalloc(&d_output, size); // 初始化输入 for (int i = 0; i < n; i++) { h_input[i] = (float)i; } // 拷贝数据 cudaMemcpy(d_input, h_input, size, cudaMemcpyHostToDevice); // 执行卷积 int blockSize = 256; int gridSize = (n + blockSize - 1) / blockSize; convolution<<<gridSize, blockSize>>>(d_input, d_output, n); // 拷贝结果 cudaMemcpy(h_output, d_output, size, cudaMemcpyDeviceToHost); // 清理 cudaFree(d_input); cudaFree(d_output); free(h_input); free(h_output); return 0;

}
常量内存 vs 全局内存
特性 常量内存 全局内存
访问权限 只读 读写
大小限制 64KB 大
缓存 有（常量缓存） 有（L1/L2）
广播优化 是 否
适用场景 常量数据、滤波器、查找表 通用数据
何时使用常量内存？
数据不会改变
所有线程读取相同数据（广播）
数据量小于 64KB
需要缓存优化
练习题 14
常量内存的大小限制是多少？
使用什么函数将数据拷贝到常量内存？
常量内存适合什么场景？
第十五课：纹理内存知识点
什么是纹理内存？
纹理内存是专门为图像处理优化的只读内存，支持：

硬件插值
边界处理
缓存优化
纹理内存特点
纹理内存特性：

• 只读
• 支持插值（线性、最近邻）
• 支持边界模式（截断、环绕、镜像）
• 2D 空间局部性缓存优化
• 适合图像处理
  使用纹理内存
  // 声明纹理引用（旧方式，CUDA 10 之前）
  texture<float, 2, cudaReadModeElementType> texRef;

// 绑定纹理
cudaBindTextureToArray(texRef, texArray);

// 在内核中读取
globalvoid kernel(float *output, 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) { // 使用纹理读取（支持插值） float val = tex2D(texRef, x, y); output[y * width + x] = val; }

}

// 解绑纹理
cudaUnbindTexture(texRef);
现代方式：纹理对象（CUDA 11+）
// 创建纹理对象
cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeLinear;
resDesc.res.linear.devPtr = devPtr;
resDesc.res.linear.desc = cudaCreateChannelDesc();
resDesc.res.linear.sizeInBytes = size;

cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = cudaAddressModeClamp;
texDesc.addressMode[1] = cudaAddressModeClamp;
texDesc.filterMode = cudaFilterModeLinear;
texDesc.readMode = cudaReadModeElementType;

cudaTextureObject_t texObj;
cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

// 在内核中使用
globalvoid kernel(cudaTextureObject_t texObj, float *output, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
output[idx] = tex1Dfetch(texObj, idx);
}
}

// 销毁纹理对象
cudaDestroyTextureObject(texObj);