Использование CUDA Occupancy Calculator

Question

Я использую калькулятор занятости, но не могу понять, как получить регистры на поток / разделяемую память на блок. Я читаю документацию. Я использую Visual Studio. Так что в свойствах проекта в CUDA build rule-> CommandСтрока -> Дополнительные параметры Я добавляю --ptxas-options = -v. Программа прекрасно компилируется. Но я ничего не вижу. Может кто-нибудь помочь?Спасибо

Answer 1

При включенном переключателе должна быть строка в окне вывода компилятора, которая сообщает вам о количестве регистров и объеме разделяемой памяти.
Вы вообще видите что-нибудь в окне вывода компилятора?Вы можете скопировать и вставить его в вопрос?
Это должно выглядеть примерно так:

ptxas info : Used 3 registers, 2084+1060 bytes smem, 40 bytes cmem[0], 12 bytes cmem[1]

Answer 2

ответ shoosh , вероятно, самый простой способ найти регистр и использование общей памяти.Убедитесь, что вы сначала смотрите на панель вывода (выберите «Вывод» в выпадающем меню «Вид»), а затем пересоберите.Компилятор должен предоставить вам информацию ptxas для всех ядер в области вывода, как вы можете видеть на рисунке ниже ...

Answer 3

Попробуйте это простое правило:

Все локальные переменные, такие как int a, float b и т. Д. В вашем ядре, хранятся в регистрах.Это только в том случае, если локальные переменные в вашем коде остаются в пределах доступных регистров в многопроцессорном режиме, См. Ограничения .Однако, если вы объявите тысячу целых чисел, таких как int и [1000] , тогда a не будет сохранено в регистрах, а будет сохранено в локальной памяти (DRAM).

Количество общей памяти, используемой в коде вашего ядра, равно Shared Memory / Block.Например, если вы определяете __shared__ float shMem[256], то вы используете 256 * 4 (размер с плавающей запятой) = 1024 байта разделяемой памяти.

В следующем примере кода (он не будет работать должным образом, например,) используется 9 32-битные регистры на поток, которые являются: int xIndex, yIndex, Idx, shY, shX, aLocX, aLocY и float t, temp .Код использует 324 байта общей памяти на блок, как BLOCK_DIM = 16.

__global__ void averageFilter (unsigned char * outImage,
                           int imageWidth,
                           int imageHeight,
                           cuviPoint2 loc){


    unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;
    unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;
    unsigned int Idx = yIndex*imageWidth + xIndex;
    float t = INC;


    if(xIndex>= imageWidth|| yIndex>=imageHeight)
        return;


    else if(xIndex==0 || xIndex== imageWidth-1 || yIndex==0 || yIndex==imageHeight-1){

          for (int i=-1; i<=1; i++)
             for (int j=-1; j<=1; j++)
                 t+= tex1Dfetch(texMem,Idx+i*imageWidth+j);
                    outImage[Idx] = t/6;

          }


    __shared__ unsigned char shMem[BLOCK_DIM+2][BLOCK_DIM+2];


    unsigned int shY = threadIdx.y + 1;
    unsigned int shX = threadIdx.x + 1;


   if (threadIdx.x==0 || threadIdx.x==BLOCK_DIM-1 || threadIdx.y==0 || threadIdx.y==BLOCK_DIM-1){


 for (int i=-1; i<=1; i++)
      for (int j=-1; j<=1; j++)
        shMem[shY+i][shX+j]=  tex1Dfetch(texMem,Idx+i*imageWidth+j);

    }
    else
    shMem[shY][shX] =  tex1Dfetch(texMem,Idx);

     __syncthreads();     



if(xIndex==0 || xIndex== imageWidth-1 || yIndex==0 || yIndex==imageHeight-1)
        return;     

  int aLocX = loc.x, aLocY = loc.y;

    float temp=INC;

      for (int i=aLocY; i<=aLocY+2; i++)
         for (int j=aLocX; j<=aLocX+2; j++)
        temp+= shMem[shY+i][shX+j];

        outImage[Idx] = floor(temp/9);

}

Answer 4

Еще один способ узнать эту информацию - использовать визуальный профилировщик или nvidia для параллельного просмотра.