我如何利用 OpenCL 本地内存进行简单的平流方案（usptream bias 或 leap frog）？

Question

我编写了一个 OpenCL 代码来使用两种不同的方案来求解平流方程：上游偏置和 Leapfrog 方案。

代码运行良好，但我想知道是否可以使用OpenCL 本地内存来优化代码。根据我的理解，当我们有可以在本地工作组内共享的内容时，本地内存很有用。但是在我的 OpenCL 内核中，除了索引之外，我没有（或者我本身想不到）任何 should/could 可以共享的东西。

上游偏置方案的内核

kernel void upstream3d(
                        const  int nx,
                        const  int ny,
                        const  int nz,
                        global float *in_p_tf,
                        global float *in_p_tn,
                        const  float u_vel,
                        const  float v_vel,
                        const  float w_vel,
                        const  float C
                      )
{
   int i    =  get_global_id(0);
   int j    =  get_global_id(1);
   int k    =  get_global_id(2);

   int idx, idx_i, idx_j, idx_k;

   int c_xi = i % nx, 
       c_yi = j % ny, 
       c_zi = k % nz,
       m_xi = (i+nx-1)%nx, 
       m_yi = (j+ny-1)%ny, 
       m_zi = (k+nz-1)%nz;


   idx      =  c_xi + c_yi * nx + c_zi * nx * ny;

   idx_i    =  m_xi + c_yi * nx + c_zi * nx * ny;
   idx_j    =  c_xi + m_yi * nx + c_zi * nx * ny;
   idx_k    =  c_xi + c_yi * nx + m_zi * nx * ny;

   in_p_tf[idx]  = in_p_tn[idx] 
                 - u_vel * C * (in_p_tn[idx] - in_p_tn[idx_i])
                 - v_vel * C * (in_p_tn[idx] - in_p_tn[idx_j])
                 - w_vel * C * (in_p_tn[idx] - in_p_tn[idx_k]);

}

Leapfrog 方案的内核

kernel void leapfrog3d(
                        const  int nx,
                        const  int ny,
                        const  int nz,
                        global float *in_p_tf,
                        global float *in_p_tn,
                        global float *in_p_tp,
                        const  float u_vel,
                        const  float v_vel,
                        const  float w_vel,
                        const  float C
                      )
{
   int i    =  get_global_id(0);
   int j    =  get_global_id(1);
   int k    =  get_global_id(2);

   int idx0, idx_i0, idx_i1, idx_j0, idx_j1, idx_k0, idx_k1;

   int p_xi = (i+1)%nx,
       p_yi = (j+1)%ny, 
       p_zi = (k+1)%nz,
       c_xi = i % nx, 
       c_yi = j % ny, 
       c_zi = k % nz,
       m_xi = (i+nx-1)%nx, 
       m_yi = (j+ny-1)%ny, 
       m_zi = (k+nz-1)%nz;


   idx0     =  c_xi + c_yi * nx + c_zi * nx * ny;

   idx_i0   =  p_xi + c_yi * nx + c_zi * nx * ny;
   idx_j0   =  c_xi + p_yi * nx + c_zi * nx * ny;
   idx_k0   =  c_xi + c_yi * nx + p_zi * nx * ny;
   
   idx_i1   =  m_xi + c_yi * nx + c_zi * nx * ny;
   idx_j1   =  c_xi + m_yi * nx + c_zi * nx * ny;
   idx_k1   =  c_xi + c_yi * nx + m_zi * nx * ny;

   in_p_tf[idx0] = in_p_tp[idx0] 
                 - u_vel * C * (in_p_tn[idx_i0] - in_p_tn[idx_i1])
                 - v_vel * C * (in_p_tn[idx_j0] - in_p_tn[idx_j1])
                 - w_vel * C * (in_p_tn[idx_k0] - in_p_tn[idx_k1]);

   in_p_tn[i + j * nx + k * nx * ny] = in_p_tn[i + j * nx + k * nx * ny] 
                                       + 0.80 * (in_p_tf[i + j * nx + k * nx * ny] 
                                       - 2.0 * in_p_tn[i + j * nx + k * nx * ny] 
                                       + in_p_tp[i + j * nx + k * nx * ny]);

}

这是我能从 OpenCL 获得的全部，还是我遗漏了什么？

谢谢。

Answer 1

有时非常有用（矩阵乘法增益约 10 倍），有时则不然。在您的情况下，它可能有点用处，但用处不大 - 您必须尝试一下。我已经在 lattice Boltzmann 中测试了一个类似的应用程序，但没有任何性能提升（可以在那里使用它进行 8/171 字节的内存传输，所以我一开始预计只有微不足道的收益）。

在 3 个平流方向中，您只能为 1 个使用共享内存，在您的情况下为 x 方向。在 upstream3d 内核中，内存传输为 4/20 字节，因此最多可以提高 20% 的性能。

工作组是沿 x 方向的一条网格单元格。每个网格单元为自己加载 in_p_tn，并从其 3 个相邻的低一个索引处加载。邻居 y-1 和 z-1 不在工作组中，因此您不能为它们使用本地内存。但是 x-1 个邻居至少是最左边的一个。所以策略是：

为工作组中的每个单元从global内存加载in_p_tn到private内存，然后用[=17=将其写入local内存阵列].最左边的线程必须额外加载其左边的邻居 x-1，因为它在工作组条带之外*。本地内存数组的维度必须为 workgroup size +1.
做一个barrier(CLK_LOCAL_MEM_FENCE);。屏障之后，local数组已经填满
从 private 内存中为当前单元格取 in_p_tn。对于 x-1 邻居，从 local 内存加载它。进行平流，将结果写入 global 内存。完成！

*如果工作组太小，这可能会消除任何性能提升。

另注：在upstream3d内核中，您当前将in_p_tn[idx]从global加载到private内存4次。制作一个private变量并加载一次，然后使用private变量。在这方面永远不要相信编译器。

我如何利用 OpenCL 本地内存进行简单的平流方案（usptream bias 或 leap frog）？

How do I exploit OpenCL local memory for simple advection scheme (usptream bias or leap frog)?

opencl