2
Я пытаюсь реализовать фильтр FIR (конечный импульсный отклик) в CUDA. Мой подход достаточно прост и выглядит примерно так:FIR-фильтр в CUDA (как 1D свертка)
#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, мой код занимает примерно в три раза больше времени для выполнения.
Есть ли у кого-нибудь идеи о том, как ускорить процесс?
EDIT: Сменили все filterLength на numeratorLength.
ли 'numeratorLength' так же, как' filterLength'? Я не вижу определения 'numatorLength' где-либо в том, что вы разместили. Эта проблема, по сути, является проблемой одномерного трафарета. Стандартная оптимизация для проблем с трафаретом заключается в том, чтобы довести часть входных данных в разделяемую память, достаточную для потоковых потоков блоков для вычисления их выходов, а затем позволить этим потокам работать из копии общей памяти. –
Если вы в конечном итоге избиваете ArrayFire, дайте нам знать! Если нет, вы всегда можете просто использовать ArrayFire, так как это быстрее :) – arrayfire
@RobertCrovella Да, numatorLength такая же, как filterLength. Я решил изменить имя, но, видимо, пропустил несколько мест. Мой плохой, извините. Я изменил исходный пост так, что есть только numatorLength. Спасибо за предложение использовать общую память. Я читал, что они значительно быстрее, чем глобальная память, но я немного уверен в том, как наилучшим образом реализовать это, поскольку общая память ограничена по размеру, а длины фильтров могут быть довольно длинными. Я буду играть с ним и посмотреть, как это происходит. – Elfendahl