OpenCL Matrix乘法Altera示例

我对OpenCL非常陌生，正在通过Altera OpenCL示例。在他们的矩阵乘法的例子中，他们已经使用了块的概念，其中输入矩阵的维数是块大小的倍数。下面的代码：OpenCL Matrix乘法Altera示例

void matrixMult(// Input and output matrices 
     __global float *restrict C, 
     __global float *A, 
     __global float *B, 
     // Widths of matrices. 
     int A_width, int B_width) 
{ 
    // Local storage for a block of input matrices A and B 
    __local float A_local[BLOCK_SIZE][BLOCK_SIZE]; 
    __local float B_local[BLOCK_SIZE][BLOCK_SIZE]; 

    // Block index 
    int block_x = get_group_id(0); 
    int block_y = get_group_id(1); 

    // Local ID index (offset within a block) 
    int local_x = get_local_id(0); 
    int local_y = get_local_id(1); 

    // Compute loop bounds 
    int a_start = A_width * BLOCK_SIZE * block_y; 
    int a_end = a_start + A_width - 1; 
    int b_start = BLOCK_SIZE * block_x; 

    float running_sum = 0.0f; 
    for (int a = a_start, b = b_start; a <= a_end; a += BLOCK_SIZE, b += (BLOCK_SIZE * B_width)) 
    { 
     A_local[local_y][local_x] = A[a + A_width * local_y + local_x]; 
     B_local[local_x][local_y] = B[b + B_width * local_y + local_x]; 
     #pragma unroll 
     for (int k = 0; k < BLOCK_SIZE; ++k) 
     { 
      running_sum += A_local[local_y][k] * B_local[local_x][k]; 
     } 
    } 

    // Store result in matrix C 
    C[get_global_id(1) * get_global_size(0) + get_global_id(0)] = running_sum; 
}

假设块大小为2，则：block_x和block_y均为0;并且local_x和local_y都是0.
然后A_local[0][0]将是A[0]和B_local[0][0]将是B[0]。
尺寸A_local和B_local其中每个插件4个元素。

在这种情况下，A_local和B_local如何在该迭代中访问块中的其他元素？
也将单独的线程/核心分配给每个local_x和local_y？

来源

2015-12-16 user3509540

A_local和B_local都由工作组的所有工作项目共享，因此，所有它们的元素被加载在平行（由工作组中的所有工作项目）在包围for循环的每个步骤。

然后每个工作项目使用一些加载的值（不一定是工作项自己加载的值）来执行其计算的共享。

最后，工作项目将其单个结果存储到全局输出矩阵中。

这是矩阵 - 矩阵乘法的经典平铺实现。不过，我真的很惊讶，看不到任何种类的内存同步函数调用，如work_group_barrier(CLK_LOCAL_MEM_FENCE)之间的加载A_local和B_local和他们在k循环中使用...但我可能很好地忽略了一些东西在这里。

来源

2015-12-17 09:06:00 Gilles

你的代码示例中肯定存在一个障碍。如果所有工作项以锁步方式执行指令，则外部for循环只会产生正确的结果，从而保证在for循环之前填充本地内存。

也许这是Altera和其他FPGA的情况，但这对于CPU和GPU是不正确的。

您应该添加障碍（CLK_LOCAL_MEM_FENCE）;如果您收到意想不到的结果，或想与其他类型的硬件兼容。

float running_sum = 0.0f; 
for (int a = a_start, b = b_start; a <= a_end; a += BLOCK_SIZE, b += (BLOCK_SIZE * B_width)) 
{ 
    A_local[local_y][local_x] = A[a + A_width * local_y + local_x]; 
    B_local[local_x][local_y] = B[b + B_width * local_y + local_x]; 

    barrier(CLK_LOCAL_MEM_FENCE); 

    #pragma unroll 
    for (int k = 0; k < BLOCK_SIZE; ++k) 
    { 
     running_sum += A_local[local_y][k] * B_local[local_x][k]; 
    } 
}

来源

2015-12-17 13:48:19 mfa

OpenCL Matrix乘法Altera示例

回答

相关问题