OpenCL 中函数的梯度

如何解决OpenCL 中函数的梯度

我正在尝试使用 OpenCL，但遇到了一个问题，可以简化如下。我确定这是一个常见问题，但我找不到很多可以向我展示通常如何完成的参考资料或示例假设例如你有一个函数（用 CStyle 语法编写）

float function(float x1,float x2,float x3,float x4,float x5)
{
   return sin(x1) + x1*cos(x2) + x3*exp(-x3) + x4 + x5;
}

我也可以将这个函数的梯度实现为

void functionGradient(float x1,float x5,float gradient[])
{
   gradient[0] = cos(x1) + cos(x2);
   gradient[1] = -sin(x2);
   gradient[2] = exp(-x3) - x3*exp(-x3);
   gradient[3] = 1.0f;
   gradient[4] = 1.0f;
}

现在我正在考虑实现一个可以做同样事情的 OpenCL C 内核函数，因为我想加快速度。我想到的唯一方法是为每个工作单元分配梯度的一个组件，但随后我需要在代码中放置一堆 if 语句来确定哪个工作单元正在计算哪个组件，哪个不是总体来说不错，因为有分歧。

因此这里的问题是，一般如何解决此类问题？例如，我知道 GPU 上的梯度下降实现，例如使用反向传播的机器学习。所以我想知道通常会做些什么来避免代码的分歧。

跟进建议

我正在考虑如下可能的 SIMD 兼容实现：

/*
Pseudo OpenCL-C code
here weight is a 5x5 array containing weights in {0,1} masking the relevant
computation
*/
__kernel void functionGradient(float x1,__global float* weight,__global* float gradient)
{
   size_t threadId = get_global_id(0);
   gradient[threadId] = 
      weight[5*threadId]*(cos(x1) + cos(x2)) +
      weight[5*threadId + 1]*(-sin(x2)) +
      weight[5*threadId + 2]*(exp(-x3) - x3*exp(x3)) +
      weight[5*threadId + 3] + weight[5*threadId + 4];
   barrier(CLK_GLOBAL_MEM_FENCE);
}

解决方法

如果您的渐变函数只有 5 个组件，那么以一个线程执行一个组件的方式对其进行并行化是没有意义的。正如您提到的，如果每个组件的数学结构不同（多指令多数据，MIMD），GPU 并行化就不起作用。

如果您需要在 100k 不同坐标处计算 5 维梯度，那么每个线程将为每个坐标执行所有 5 个组件，并且并行化将高效工作。

在反向传播示例中，您有一个具有数千维的梯度函数。在这种情况下，您确实会并行化梯度函数本身，以便一个线程计算梯度的一个分量。然而，在这种情况下，所有梯度分量都具有相同的数学结构（全局内存中具有不同的权重因子），因此不需要分支。每个梯度分量都是相同的等式，但编号不同（单指令多数据，SIMD）。 GPU 旨在仅处理 SIMD；这也是为什么与 CPU（可以执行 MIMD，~2-3TFLOPs @ 150W）相比，它们如此节能（~30TFLOPs @ 300W）。

最后，请注意反向传播/神经网络专门设计为 SIMD。并非您遇到的每个新算法都可以以这种方式并行化。

回到您的 5 维梯度示例：有多种方法可以使其与 SIMD 兼容而无需分支。特别是 bit-maskimg：您将计算 2 个余弦（对于组件 1 通过余弦表示正弦）和一个指数，并将所有项相加，并在每个项前面加上一个因子。不需要的项乘以因子 0。最后，因子是组件 ID 的函数。但是，如上所述，这仅在您有数千到数百万个维度时才有意义。

编辑：这里是带有位掩码的 SIMD 兼容版本：

kernel void functionGradient(const global float x1,const global float x2,const global float x3,const global float x4,const global float x5,global float* gradient) {
    const float gid = get_global_id(0);
    const float cosx1 = cos(x1);
    const float cosx2 = cos((gid!=1)*x2+(gid==1)*3.1415927f);
    const float expmx3 = exp(-x3);
    gradient[gid] = (gid==0)*cosx1 + (gid<=1)*cosx2 + (gid==2)*(expmx3-x3*expmx3) + (gid>=3);
}

请注意，没有额外的全局/本地内存访问，并且所有（互斥的）加权因子都是全局 ID 的函数。每个线程计算完全相同的东西（2 cos、1 exp 和 fes 乘法/加法），没有任何分支。三角函数/除法比乘法/加法花费更多的时间，因此通过预先计算的术语应尽可能少地使用。