将神经网络层内核组合成一个内核 CUDA答案

【问题标题】：Combine neural network layer kernels into one kernel CUDA将神经网络层内核组合成一个内核 CUDA
【发布时间】：2021-12-11 02:32:12
【问题描述】：

我正在研究神经网络的 CUDA 实现，我想知道如何进一步优化全连接层中的计算。

我当前用于神经网络中全连接层的 CUDA 内核包括以下步骤：

将输出神经元累加器 (input) 设置为 0
将上一层 (in) 的输出数据与当前层的 weights 相乘，并将结果相加到累加器中
通过对累积的数据应用激活函数来计算当前层 (out) 的输出

这些是单层神经网络中的一般步骤，但目前（见下文）实现为单独的内核。对于较小的输出大小（例如outSizeX 等于 10），第一步和第三步相对较慢，尤其是结合启动三个内核。

因此，我的问题是：如何将这三个内核组合成一个内核来执行上述所有三个步骤？

// Step 1
__global__ void set_to_zero_cuda(float *__restrict__ input, int outSizeX)
{
  int i = threadIdx.x + blockIdx.x * blockDim.x;

  if (i >= outSizeX)
    return;

  input[i] = 0;
}

// Step 2
__global__ void activate_cuda_fc(const float *__restrict__ in, float *__restrict__ input, const float *__restrict__ weights,
                                 int totalInSize, int outSizeX, int weightSizeX)
{
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  int y = blockIdx.y * blockDim.y + threadIdx.y;

  int nx = blockDim.x * gridDim.x;
  int ny = blockDim.y * gridDim.y;

  for (int n = x; n < outSizeX; n += nx)
  {
    for (int i = y; i < totalInSize; i += ny)
    {
      atomicAdd(&input[n], in[i] * weights[i + n * weightSizeX]);
    }
  }
}

// Step 3
__global__ void perform_activation_function_cuda_fc(float *__restrict__ out, float *input,
                                                    int outSizeX)
{
  int i = threadIdx.x + blockIdx.x * blockDim.x;

  if (i >= outSizeX)
    return;

  out[i] = activator_function_cuda(input[i]);
}

作为参考，当前配置文件如下所示：

【问题讨论】：

标签： c++ machine-learning deep-learning neural-network cuda

【解决方案1】：

因此，我的问题是：如何将这三个内核组合成一个内核来执行上述所有三个步骤？

除非您使用线性激活函数，否则您不能像这样“折叠”一系列完全连接的层。

将权重和偏差应用于每一层的输入正是一种可并行化的线性代数运算，它是 GPU 的基础。但是，为了有效地工作，您需要在启动图层之前准备好图层的所有输入。阻止批量执行该操作的任何操作都会立即损害性能。

同时，由于大多数激活函数都引入了非线性，它们不能直接嵌入到线性代数过程中，因此您别无选择，只能分别执行。

但是，您发布的代码仍有很多好处。正如我所说，应用权重和偏差是 GPU 的基础。事实上，它实际上与通过矩阵转换向量完全相同，但您是以一种相当迂回的方式进行的。使用现成的函数 M*V 函数，例如 cublasSgemv()，很可能会给您带来一些立竿见影的好处。

附录：

如果您正在使用线性激活函数，那么您实际上是在执行y = A3 * L3 * A2 * L2 * A1 * L1 * x，其中Ln 是与层关联的矩阵，而激活函数An 只是标量。您可以提前将所有A 和L 预乘在一起，并将其视为一个大矩阵乘法。

【讨论】：