Улучшение ядра OpenCL для нейронной сети Perceptron

Как вы сказали, узким местом является взвешенная сумма. Это нетрудно, так как на каждом уровне каждый WI (рабочий элемент) выполняет множество операций ввода-вывода по сравнению с количеством арифметических операций. У меня нет опыта работы с нейронными сетями, но для меня проблема выглядит как плохой паттерн доступа к памяти на GPU.

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

void kernel Kernel(
__global const int  in_layer_size, 
__global const int  out_layer_size, 
__global const float    *in_value, 
__global const float    *in_weights, 
__global float      *out_values){

__local float buffer[SOME_SIZE];
__global const float* p_in  = in_value;
__global float* p_out = out_values;

const int 
    global_id   = get_global_id(0),
    local_id    = get_local_id(0),
    num_buffers = in_layer_size / SOME_SIZE,
    offset      = out_layer_size * global_id;

float sum = 0.0f;
for(int i=0; i < num_buffers; i++){
    buffer[local_id] = p_in[local_id];
    barrier(CLK_LOCAL_MEM_FENCE);

    //Process all data inside buffer by every WI in WG
    //...

    p_in += SOME_SIZE;
    out_values += SOME_SIZE;
    }

//...
return;

}

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

Вы можете сделать это разными способами. Но самый общий способ, не меняя поведение вашего ядра, - это повторно использовать размер вашей рабочей группы (независимо от того, что вы выбрали, или по умолчанию) и повторно использовать доступ к памяти из группы.

Я бы предложил что-то вроде этого:

ПРИМЕЧАНИЕ. Я удалил уродливые указатели для одиночных значений. OpenCL поддерживает это, и это намного проще. Зону памяти создавать не нужно, достаточно сделать clSetKernelArg(kernel, arg_index, sizeof(cl_float), &size); Где cl_float size = the_size;.

#define IN_LOCAL_SIZE 4096 //Because 16KB/4B (for each float)

void kernel perceptron(global const int in_layer_size, global const int out_layer_size, global const float *in_value, global const float* in_weights, global float* out_values)
{
    const int global_id = get_global_id(0);
    __local float in_buffer[IN_LOCAL_SIZE];

    float sum = 0.0f;
    event_t ev;
    int j;
    //For each full buffer
    for(j=0; j < (in_layer_size/IN_LOCAL_SIZE)-1; i++) {
        ev = async_work_group_copy(in_buffer, in_value+j*IN_LOCAL_SIZE, IN_LOCAL_SIZE, ev);
        wait_group_events(1,&ev);
        barrier(CLK_LOCAL_MEM_FENCE);
        for(int i=0; i < IN_LOCAL_SIZE; i++) {
            sum += in_weights[(i+j*IN_LOCAL_SIZE)*out_layer_size+global_id] * in_buffer[i];
        }
    }
    //Last one
    ev = async_work_group_copy(in_buffer, in_value+j*IN_LOCAL_SIZE, in_layer_size%IN_LOCAL_SIZE, ev);
    wait_group_events(1,&ev);
    barrier(CLK_LOCAL_MEM_FENCE);
    for(int i=0; i < in_layer_size%IN_LOCAL_SIZE; i++) {
        sum += in_weights[(i+j*IN_LOCAL_SIZE)*out_layer_size+global_id] * in_buffer[i];
    }
    out_values[global_id] = sum;
}

Однако, если выходной размер небольшой (100k, 250k, 500), то у вас будет всего 500 рабочих элементов, что не оптимально. В этом случае вам следует изменить алгоритм.

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

Но, возможно, другие подходы лучше подходят для вашей проблемы.

Вы можете добиться больших улучшений, кэшируя in_values ​​в локальной памяти. Чем меньше раз вам придется читать каждый элемент in_values ​​из глобальной памяти, тем лучше.

