我们可以在 gpu 的 l1 缓存上有脏数据吗？

Question

我已经阅读了 GPU 微体系结构中的一些常见写入策略。大部分GPU写的策略和下图一样（图片来自the gpgpu-sim manual）。根据下图，我有一个问题。我们可以在 l1 缓存上有脏数据吗？

Answer 1

某些 GPU 架构上的 L1 is a write-back cache 用于全局访问。请注意，本主题因 GPU 架构而异，例如全局activity是否缓存在L1.

一般来说，是的，您可能会有脏数据。我的意思是 L1 缓存中的数据已修改（与全局 space 或 L2 缓存中的数据相比）并且尚未“刷新”或更新到 L2 缓存中。（您也可以拥有“陈旧”数据 - L1 中的数据未被修改，但与 L2 不一致。）

我们可以为此（脏数据）创建一个简单的证明点。

以下代码在 cc7.0 设备（可能还有其他一些架构）上执行时不会给出预期的答案 1024。

这是因为 L1（每个 SM 是一个单独的实体）不会立即刷新到 L2。因此，根据上述定义，它具有“脏数据”。

（代码因此被破坏，请勿使用此代码，它只是一个证明点。）

#include <iostream>
#include <cuda_runtime.h>

constexpr int num_blocks = 1024;
constexpr int num_threads = 32;

struct Lock {
  int *locked;

  Lock() {
    int init = 0;
    cudaMalloc(&locked, sizeof(int));
    cudaMemcpy(locked, &init, sizeof(int), cudaMemcpyHostToDevice);
  }

  ~Lock() {
    if (locked) cudaFree(locked);
    locked = NULL;
  }

  __device__ __forceinline__ void acquire_lock() {
    while (atomicCAS(locked, 0, 1) != 0);
  }

  __device__ __forceinline__ void unlock() {
    atomicExch(locked, 0);
  }
};

__global__ void counter(Lock lock, int *total) {
  if (threadIdx.x == 1) {
    lock.acquire_lock();
    *total = *total + 1;
//    __threadfence();  uncomment this line to fix
    lock.unlock();
  }
}

int main() {
  int *total_dev;
  cudaMalloc(&total_dev, sizeof(int));
  int total_host = 0;
  cudaMemcpy(total_dev, &total_host, sizeof(int), cudaMemcpyHostToDevice);
  {
  Lock lock;
  counter<<<num_blocks, num_threads>>>(lock, total_dev);
  cudaDeviceSynchronize();
  cudaMemcpy(&total_host, total_dev, sizeof(int), cudaMemcpyDeviceToHost);
  std::cout << total_host << std::endl;
  }
  cudaFree(total_dev);
}

如果对这是否是一个适当的证明有任何进一步的疑问（例如，为了消除关于事物被“优化到寄存器中”等的争论），我们可以研究结果 sass 代码。上面内核的末尾有这样的代码：

    /*0130*/                   LDG.E.SYS R0, [R4] ;                           /* 0x0000000004007381 */
                               // load *total                                               /* 0x000ea400001ee900 */
    /*0140*/                   IADD3 R7, R0, 0x1, RZ ;                        /* 0x0000000100077810 */
                               // add 1                                               /* 0x004fd00007ffe0ff */
    /*0150*/                   STG.E.SYS [R4], R7 ;                           /* 0x0000000704007386 */
                               // store *total                                               /* 0x000fe8000010e900 */
    /*0160*/                   ATOMG.E.EXCH.STRONG.GPU PT, RZ, [R2], RZ ;     /* 0x000000ff02ff73a8 */
                               //lock.unlock                                               /* 0x000fe200041f41ff */
    /*0170*/                   EXIT ;

由于结果寄存器肯定已经存储到全局space，我们可以推断如果另一个线程（在另一个SM中）在全局space中读取了一个意外的值for *total 这一定是因为来自另一个 SM 的存储还没有到达 L2，即还没有到达设备范围 consistency/coherency。因此，其他一些 SM 中的数据是“脏”的。我们可以（大概）在这里排除“陈旧”的情况（另一个 L1 中的数据是 写入的，但我的 L1 中有“旧”数据）因为上面指出的全局负载在 SM 中获取锁之前不会发生。

请注意，上面的代码在 cc7.0 设备（可能还有一些其他设备架构）上“失败”。它不一定会在您使用的 GPU 上失败。但还是“坏”了。

我们可以在 gpu 的 l1 缓存上有脏数据吗？

Can we have dirty data on l1 cache in gpu?

gpu

gpgpu

nvidia