Ошибка постоянной памяти CUDA

Я пытаюсь сделать пример кода с постоянной памятью с помощью CUDA 5.5. У меня есть 2 постоянных массива размером 3000 каждый. У меня есть еще один глобальный массив X размера N. Я хочу вычислить

Y[tid] = X[tid]*A[tid%3000] + B[tid%3000]

Вот код.

#include <iostream>
#include <stdio.h>
using namespace std;

#include <cuda.h>



__device__ __constant__ int A[3000];
__device__ __constant__ int B[3000];


__global__ void kernel( int *dc_A, int *dc_B, int *X, int *out, int N)
{
    int tid = threadIdx.x + blockIdx.x*blockDim.x;
    if( tid<N )
    {
        out[tid] = dc_A[tid%3000]*X[tid] + dc_B[tid%3000];
    }

}

int main()
{
    int N=100000;

    // set affine constants on host
    int *h_A, *h_B ; //host vectors
    h_A = (int*) malloc( 3000*sizeof(int) );
    h_B = (int*) malloc( 3000*sizeof(int) );
    for( int i=0 ; i<3000 ; i++ )
    {
        h_A[i] = (int) (drand48() * 10);
        h_B[i] = (int) (drand48() * 10);
    }

    //set X and Y on host
    int * h_X = (int*) malloc( N*sizeof(int) );
    int * h_out = (int *) malloc( N*sizeof(int) );
    //set the vector
    for( int i=0 ; i<N ; i++ )
    {
        h_X[i] = i;
        h_out[i] = 0;
    }

    // copy, A,B,X,Y to device
    int * d_X, *d_out;
    cudaMemcpyToSymbol( A, h_A, 3000 * sizeof(int) ) ;
    cudaMemcpyToSymbol( B, h_B, 3000 * sizeof(int) ) ;

    cudaMalloc( (void**)&d_X, N*sizeof(int) ) );
    cudaMemcpy( d_X, h_X, N*sizeof(int), cudaMemcpyHostToDevice ) ;
    cudaMalloc( (void**)&d_out, N*sizeof(int) ) ;



    //call kernel for vector addition
    kernel<<< (N+1024)/1024,1024 >>>(A,B, d_X, d_out, N);
    cudaPeekAtLastError() ;
    cudaDeviceSynchronize() ;


    // D --> H
    cudaMemcpy(h_out, d_out, N * sizeof(int), cudaMemcpyDeviceToHost ) ;


    free(h_A);
    free(h_B);


    return 0;
}

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

Coalescing of the CUDA commands output is off.
[Thread debugging using libthread_db enabled]
[New Thread 0x7ffff5c5b700 (LWP 31200)]

Помогите пожалуйста с постоянной памятью


c++ cuda cuda-gdb gpu-constant-memory
person mkuse    schedule 07.10.2013    source источник
comment
Показанный вами отладчик не является ошибкой. Это абсолютно нормальный вывод при запуске нового отладчика.   -  person talonmies    schedule 07.10.2013

Ответы (1)


arrow_upward
4
arrow_downward

Здесь есть несколько проблем. Вероятно, проще начать с демонстрации «правильного» способа использования этих двух константных массивов, а затем объяснить, почему то, что вы сделали, не работает. Итак, ядро ​​должно выглядеть так:

__global__ void kernel(int *X, int *out, int N)
{
    int tid = threadIdx.x + blockIdx.x*blockDim.x;
    if( tid<N )
    {
        out[tid] = A[tid%3000]*X[tid] + B[tid%3000];
    }
}

т.е. не пытайтесь передавать A и B ядру. Причины следующие:

  1. Несколько сбивает с толку то, что A и B в коде хоста не являются допустимыми адресами памяти устройства. Они являются хост-символами, которые предоставляют перехватчики для поиска символа устройства во время выполнения. Незаконно передавать их ядру. Если вам нужен адрес памяти их устройства, вы должны использовать cudaGetSymbolAddress для его получения во время выполнения.
  2. Даже если вы вызвали cudaGetSymbolAddress и получили адреса символьных устройств в постоянной памяти, вам не следует передавать их ядру в качестве аргумента, потому что это не даст унифицированного доступа к памяти в работающем ядре. Правильное использование постоянной памяти требует, чтобы компилятор выдавал специальные инструкции PTX, и компилятор будет делать это только тогда, когда он знает, что конкретная ячейка глобальной памяти находится в постоянной памяти. Если вы передаете постоянный адрес памяти по значению в качестве аргумента, свойство __constant__ теряется, и компилятор не может знать, как создать правильные инструкции загрузки.

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

person talonmies    schedule 07.10.2013
comment
Большое спасибо, это сработало. Также спасибо, что поделились своим мнением о производительности. На самом деле, я пытаюсь научиться использовать постоянную память. Но запомнит ваш совет для производственных кодов - person mkuse; 08.10.2013