OpenCL 矩阵乘法 Altera 示例答案

【问题标题】：OpenCL Matrix Multiplication Altera ExampleOpenCL 矩阵乘法 Altera 示例
【发布时间】：2015-12-17 21:48:19
【问题描述】：

我是 OpenCL 的新手，正在浏览 Altera OpenCL 示例。在他们的矩阵乘法示例中，他们使用了块的概念，其中输入矩阵的维度是块大小的倍数。代码如下：

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 分配单独的线程/内核吗？

【问题讨论】：

标签： opencl matrix-multiplication intel-fpga

【解决方案1】：

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

Altera 和其他 FPGA 可能就是这种情况，但对于 CPU 和 GPU 来说却不是这样。

如果您得到意想不到的结果，或者想要与其他类型的硬件兼容，您应该添加 barrier(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];
    }
}

【讨论】：

【解决方案2】：

A_local 和B_local 都由工作组的所有工作项共享，因此它们的所有元素在每一步都并行（由工作组的所有工作项）加载包含for 循环。

然后每个工作项使用一些加载的值（不一定是工作项本身加载的值）来完成它的计算份额。

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

它是矩阵-矩阵乘法的经典平铺实现。但是，我真的很惊讶没有看到任何形式的内存同步函数调用，例如work_group_barrier(CLK_LOCAL_MEM_FENCE) 在A_local 和B_local 的负载之间以及它们在k 循环中的使用......但我很可能在这里忽略了一些东西。

【讨论】：