OpenCL: примеры сокращения и сохранения объектов памяти / преобразования кода cuda в openCL

Question

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

__kernel void sum(__global const short *A,__global unsigned long  *C,uint size, __local unsigned long *L) {
            unsigned long sum=0;
            for(int i=get_local_id(0);i<size;i+=get_local_size(0))
                    sum+=A[i];
            L[get_local_id(0)]=sum;

            for(uint c=get_local_size(0)/2;c>0;c/=2)
            {
                    barrier(CLK_LOCAL_MEM_FENCE);
                    if(c>get_local_id(0))
                            L[get_local_id(0)]+=L[get_local_id(0)+c];

            }
            if(get_local_id(0)==0)
                    C[0]=L[0];
            barrier(CLK_LOCAL_MEM_FENCE);
}

Это выглядит правильно?Третий аргумент "размер", это должен быть локальный размер работы или глобальный размер работы?

Я установил свои аргументы следующим образом:

clSetKernelArg(ocReduce, 0, sizeof(cl_mem), (void*) &DevA);
clSetKernelArg(ocReduce, 1, sizeof(cl_mem), (void*) &DevC); 
clSetKernelArg(ocReduce, 2, sizeof(uint),   (void*) &size);  
clSetKernelArg(ocReduce, 3, LocalWorkSize * sizeof(unsigned long), NULL); 

Первый аргумент, который является входом, я пытаюсь сохранить из вывода ядра, запущенного до него.

clRetainMemObject(DevA);
clEnqueueNDRangeKernel(hCmdQueue[Plat-1][Dev-1], ocKernel, 1, NULL, &GlobalWorkSize, &LocalWorkSize, 0, NULL, NULL);
//the device memory object DevA now has the data to be reduced

clEnqueueNDRangeKernel(hCmdQueue[Plat-1][Dev-1], ocReduce, 1, NULL, &GlobalWorkSize, &LocalWorkSize, 0, NULL, NULL);
clEnqueueReadBuffer(hCmdQueue[Plat-1][Dev-1],DevRE, CL_TRUE, 0, sizeof(unsigned long)*512,(void*) RE , 0, NULL, NULL);

Сегодня я планирую попытаться преобразовать следующий пример сокращения cuda в openCL.

