如何向CUDA发送提示，异步将指定变量预取至L1缓存？

问题背景 #

当使用非LRU友好的确定性数据访问模式（如自定义的crazy_access_order函数）时，CUDA默认的LRU缓存淘汰策略无法有效利用缓存，导致数据加载延迟过高。尝试通过循环展开等编译器优化手段引导自动预取，但效果不佳，推测编译器未生成预期的预取指令。

解决方案：手动PTX预取 #

CUDA提供PTX层面的prefetch指令（编译后对应SASS架构的CCTL或CCTLL指令），可异步触发指定数据加载至L1缓存，且指令会立即返回，无需等待预取完成，完全匹配需求中的async_prefetch_to_l1功能。

实现自定义预取函数 #

通过内联PTX代码实现异步预取到L1的函数：

__device__ inline void async_prefetch_to_l1(const void* addr) {
    // PTX prefetch.global.L1 指令：将指定地址的数据预取到L1缓存
    // 参数说明：[%0]为预取地址，"l"表示64位地址寄存器
    asm volatile("prefetch.global.L1 [%0];" : : "l"(addr));
}

适配访问模式的预取代码示例 #

结合你的crazy_access_order访问模式，按预取距离提前加载数据：

#define PREFETCH_DISTANCE 4

// 自定义非连续访问顺序函数
__device__ int crazy_access_order(const int i) {
    // 示例：非连续的索引计算逻辑
    return (i * 7) % N;
}

const __device__ MyType data[N] = {...};

__global__ void kernel() {
    // 预热：提前预取前PREFETCH_DISTANCE-1个数据
    for (int i = 0; i < PREFETCH_DISTANCE - 1; ++i) {
        int idx = crazy_access_order(i);
        async_prefetch_to_l1(&data[idx]);
    }

    // 主计算循环：提前预取后续数据
    for (int i = 0; i < N; ++i) {
        // 预取PREFETCH_DISTANCE步后的目标数据（避免越界）
        if (i + PREFETCH_DISTANCE < N) {
            int prefetch_idx = crazy_access_order(i + PREFETCH_DISTANCE);
            async_prefetch_to_l1(&data[prefetch_idx]);
        }

        // 访问已预取到L1的数据
        int curr_idx = crazy_access_order(i);
        const MyType datum = data[curr_idx];

        // 执行数据计算操作
        // ...
    }
}

效果验证 #

手动使用PTXprefetch指令后，性能直接提升2倍，且Long Scoreboard停顿显著减少，证明预取策略有效缓解了缓存未命中带来的延迟。

内容的提问来源于stack exchange，提问作者emchristiansen