性能 – 改进Perceptron神经网络的OpenCL内核

我以前做过很多OpenGL和着色器,现在,我决定试一试OpenCL.我看了一些在线教程,并开始阅读有关该主题的书籍.为了更好地理解,并且因为我认为最好的学习方法是通过智能地尝试和学习这样做时出现的问题,我决定开始为完全连接的感知器实现内核.

对于那些不知道那是什么的人,我会解释基本的想法.它是一个神经网络,其中一个层的每个神经元连接到下一层的每个神经元.每个神经元只有一个动作要执行：执行前一层所有神经元的总和,并按每个神经元的不同值加权.

这看起来很简单,在阅读了“使用OpenCL进行并行神经网络训练”之后,我按照以下方式实现了它

>每个图层都依赖于前一个图层,它们由主机按顺序运行
>为了计算一个图层,我运行我的内核,其中全局工作大小为图层内的神经元数量(例如,可能非常大,数万个).这使得所有神经元都能彼此独立地执行其总和.
>每个神经元(由其global_work_id标识)对前一层的所有神经元执行加权和.

这是我的全功能opencl内核：

/**
* @brief Computes one layer of the perceptron given the previous one and the
* weights
* The kernel is run once for each layer.
* The work items are each tasked with computing the output of a single neuron
* of the out layer.
*
* @param out_layer_size
*   Size of the output layer (number of elements in the output array that will
*   contain the result for each neuron).
* @param in_layer_size
*   Number of elements of the input layer
* @param in_value
*   Values of the neuron in the previous layer
* @param in_weights
*   Array containing the weights for each input neuron. It is organised as a
*   two dimensional matrix, written by concatenating each line in the array
*   [ w11, w12, w13, ...
*     w21, w22, w23, ...
*     ..., ..., ..., ...
*   ]
*   Where wij is the weight linking the neuron i of the input layer to the
*   neuron j of the output layer
* @param out_values
*   Computed values for the current layer
*/
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 float sum = 0.;
    for(int i=0; i < in_layer_s; i++) {
        sum += in_weights[i*out_layer_s+global_id] * in_value[i];
    }
    //out_values[global_id] = sigma(sum);
    out_values[global_id] = sum;
}

以下是我如何调用它：

queue.enqueueNDRangeKernel(kernel, cl::NullRange,cl::NDRange(number of neurons within layer),cl::NullRange);

我意识到这个内核的瓶颈是加权和的实现.如果有人可以解释我如何改进这一点以使其更快,那将是非常有帮助的.

我可能没有正确使用不同的内存区域,我基本上考虑的是我甚至不使用的本地内存.

只是为了让您了解性能(即在Nvidia GTX 660M上),我将向您展示我实现的一些时间.每个值是每层神经元的数量：

> 2500,10000,2500：0.018s~60FPS.它比我的处理器(运行在2.40GHz的Intel Core i7)快4到5倍
> 100 000,150 000,500：140s – >因为第二层中的每个神经元必须执行100 000个元素的加权和,所以我认为这并不令人惊讶.在我的处理器上运行它会产生大致相同的结果.

如你所知,瓶颈是加权汇总.这并不难,因为在每一层,每个WI(工作项)与算术运算的数量相比,正在进行大量的IO操作.我没有神经网络的经验,但对我来说问题看起来像GPU上的内存访问模式不佳.

潜在地,可以通过将WI组织到本地工作组(工作组)来解决这个问题.因为每个WI都需要处理来自prev的所有数据.层,我想WG中的所有WI都可以将一些数据加载到本地内存中,处理它们而不是下一堆数据.这将使您的算法更加缓存友好.内核的伪代码如下：

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;

}

所以,你正在用固定尺寸的窗户滑动.在&内计算数据然后进入下一个窗口. Al数据操作是独立完成的,工作项只能同时使用相同的数据.本地组的最佳大小取决于设备和内核.