performance - 改进感知器神经网络的 OpenCL 内核

Question

我之前一直在做很多 OpenGL 和着色器，现在，我决定尝试一下 OpenCL。我看了一些在线教程，并开始阅读有关该主题的书籍。为了更好地理解，并且因为我相信最好的学习方式是聪明地尝试并从这样做时出现的问题中学习，我决定开始为一个完全连接的感知器实现一个内核。

对于那些不知道那是什么的人，我将解释其基本思想。它是一个神经网络，其中一层的每个神经元都连接到下一层的每个神经元。每个神经元只有一个 Action 要执行:执行来自前一层的所有神经元的总和，由每个神经元的不同值加权。

这看起来很容易实现，在阅读了论文“使用 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 的英特尔酷睿 i7)快 4 到 5 倍

100 000, 100 000, 500: 140s -> 我想这并不奇怪，因为第二层中的每个神经元都必须执行 100 000 个元素的加权和。在我的处理器上运行它会产生大致相同的结果。

Answer 1

正如您所说，瓶颈是加权总和。这并不难，因为在每一层，与算术运算的数量相比，每个 WI(工作项)都在执行大量 IO 操作。我没有神经网络方面的经验，但对我来说，问题看起来像是 GPU 上糟糕的内存访问模式。

潜在地，这可以通过将您的 WI 组织到本地 WG(工作组)中来解决。由于每个 WI 都需要处理上一个的所有数据。层，我猜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;

}

因此，您正在滑动固定大小的窗口并计算其中的数据，然后转到下一个窗口。所有数据操作都是独立完成的，工作项只同时使用相同的数据。本地组的最佳大小取决于设备和内核。