Бикубическая интерполяция является одним из распространенных методов интерполяции, но я не могу найти ни одной работающей реализации на OpenCL. Было решено самому написать бикубическую интерполяцию на OpenCL, но...
У меня проблема с программой ядра.
Когда я запускаю выполнение ядра, программа завершается с ошибкой CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST. Никакой другой информации о причине ошибки нет. Я использую код Google для привязки javacl: http://code.google.com/p/javacl, AMD Accelerated Parallel Processing SDK 2.3 на Ubuntu linux 10.10, аппаратное обеспечение AMD Radeon 5xxxHD
У меня нет отладчика opencl на ubuntu для AMD APP SDK (
Если я раскомментирую float4 val=read_imagef(signal, sampler, (float2)(x+iX,y+iY)); и комментарий вычисления бикубической интерполяции "float4 val=..." все работает без ошибок (но с использованием билинейной интерполяции). Я думаю, что эта ошибка из-за недопустимого доступа к памяти или переполнения памяти регистра.
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_LINEAR | CLK_ADDRESS_CLAMP_TO_EDGE;
const float CATMULL_ROM[16]={-0.5F,1.5F,-1.5F,0.5F,1.0F,-2.5F,2.0F,-0.5F,-0.5F,0.0F,0.5F,0.0F,0.0F,1.0F,0.0F,0.0F};
__kernel void bicubicUpscale(int scale,read_only image2d_t signal, write_only image2d_t upscale) {
int x = get_global_id(0)-2, y = get_global_id(1)-2;
float C[16];
float T[16];
for (int i = 0; i < 16; i++)
{
C[i]=0.0F;
T[i]=0.0F;
}
for (int i = 0; i < 4; i++)
for (int j = 0; j < 4; j++)
for (int k = 0; k < 4; k++)
{
T[4*i+j] += read_imagef(signal, sampler, (int2)(x+k,y+i)).x * CATMULL_ROM[4*j+k];
}
for (int i = 0; i < 4; i++)
for (int j = 0; j < 4; j++)
for (int k = 0; k < 4; k++)
{
C[4*i+j] += CATMULL_ROM[4*i+k] * T[4*k+j];
}
for (int i = 0; i < scale; i++)
{
for (int j = 0; j < scale; j++)
{
float iX=(float)j/(float) scale;
float iY=(float)i/(float) scale;
//float4 val=read_imagef(signal, sampler, (float2)(x+iX,y+iY));
float val= iX * (iX * (iX * (iY * (iY * (iY * C[0] + C[1]) + C[2]) + C[3])
+ (iY * (iY * (iY * C[4] + C[5]) + C[6]) + C[7]))
+ (iY * (iY * (iY * C[8] + C[9]) + C[10]) + C[11]))
+ (iY * (iY * (iY * C[12] + C[13]) + C[14]) + C[15]);
write_imagef(upscale, (int2)(x*scale+j, y*scale+i), val);
}
}
}
Я переписываю эту программу для использования локальной памяти, но она все еще работает неправильно
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_LINEAR | CLK_ADDRESS_CLAMP_TO_EDGE;
const float CATMULL_ROM[]={-0.5F,1.5F,-1.5F,0.5F,1.0F,-2.5F,2.0F,-0.5F,-0.5F,0.0F,0.5F,0.0F,0.0F,1.0F,0.0F,0.0F};
__kernel void bicubicUpscale(local float* sharedBuffer,int scale,read_only image2d_t signal, write_only image2d_t upscale) {
int x = get_global_id(0)-2, y = get_global_id(1)-2;
//int locX=get_local_id(0);
int offsetT = (y+2)*512+(x+2)*32+16;
int offsetC = (y+2)*512+(x+2)*32;
global float* C=&sharedBuffer[offsetT];
global float* T=&sharedBuffer[offsetT];
for (int i = 0; i < 32; i++){
sharedBuffer[offsetC+ i]=0.0F;
}
for (int i = 0; i < 4; i++)
for (int j = 0; j < 4; j++)
for (int k = 0; k < 4; k++){
//T[4*i+j] = mad(read_imagef(signal, sampler, (int2)(x+k,y+i)).x,CATMULL_ROM[4*j+k],T[4*i+j]);
T[i+j] += read_imagef(signal, sampler, (int2)(x+k,y+i)).x * CATMULL_ROM[4*j+k];
}
for (int i = 0; i < 4; i++)
for (int j = 0; j < 4; j++)
for (int k = 0; k < 4; k++){
//C[4*i+j] = mad(CATMULL_ROM[4*i+k],T[4*k+j],C[4*i+j]);
sharedBuffer[offsetC +4*i+j] += CATMULL_ROM[4*i+k] * sharedBuffer[offsetT + 4*k+j];
}
barrier (CLK_GLOBAL_MEM_FENCE);
for (int i = 0; i < scale; i++)
for (int j = 0; j < scale; j++)
{
float iX=(float)j/(float) scale;
float iY=(float)i/(float) scale;
float4 val= iX * (iX * (iX * (iY * (iY * (iY * C[0] + C[1]) + C[2]) + C[3])
+ (iY * (iY * (iY * C[4] + C[5]) + C[6]) + C[7]))
+ (iY * (iY * (iY * C[8] + C[9]) + C[10]) + C[11]))
+ (iY * (iY * (iY * C[12] + C[13]) + C[14]) + C[15]);
write_imagef(upscale, (int2)(x*scale+j, y*scale+i), val);
}
}
Знаете ли вы какое-либо решение этой проблемы.
Исходники Java + сборка maven2. Используйте команду "mvn clean compile exec:java" для компиляции и запуска демо.
С уважением, Игорь