определить переменный размер массива в локальной памяти, используя CUDA

Question

Возможно ли создать список, массив, что-то в функции device с размером списка / массива, являющегося параметром в вызове ... или глобальной переменной, которая инициализируется во время вызова?

Я бы хотел, чтобы что-то вроде этого списка работало:

unsigned int size1;

__device__ void function(int size2) {

    int list1[size1];
    int list2[size2];
}

Можно ли сделать что-то умное, чтобы сделать что-то подобное этой работе?

Answer 1

Существует 1 способ выделить динамический объем разделяемой памяти - использовать третий параметр ядра запуска:

__global__ void kernel (int * arr) 
{
    extern __shared__ int buf []; // size is not stated
    // copy data to shared mem:
    buf[threadIdx.x] = arr[blockIdx.x * blockDim.x + threadIdx.x];
    // . . . 
}
// . . . 
// launch kernel, set size of shared mem in bytes (k elements in buf):
kernel<<<grid, threads, k * sizeof(int)>>> (arr);

Для многих массивов есть хак:

__device__ void function(int * a, int * b, int k) // k elements in first list
{
    extern __shared__ int list1 [];
    extern __shared__ int list2 []; // list2 points to the same point as list1 does

    list1 [threadIdx.x] = a[blockIdx.x * blockDim.x + threadIdx.x];
    list2 [k + threadIdx.x] = b[blockIdx.x * blockDim.x + threadIdx.x];
    // . . .
}

Необходимо учитывать: память выделена для всего блока.

Answer 2

Конечно, это возможно!

Взгляните на исходный код проекта: http://code.google.com/p/cuda-grayscale/

Эта функция вызывается из main () и выполняет преобразование оттенков серого на * 1006.* gpu_image в зависимости от ширины и высоты: cuda_grayscale (gpu_image, ширина, высота, сетка, блок);

Если вы немного покопаетесь, вы найдете реализациюkernel_gpu.cu:

__global__ void grayscale(float4* imagem, int width, int height)
{
    const int i = blockIdx.x * (blockDim.x * blockDim.y) + blockDim.x * threadIdx.y + threadIdx.x;

    if (i < width * height)
    {
        float v = 0.3 * imagem[i].x + 0.6 * imagem[i].y + 0.1 * imagem[i].z;
        imagem[i] = make_float4(v, v, v, 0);
    }
}

Answer 3

, если вы знаете, какие значения размера вы можете ожидать, подумайте об использовании шаблонов C ++.Вместе с препроцессором boost вы можете легко создавать несколько экземпляров / точек входа.

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

Я могу предоставить вам ссылку, если вы хотите увидеть пример