OpenCL 1D 跨步卷积性能

Question

为了对信号进行下采样，我使用了 FIR 滤波器 + 抽取阶段（这是一个实用的跨步卷积）。结合过滤和抽取的最大优势是降低了计算成本（通过抽取因子）。

使用直接的 OpenCL 实现，我无法从抽取中受益。恰恰相反：抽取因子为4的卷积比全卷积慢25%

内核代码：

__kernel void decimation(__constant float *input,
                         __global   float *output,
                         __constant float *coefs,
                         const int taps,
                         const int decimationFactor) {

    int posOutput = get_global_id(0);
    float result = 0;

    for (int tap=0; tap<taps; tap++) {
        int posInput = (posOutput * decimationFactor) - tap;
        result += input[posInput] * coefs[tap];
    }

    output[posOutput] = result;
}

我猜这是由于未合并的内存访问。虽然我想不出解决问题的方法。有什么想法吗？

编辑： 我尝试了 Dithermaster 的解决方案，将问题拆分为合并读取到共享本地内存和从本地内存进行卷积：

__kernel void decimation(__constant float *input,
                        __global   float *output,
                        __constant float *coefs,
                        const int taps,
                        const int decimationFactor,
                        const int bufferSize,
                        __local float *localInput) {

    const int posOutput = get_global_id(0);
    const int localSize = get_local_size(0);
    const int localId   = get_local_id(0);
    const int groupId   = get_group_id(0);

    const int localInputOffset  = taps-1;
    const int localInputOverlap = taps-decimationFactor;
    const int localInputSize    = localInputOffset + localSize * decimationFactor;

    // 1. transfer global input data to local memory
    // read global input to local input (only overlap)
    if (localId < localInputOverlap) {
        int posInputStart = ((groupId*localSize) * decimationFactor) - (taps-1);
        int posInput      = posInputStart + localId;
        int posLocalInput = localId;

        localInput[posLocalInput] = 0.0f;
        if (posInput >= 0)
            localInput[posLocalInput] = input[posInput];
    }

    // read remaining global input to local input
    // 1. alternative: strided read
    // for (int i=0; i<decimationFactor; i++) {
    //     int posInputStart = (groupId*localSize) * decimationFactor;
    //     int posInput      = posInputStart    + localId * decimationFactor - i;
    //     int posLocalInput = localInputOffset + localId * decimationFactor - i;

    //     localInput[posLocalInput] = 0.0f;
    //     if ((posInput >= 0) && (posInput < bufferSize*decimationFactor))
    //         localInput[posLocalInput] = input[posInput];
    // }

    // 2. alternative: coalesced read (in blocks of localSize)
    for (int i=0; i<decimationFactor; i++) {
        int posInputStart = (groupId*localSize) * decimationFactor;
        int posInput      = posInputStart    - (decimationFactor-1) + i*localSize + localId;
        int posLocalInput = localInputOffset - (decimationFactor-1) + i*localSize + localId;

        localInput[posLocalInput] = 0.0f;
        if ((posInput >= 0) && (posInput < bufferSize*decimationFactor))
            localInput[posLocalInput] = input[posInput];
    }

    // 2. wait until every thread completed
    barrier(CLK_LOCAL_MEM_FENCE);

    // 3. convolution
    if (posOutput < bufferSize) {
        float result = 0.0f;
        for (int tap=0; tap<taps; tap++) {
            int posLocalInput = localInputOffset + (localId * decimationFactor) - tap;

            result += localInput[posLocalInput] * coefs[tap];
        }

        output[posOutput] = result;
    }
}

大进步！但是，性能仍然与整体操作无关（与抽取因子不成比例）：

与第一种方法相比，全卷积加速：~12%
与全卷积相比，抽取的计算时间：
- 抽取因子 2：61 %
- 抽取因子 4：46 %
- 抽取因子 8：53 %
- 抽取因子 16：68 %

抽取系数为 4 时性能最佳。这是为什么呢？有进一步改进的想法吗？

编辑 2： 具有共享本地内存的图表：

编辑 3： 3 种不同实现的性能比较

Answer 1

由于数据重叠量 (66%)，这可能会受益于在工作组内的工作项之间共享从内存中读取的数据。您可以摆脱冗余读取并也进行合并读取。将内核分为两部分：第一部分将工作组内所需的所有数据合并读取到共享本地内存中。然后内存屏障同步。然后在第二部分使用从共享本地内存读取的内容进行卷积。

P.S。感谢您提供图表，它帮助我比尝试阅读代码更快地理解您的目标。

OpenCL 1D 跨步卷积性能

OpenCL 1D strided convolution performance

filtering

gpu

signal-processing

downsampling

opencl