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

Question

Постоянная ошибка памяти 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)]

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

1

c++cuda cuda-gdb gpu-constant-memory

Решение

Другие решения

Других решений пока нет …

Источник

Accepted Answer

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

__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 ядру. Причины следующие:

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

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

4