理解nvprof事件"l2_subp0_write_sector_misses“和"l2_subp1_write_sector_misses”的结果

Question

通过查看l2_subp0_read_sector_misses帖子，我能够理解"l2_subp1_read_sector_misses“和”this“。现在，对于事件"l2_subp0_write_sector_misses“和"l2_subp1_write_sector_misses”，我有一个类似的问题。

让我们首先以给定链接(向量添加)中的相同示例为例。

内核代码：

__global__ void AddVectors(const float* A, const float* B, float* C, int N)
{
    int blockStartIndex  = blockIdx.x * blockDim.x * N;
    int threadStartIndex = blockStartIndex + threadIdx.x;
    int threadEndIndex   = threadStartIndex + ( N * blockDim.x );
    int i;

    for( i=threadStartIndex; i<threadEndIndex; i+=blockDim.x ){
        C[i] = A[i] + B[i];
    }
}

在这里，我还将数组C从主机复制到设备。因此，C数组必须在L2缓存中(3个数组的总大小小于L2缓存的大小)。但是，我仍然看到，根据nvprof的结果，所有对C的写访问都是由于L2缓存失败而导致的。

这是预期的行为吗？是否有任何情况可以预期L2写缓存命中，或者L2缓存写访问总是被忽略？

谢谢。

Answer 1

发现L2是一个通过缓存写的，因此L2的所有写访问都被报告为L2丢失。