Cuda - Инициализация большого массива

Question

Каков наилучший подход (эффективно) для инициализации большого массива целых чисел для gpu? Мне нужно назначить 1 для первых двух элементов и 0 для других (для сита эратосфенов).

1. cudaMemcpy
2. cudaMemset + заданное значение 2 первых элементов в ядре
3. инициализации прямых в ядре
4. СТГ еще

Примечание: размер массива является динамическим (п передается как аргумент).

Моя текущая версия:

int array = (int*) malloc(array_size); 
array[0] = 1; 
array[1] = 1; 
for (int i = 2; i < n; i++) { 
    array[i] = 0; 
} 
HANDLE_ERROR(cudaMemcpy(dev_array, array, array_size, cudaMemcpyHostToDevice)); 
kernel<<<10, 10>>>(dev_array); 

Я был бы благодарен за пример.

Answer 1

Одна возможность состоит в том, чтобы непосредственно инициализировать __device__ массив на GPU, если он имеет постоянный размер, добавив следующую декларацию в области видимости файла (то есть, вне какой-либо функции):

__device__ int dev_array[SIZE] = {1, 1}; 

Остальные элементы будут инициализирован с нули (вы можете проверить сборку PTX, чтобы убедиться в этом).

тогда, он может быть использован в ядре, как:

__global__ void kernel(void) 
{ 
    int tid = ...; 
    int elem = dev_array[tid]; 
    ... 
} 

В случае переменного размера, вы можете комбинировать cudaMalloc() с cudaMemset():

int array_size = ...; 
int *dev_array; 

cudaMalloc((void **) &dev_array, array_size * sizeof(int)); 
cudaMemset(dev_array, 0, array_size * sizeof(int)); 

затем установите первые два элемента, как те:

int helper_array[2] = {1, 1}; 
cudaMemcpy(dev_array, helper_array, 2 * sizeof(int), cudaMemcpyHostToDevice); 

Начиная с вычислительной способности 2.0 вы можете также выделить весь массив непосредственно внутри ядра с помощью функции malloc() устройства:

__global__ void kernel(int array_size) 
{ 
    int *dev_array; 
    int tid = ...; 

    if (tid == 0) { 
     dev_array = (int *) malloc(array_size * sizeof(int)); 
     if (dev_array == NULL) { 
      ... 
     } 
     memset(dev_array, 0, array_size * sizeof(int)); 
     dev_array[0] = dev_array[1] = 1; 
    } 
    __syncthreads(); 

    ... 
} 

Обратите внимание, что нити из разных блоков не знают о синхронизации барьера.

Из CUDA C Programming Guide:

CUDA в ядре malloc() функция выделяет, по меньшей мере size байт из кучи устройства и возвращает указатель на выделенную память или NULL, если существует недостаточно памяти для выполнения запроса , Возвращаемый указатель должен быть выровнен с 16-байтной границей.

К сожалению, функция calloc() не реализована, поэтому вам все равно нужно ее сменить.Выделенная память имеет срок службы CUDA контекста, но вы можете явно назвать free() от этого или последующего ядра в любой момент:

память выделяется с помощью данной CUDA нити через malloc() остается выделенным для жизни контекста CUDA, или до тех пор, пока не будет , явно выпущенный по вызову free(). Он может использоваться любыми другими потоками CUDA даже после запуска ядра.

Со всем, что сказал, я не возражал бы, что много о дополнительном cudaMemcpy(), так как это только два элемента, чтобы скопировать и было бы вероятно занимает менее 0,01% от общего времени выполнения (это легко профиль). Выберите любой способ, который сделает код четким. В противном случае это premature optimization.