filter - CUDA 中的 FIR 滤波器(作为一维卷积)

Question

我正在尝试在 CUDA 中实现 FIR(有限脉冲响应)滤波器。我的方法很简单，看起来有点像这样:

#include <cuda.h>

__global__ void filterData(const float *d_data,
                           const float *d_numerator, 
                           float *d_filteredData, 
                           const int numeratorLength,
                           const int filteredDataLength)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;

    float sum = 0.0f;

    if (i < filteredDataLength)
    {
        for (int j = 0; j < numeratorLength; j++)
        {
            // The first (numeratorLength-1) elements contain the filter state
            sum += d_numerator[j] * d_data[i + numeratorLength - j - 1];
        }
    }

    d_filteredData[i] = sum;
}

int main(void)
{
    // (Skipping error checks to make code more readable)

    int dataLength = 18042;
    int filteredDataLength = 16384;
    int numeratorLength= 1659;

    // Pointers to data, filtered data and filter coefficients
    // (Skipping how these are read into the arrays)
    float *h_data = new float[dataLength];
    float *h_filteredData = new float[filteredDataLength];
    float *h_filter = new float[numeratorLength];


    // Create device pointers
    float *d_data = nullptr;
    cudaMalloc((void **)&d_data, dataLength * sizeof(float));

    float *d_numerator = nullptr;
    cudaMalloc((void **)&d_numerator, numeratorLength * sizeof(float));

    float *d_filteredData = nullptr;
    cudaMalloc((void **)&d_filteredData, filteredDataLength * sizeof(float));


    // Copy data to device
    cudaMemcpy(d_data, h_data, dataLength * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_numerator, h_numerator, numeratorLength * sizeof(float), cudaMemcpyHostToDevice);  

    // Launch the kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (filteredDataLength + threadsPerBlock - 1) / threadsPerBlock;
    filterData<<<blocksPerGrid,threadsPerBlock>>>(d_data, d_numerator, d_filteredData, numeratorLength, filteredDataLength);

    // Copy results to host
    cudaMemcpy(h_filteredData, d_filteredData, filteredDataLength * sizeof(float), cudaMemcpyDeviceToHost);

    // Clean up
    cudaFree(d_data);
    cudaFree(d_numerator);
    cudaFree(d_filteredData);

    // Do stuff with h_filteredData...

    // Clean up some more
    delete [] h_data;
    delete [] h_filteredData;
    delete [] h_filter;
}

过滤器有效，但由于我是 CUDA 编程的新手，所以我不确定如何优化它。

我看到的一个小问题是，在我打算使用的应用程序中，dataLength、filteredDataLength 和 numeratorLength 是事先不知道的此外，即使 dataLength 在上述代码中是 32 的倍数，但不能保证在最终应用程序中就是这样。

当我将上面的代码与 ArrayFire 进行比较时，我的代码执行时间大约是原来的三倍。

有没有人对如何加快速度有任何想法？

编辑:已将所有 filterLength 更改为 numeratorLength。

Answer 1

我可以提出以下建议来加速您的代码:

1. 使用共享内存:它是一个很小的类似缓存的内存，但非常比全局卡内存更快。你可以找到更多关于它的信息在 CUDA 文档中寻找 __shared__ 关键字。为了例如，您可以预取过滤器分子和大块共享内存中的数据，这将显着提高您的表现。您需要格外注意数据在这种情况下对齐，因为它真的很重要，它可能会减慢速度你的代码。
2. 考虑展开分子的 for 循环和。您可以查看 CUDA 中的 reduce-vector 示例文档。
3. 您还可以考虑将分子循环本身。这可以通过向线程 block 添加一个额外的维度(比如“y”)来完成。您还需要使 sum 成为具有 numeratorLength 维度的共享向量。您还可以查看 reduce vector 示例，了解如何在最后快速取这个向量的总和。