Я придумал решение, которое кэширует максимальное количество входных значений и считывает каждый элемент из глобальной памяти только один раз для каждой рабочей группы. Это делается путем копирования блока in_values ​​за раз, обработки его со всеми out_values ​​и перехода к следующему блоку. Существует также локальный массив с плавающей запятой, используемый для уменьшения сумм рабочих элементов каждого блока.

псевдокод:

  output elements assumed to be set to 0 already
  for each block of input values:
    cache the input block
    for each target output value:
      reset local sum to 0
      for each element this work item is responsible for:
        read the weight, multiply, and add to sum
      reduce sums to a single value, ADD value to output element

У меня еще не было возможности запустить это через профилировщик или отладчик, но я попробую, когда вернусь на свой домашний ПК. (нет инструментов opencl на моей офисной рабочей станции). Убедитесь, что ядро ​​поставлено в очередь с размером группы, равным константе GROUP_SIZE. Кроме того, создавайте только одну группу для каждого вычислительного блока на вашем устройстве.

реальный код:

//experiment with GROUP_SIZE to discover the optimal value for your device
//this needs to be equal to local_work_size passed into clEnqueueNDRangeKernel
//use a multiple of CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
//max. for most devices is 256
#define GROUP_SIZE = 64;

// IN_VALUE_CACHE_SIZE is the number of floats from in_value to copy to local memory at a time
//assuming GROUP_SIZE can be up to 256, sizeof(float)=4,  and local memory size is 32kb, full saturation can be achieved with the following:
//(32768 - (256 * 4)) /4 = 7936
//try another multiple of 1024 (6144, 4096... )if there is trouble with this value
#define IN_VALUE_CACHE_SIZE = 7936;

void kernel perceptron(global const int* in_layer_size, global const int* out_layer_size, global const float *in_value, global const float* in_weights, global float* out_values)
{
    private const int global_id = get_global_id(0);
    private const int out_layer_s = *out_layer_size;
    private const int in_layer_s = *in_layer_size;
    private const int offset = out_layer_s * global_id;

    private const int item_id = get_local_id(0);    
    private const int group_id = get_group_id(0);   
    private const int group_count = get_num_groups(0);  


    local float result_buffer[GROUP_SIZE];

    local float in_value_cache[IN_VALUE_CACHE_SIZE];
    int i,j,k;

    //init the block to 0, in case there are fewer than IN_VALUE_CACHE_SIZE values in total
    for(i=item_id; i<IN_VALUE_CACHE_SIZE; i+= GROUP_SIZE){
        in_value_cache[i] = 0.0;
    }
    barrier(CL_LOCAL_MEM_FENCE);


    private float sum = 0.0;
    event_t e;
    int copy_total = 0;
    int copy_offset;

    for(i=0; i<in_layer_s; i+=IN_VALUE_CACHE_SIZE){
        //cap the number of values to copy to local memory if loop is near the end of the input data
        copy_total = IN_VALUE_CACHE_SIZE;
        if((copy_total + i*IN_VALUE_CACHE_SIZE) > in_layer_s){
            copy_total = in_layer_s - i*IN_VALUE_CACHE_SIZE;
        }           
        //copy the next block of values
        e = async_work_group_copy(in_value_cache, in_value + i * 4, copy_total, 0);
        wait_group_events(1, &e);

        for(j=group_id; j<out_layer_s; j+=group_count){
            sum = 0.0;

            //need to reset result_buffer[item_id] as well
            //this is in case there are fewer than GROUP_SIZE input values remaining  ie copy_total < GROUP_SIZE
            result_buffer[item_id] = 0.0;

            for(k=item_id; k<copy_total; k+=GROUP_SIZE){
                sum += in_value_cache[k] * in_weights[(k+i) + j * out_layer_s];
            }
            result_buffer[item_id] = sum;

            //simple O(n) reduction can be optimized further
            if(item_id == 0){
                for(k=1;k<GROUP_SIZE;k++){
                    sum += result_buffer[k];
                }
                out_values[j] += sum;
            }
            barrier(CL_LOCAL_MEM_FENCE);
        }

    }
}

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