OpenCL 中函数的梯度

Gradient of a function in 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 x2, float x3, float x4, 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, float x2, float x3, float x4, float x5, __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);
}

请注意,没有额外的 global/local 内存访问,所有(互斥的)权重因子都是全局 ID 的函数。每个线程计算完全相同的东西(2 个 cos、1 个 exp 和一个 fes multiplications/additions),没有任何分支。三角函数/除法比multiplications/additions花费更多的时间,因此预计算项应尽可能少地使用。