понимание результатов 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). Но все же я вижу, что все обращения на запись к C, как кэш L2, пропускаются в соответствии с результатами nvprof.

Является ли это ожидаемым поведением? Есть ли ситуации, когда мы можем ожидать, что L2 записывает кеширование кэша или L2-доступ к записи в кэш всегда становится промахом?

Спасибо.

Answer 1

Обнаружено, что L2 является кэшем для записи, поэтому все обращения на запись к L2 сообщаются как пропуски L2.