__global__ voidreduce1(int*g_idata, int*g_odata){
extern __shared__ intsdata[];

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;
sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
__syncthreads();


for(unsigned int s=blockDim.x/2; s>0; s>>=1) {
if (tid < s) {
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}

// write result for this block to global mem
if(tid == 0) g_odata[blockIdx.x] = sdata[0];
}

Существует более оптимизированный (полностью развернутый + несколько элементов на поток).

http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/reduction/doc/reduction.pdf

Возможно ли это с помощью openCL?

Гризли дал мне этот совет на днях:

"... использовать ядро ​​редукции, которое работает с n элементом и сводит их к чему-то вроде n / 16 (или любому другомучисло). Затем вы итеративно вызываете это ядро, пока не получите один элемент, который является вашим результатом "

Я тоже хочу попробовать это, но я не знаю точно, с чего начать,и я хочу сначала просто заставить что-то работать.

Answer 1

Первый код сокращения, который вы дали, должен работать, пока над сокращением работает только одна рабочая группа (поэтому get_global_size(0) == get_local_size(0)). В этом случае аргументом size ядра будет число элементов в A (которое не имеет реальной корреляции ни с глобальным, ни с локальным рабочим размером). Несмотря на то, что это работоспособное решение, кажется невероятно расточительным позволить большей части gpu бездействовать при выполнении сокращения, именно поэтому я предложил итеративно вызывать ядро ​​сокращения. Это стало бы возможным только с небольшими изменениями в коде:

__kernel void sum(__global const short *A, __global unsigned long  *C, uint size, __local unsigned long *L) {
        unsigned long sum=0;
        for(int i=get_global_id(0); i < size; i += get_global_size(0))
                sum += A[i];
        L[get_local_id(0)]=sum;

        for(uint c=get_local_size(0)/2;c>0;c/=2)
        {
                barrier(CLK_LOCAL_MEM_FENCE);
                if(c>get_local_id(0))
                        L[get_local_id(0)]+=L[get_local_id(0)+c];

        }
        if(get_local_id(0)==0)
                C[get_group_id(0)]=L[0];
        barrier(CLK_LOCAL_MEM_FENCE);
}

Вызов этого значения с GlobalWorkSize меньшим, чем size (например, 4) уменьшит ввод в A с коэффициентом 4*LocalWorkSize, который можно повторять (используя выходной буфер в качестве входного для следующий вызов sum с другим выходным буфером. На самом деле это не совсем так, поскольку вторая (и все последующие) итерации требуют, чтобы A был типа global const unsigned long*, так что вам действительно понадобятся ядра , но вы поняли.

Что касается примера сокращения cuda: зачем вам его конвертировать, он работает в основном точно так же, как и версия opencl, которую я разместил выше, за исключением уменьшения только на жестко заданный размер за итерацию (2*LocalWorkSize вместо size/GlobalWorkSize*LocalWorkSize).

Лично я использую практически тот же подход для сокращения, хотя я разделил ядро ​​на две части и использую путь только с использованием локальной памяти для последней итерации:

__kernel void reduction_step(__global const unsigned long* A, __global unsigned long  * C, uint size) {
        unsigned long sum=0;
        for(int i=start; i < size; i += stride)
                sum += A[i];
        C[get_global_id(0)]= sum;
}

На последнем шаге была использована полная версия, которая сокращает число внутри рабочей группы. Конечно, вам понадобится вторая версия reduction step, взяв global const short*, и этот код является непроверенной адаптацией вашего кода (я не могу опубликовать свою собственную версию, к сожалению). Преимущество этого подхода заключается в значительно меньшей сложности ядра, выполняющего большую часть работы, и меньшем количестве wasted work из-за расходящихся ветвей. Что сделало это немного быстрее, чем другой вариант. Однако у меня нет результатов ни для самой новой версии компилятора, ни для новейшего оборудования, так что эта точка может быть или не быть правильной больше (хотя я подозреваю, что это может произойти из-за уменьшенного количества расходящихся ветвей).

Теперь для статьи, на которую вы ссылаетесь: Конечно, можно использовать оптимизацию, предложенную в этой статье в opencl, за исключением использования шаблонов, которые не поддерживаются в opencl, поэтому размеры блоков должны быть жестко закодированы. Конечно, версия opencl уже делает несколько добавлений для каждого ядра, и, если вы будете придерживаться подхода, который я упомянул выше, на самом деле не будет выгоды от развертывания сокращения через локальную память, поскольку это делается только на последнем шаге, который не должен занимать значительная часть всего времени расчета для достаточно большого вклада. Кроме того, я нахожу отсутствие синхронизации в развернутой реализации немного проблематичным. Это работает только потому, что все потоки в этой части принадлежат одной и той же основе. Это, однако, не обязательно верно при выполнении на любом оборудовании, кроме текущих карт nvidia (будущих карт nvidia, карт памяти amd и процессоров (хотя я думаю, что это должно работать для текущих карт памяти amd и текущих реализаций процессора, но я не обязательно буду рассчитывать на это)), так что я бы держался подальше от этого, если бы мне не требовался абсолютный последний бит скорости для сокращения (а затем все еще предоставлял бы общую версию и переключался на нее, если я не узнаю аппаратное обеспечение или что-то в этом роде).

Answer 2

Ядро редукции выглядит правильным для моих глаз.При уменьшении размер должен быть числом элементов входного массива A.Код накапливает частичную сумму для каждого потока в sum, затем выполняет уменьшение локальной памяти (совместно используемой памяти) и сохраняет результат в C.Вы получите одну частичную сумму в C за местную рабочую группу.Либо вызовите ядро ​​второй раз с одной рабочей группой, чтобы получить окончательный ответ, либо накопите частичные результаты на хосте.