1

取自样本的代码。我用它创建了一个项目，它可以工作，但我不明白某些部分。

举例来说，假设我有一个 32x32 矩阵，有 36 个工作项，因此我假设 get_global_id(0) 从 0 -> 35 变为 size = MATRIX_DIM/4 = 8。

__kernel void transpose(__global float4 *g_mat,
   __local float4 *l_mat, uint size) {

   __global float4 *src, *dst;

   /* Determine row and column location */
   int col = get_global_id(0);
   int row = 0;
   while(col >= size) {
      col -= size--;
      row++;
   }
   col += row;
   size += row;

   /* Read source block into local memory */
   src = g_mat + row * size * 4 + col;
   l_mat += get_local_id(0)*8;

在clEnqueueNDRangeKernel调用中，根据手册将 arglocal_work_size设置为NULL让编译器或其他东西弄清楚：

local_work_size can also be a NULL value in which case the OpenCL implementation will determine how to be break the global work-items into appropriate work-group instances.

但我不明白乘以 8，它为我想的工作组提供了本地内存的地址偏移量。有人可以解释一下吗？

   l_mat[0] = src[0];
   l_mat[1] = src[size];
   l_mat[2] = src[2*size];
   l_mat[3] = src[3*size];

   /* Process block on diagonal */
   if(row == col) {
      src[0] =
         (float4)(l_mat[0].x, l_mat[1].x, l_mat[2].x, l_mat[3].x);
      src[size] =
         (float4)(l_mat[0].y, l_mat[1].y, l_mat[2].y, l_mat[3].y);
      src[2*size] =
         (float4)(l_mat[0].z, l_mat[1].z, l_mat[2].z, l_mat[3].z);
      src[3*size] =
         (float4)(l_mat[0].w, l_mat[1].w, l_mat[2].w, l_mat[3].w);
   }
   /* Process block off diagonal */
   else {
      /* Read destination block into local memory */
      dst = g_mat + col * size * 4 + row;
      l_mat[4] = dst[0];
      l_mat[5] = dst[size];
      l_mat[6] = dst[2*size];
      l_mat[7] = dst[3*size];

      /* Set elements of destination block */
      dst[0] =
         (float4)(l_mat[0].x, l_mat[1].x, l_mat[2].x, l_mat[3].x);
      dst[size] =
         (float4)(l_mat[0].y, l_mat[1].y, l_mat[2].y, l_mat[3].y);
      dst[2*size] =
         (float4)(l_mat[0].z, l_mat[1].z, l_mat[2].z, l_mat[3].z);
      dst[3*size] =
         (float4)(l_mat[0].w, l_mat[1].w, l_mat[2].w, l_mat[3].w);

      /* Set elements of source block */
      src[0] =
         (float4)(l_mat[4].x, l_mat[5].x, l_mat[6].x, l_mat[7].x);
      src[size] =
         (float4)(l_mat[4].y, l_mat[5].y, l_mat[6].y, l_mat[7].y);
      src[2*size] =
         (float4)(l_mat[4].z, l_mat[5].z, l_mat[6].z, l_mat[7].z);
      src[3*size] =
         (float4)(l_mat[4].w, l_mat[5].w, l_mat[6].w, l_mat[7].w);
   }
}
4

2 回答 2

1

l_mat用作临时缓冲区，用于存储要为所有工作项反转的两个矩阵分量。

因此，对于每个工作项，它需要存储2 * 4 个float4s，因此：offset = get_local_id(0)*2*4 = get_local_id(0)*8。

于 2012-11-07T18:20:32.940 回答
1

l_mat被用作工作组中线程的本地存储。具体来说，它之所以被使用，是因为访问本地内存比访问全局内存要快几个数量级。

每个线程需要 8float4秒。进行以下指针运算

l_mat += get_local_id(0)*8;

移动l_mat每个线程的指针，使其不与其他线程的数据重叠。

这可能会导致错误，因为未指定 local_size 并且我们无法确保 的大小l_mat足以存储每个线程的值。

于 2012-11-07T18:25:32.310 回答