Для понижающей дискретизации сигнала я использую КИХ-фильтр + этап прореживания (это практическая ступенчатая свертка).Большим преимуществом сочетания фильтрации и прореживания является снижение вычислительных затрат (на коэффициент прореживания).
С прямой реализацией 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 различных реализаций 