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

Question

Я пытаюсь сделать пример кода с постоянной памятью с 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)]

Может кто-нибудь, пожалуйста, помогите мне с постоянной памятью

Answer 1

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

__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 и извлекать адреса устройств символов в постоянной памяти, вы не должны передавать их ядру в качестве аргумента, потому что выполнение do не приведет к равномерному доступу к памяти в работающем ядре. Правильное использование постоянной памяти требует, чтобы компилятор выдавал специальные инструкции PTX, и компилятор будет делать это только тогда, когда он знает, что определенное место в глобальной памяти находится в постоянной памяти. Если вы передаете постоянный адрес памяти по значению в качестве аргумента, свойство __constant__ теряется, и компилятор не может знать, чтобы получить правильные инструкции загрузки

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