Доступ к глобальной памяти в CUDA происходит медленно

У меня есть ядро CUDA, выполняющее некоторые вычисления на локальной переменной (в регистре), и после того, как оно вычисляется, его значение записывается в глобальный массив p:

__global__ void dd( float* p, int dimX, int dimY, int dimZ  )
{
    int 
        i = blockIdx.x*blockDim.x + threadIdx.x,
        j = blockIdx.y*blockDim.y + threadIdx.y,
        k = blockIdx.z*blockDim.z + threadIdx.z,
        idx = j*dimX*dimY + j*dimX +i;   

    if (i >= dimX || j >= dimY || k >= dimZ)
    {
        return;
    }   

    float val = 0;

    val = SomeComputationOnVal();

     p[idx ]=  val;
    __syncthreads();    

} 

К сожалению, эта функция выполняется очень медленно.

Тем не менее, он работает очень быстро, если я делаю это:

  __global__ void dd(   float* p, int dimX, int dimY, int dimZ  )
    {
        int 
            i = blockIdx.x*blockDim.x + threadIdx.x,
            j = blockIdx.y*blockDim.y + threadIdx.y,
            k = blockIdx.z*blockDim.z + threadIdx.z,
            idx = j*dimX*dimY + j*dimX +i;   

        if (i >= dimX || j >= dimY || k >= dimZ)
        {
            return;
        }   

        float val = 0;

        //val = SomeComputationOnVal();

         p[idx ]=  val;
        __syncthreads();    

    } 

Он также работает очень быстро, если я делаю это:

__global__ void dd( float* p, int dimX, int dimY, int dimZ  )
{
    int 
        i = blockIdx.x*blockDim.x + threadIdx.x,
        j = blockIdx.y*blockDim.y + threadIdx.y,
        k = blockIdx.z*blockDim.z + threadIdx.z,
        idx = j*dimX*dimY + j*dimX +i;   

    if (i >= dimX || j >= dimY || k >= dimZ)
    {
        return;
    }   

    float val = 0;

    val = SomeComputationOnVal();

  //   p[idx ]=  val;
    __syncthreads();    

} 

Так что я в замешательстве, и понятия не имею, как решить эту проблему. Я использовал шаг NSight и не обнаружил нарушений доступа.

Вот как я запускаю ядро (dimX: 924; dimY: 16: dimZ: 1120):

dim3 
      blockSize(8,16,2),
      gridSize(dimX/blockSize.x+1,dimY/blockSize.y, dimZ/blockSize.z);
float* dev_p;       cudaMalloc((void**)&dev_p, dimX*dimY*dimZ*sizeof(float));

dd<<<gridSize, blockSize>>>(     dev_p,dimX,dimY,dimZ);

Кто-нибудь может дать несколько советов? Потому что для меня это не имеет большого смысла. Все вычисления val быстры ,и последний шаг должен двигаться valpв. pникогда не участвует в вычислениях, и он появляется только один раз. Так почему же это так медленно?

Вычисления в основном представляют собой цикл по матрице 512 X 512. Я бы сказал, что это довольно приличный объем вычислений.

1 ответ

  1. Вычисления, выполняемые в SomeComputationOnVal, чрезвычайно дороги. Каждый поток считывает по крайней мере 1 МБ данных, которые находятся вне кэша (или в L2 в лучшем случае для небольшой части должны варьироваться в небольшом диапазоне), который составляет для вашего запуска около 16 ТБ данных. Даже на высокопроизводительном gpu, это займет около 2 минут, чтобы запустить, как минимум. Не говоря уже обо всем, что могло бы замедлить это.

    Ваша функция не записывает данные в глобальную память и не имеет эффекта границ. Компилятор может решить оптимизировать вызов метода, если выходные данные не используются.

    Следовательно, случаи два и три, не выполняющие вычисления, очень быстры. Запись 64 Мб на gpu памяти, с coesced потоков очень быстро (диапазон миллисекунд).

    Вы можете проверить сгенерированный ptx, чтобы увидеть, оптимизируется ли код. Используйте опцию —keep в nvcc и ищите ptx файлы.