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

Posted 2023-03-16

【中文标题】我们可以在 gpu 的 l1 缓存上存储脏数据吗？

【英文标题】：Can we have dirty data on l1 cache in gpu?

【发布时间】：2021-09-21 16:07:07

【问题描述】：

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

【参考方案1】：

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

一般来说，是的，您可能有脏数据。我的意思是 L1 缓存中的数据被修改（与全局空间或 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 ; 

既然结果寄存器肯定已经存入了全局空间，我们可以推断，如果另一个线程（在另一个 SM 中）在全局空间中为*total 读取了一个意想不到的值，那一定是因为存储来自另一个 SM 尚未达到 L2，即尚未达到设备范围的一致性/一致性。因此，其他一些 SM 中的数据是“脏”的。我们可以（大概）排除这里的“陈旧”情况（其他 L1 中的数据已写入，但我的 L1 中有“旧”数据），因为上面指出的全局负载不会发生直到在 SM 中获得锁。

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