Как Nsight Compute определяет / отображает метрики общей памяти?

Question

Я изучаю __shared__ память в CUDA, и я не понимаю, как Nsight Compute показывает статистику совместной памяти.

Я просматриваю эту статью (code доступно на github Nvidia здесь , но скопировано ниже для справки).

#include <stdio.h>

__global__ void staticReverse(int *d, int n)
{
  __shared__ int s[64];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

__global__ void dynamicReverse(int *d, int n)
{
  extern __shared__ int s[];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

int main(void)
{
  const int n = 64;
  int a[n], r[n], d[n];

  for (int i = 0; i < n; i++) {
    a[i] = i;
    r[i] = n-i-1;
    d[i] = 0;
  }

  int *d_d;
  cudaMalloc(&d_d, n * sizeof(int)); 

  // run version with static shared memory
  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
  staticReverse<<<1,n>>>(d_d, n);
  cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost);
  for (int i = 0; i < n; i++) 
    if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);

  // run dynamic shared memory version
  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
  dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n);
  cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);
  for (int i = 0; i < n; i++) 
    if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
}

Когда я запускаю Nsight Compute, я вижу следующую диаграмму для ядра staticReverse (dynamicReverse ядро практически идентично):

Вопрос 1 : На диаграмме показан 1 запрос к и 1 запрос из общей памяти, но почему он также показывает 0 инструкций по общей памяти? Запрос не считается инструкцией? С точки зрения этой диаграммы, что считается инструкцией к общей памяти?

Затем в исходном виде Nsight Compute показывает построчное количество различных показателей:

Вопрос 2 : Почему в строке 8 и 10 в строке 8 и 10 отображается 0 для общих транзакций L1? Я ожидал увидеть:

• Строка 8: равное количество [транзакций загрузки из глобальной памяти] и [сохранения транзакций в общую память]
• Строка 10: равное количество [загрузить транзакции из общей памяти] и [сохранить транзакции в глобальной памяти]

Вопрос 3 : Почему для каждой строки 8 и 10 используется 8 транзакций памяти?

Моя система:

• Ubuntu 18.04 LTS
• GeForce 1070 (Pascal)
• Версия CUDA: 10.2
• Версия драйвера: 440.64 0,00

Answer 1

Было бы хорошо, если бы вы могли проверить (и показать здесь) низкоуровневое представление SASS страницы Source вместе с высокоуровневым представлением CUDA- C. Исходные метрики собираются по инструкции SASS (сборка), а затем агрегируются в представлении CUDA- C. Проверка фактической сборки может быть информативной для типа инструкций, генерируемых компилятором, и может лучше объяснить данные, которые вы видите.

Не считается ли запрос инструкцией? С точки зрения этой диаграммы, что считается инструкцией совместно используемой памяти?

Запросы и инструкции - это не одно и то же. Инструкции - это фактически выполняемые инструкции сборки SASS. Запросы генерируются HW в результате выполненных инструкций, и количество запросов может варьироваться в зависимости от того, насколько хорошо ведет себя код.