OpenCL双三次插值内核失败，错误CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST
|
双立方插值是常见的插值方法之一，但我在OpenCL上找不到任何有效的实现方法。我决定自己在OpenCL上写双三次插值，但......
我的内核程序有些问题。
当我运行内核执行时，程序失败，错误为CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST。没有关于错误原因的任何其他信息。我正在使用javacl绑定表格google代码：http：//code.google.com/p/javacl,AMED加速并行处理SDK 2.3在Ubuntu linux 10.10上，硬件AMD Radeon 5xxxHD
我在ubuntu上没有用于AMD APP SDK的opencl调试器（
如果我取消注释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);
}
}
没有找到相关结果
已邀请:
1 个回复
了驳
kernelBicubic.getKernel().setArgs(scaleFactor, inImage, imageOut); lastEvent=kernelBicubic.getKernel().enqueueNDRange(queue, new int[]{(int) inImage.getWidth()+1,(int) inImage.getHeight()+1},lastEvent);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}; inlie float calcT(image2d_t signal,int x,int y,int i,int j){ return read_imagef(signal, sampler, (int2)(x ,y+i)).x * CATMULL_ROM[4*j] +read_imagef(signal, sampler, (int2)(x+1,y+i)).x * CATMULL_ROM[4*j+1] +read_imagef(signal, sampler, (int2)(x+2,y+i)).x * CATMULL_ROM[4*j+2] +read_imagef(signal, sampler, (int2)(x+3,y+i)).x * CATMULL_ROM[4*j+3]; } inline float C(image2d_t signal,int x,int y,int i,int j){ return CATMULL_ROM[4*i ] * calcT(signal,x,y,0,j) +CATMULL_ROM[4*i+1] * calcT(signal,x,y,1,j) +CATMULL_ROM[4*i+2] * calcT(signal,x,y,2,j) +CATMULL_ROM[4*i+3] * calcT(signal,x,y,3,j); } __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 C0 =C(signal,x,y,0,0); float C1 =C(signal,x,y,0,1); float C2 =C(signal,x,y,0,2); float C3 =C(signal,x,y,0,3); float C4 =C(signal,x,y,1,0); float C5 =C(signal,x,y,1,1); float C6 =C(signal,x,y,1,2); float C7 =C(signal,x,y,1,3); float C8 =C(signal,x,y,2,0); float C9 =C(signal,x,y,2,1); float C10=C(signal,x,y,2,2); float C11=C(signal,x,y,2,3); float C12=C(signal,x,y,3,0); float C13=C(signal,x,y,3,1); float C14=C(signal,x,y,3,2); float C15=C(signal,x,y,3,3); float xOff=scale*1.5F + x*scale; float yOff=scale*1.5F + y*scale; for (int i = 0; i < scale; i++) { for (int j = 0; j < scale; j++) { float iY=(float)j/(float) scale; float iX=(float)i/(float) scale; float val= iX * (iX * (iX * (iY * (iY * (iY * C0 + C1) + C2) + C3) + (iY * (iY * (iY * C4 + C5) + C6) + C7)) + (iY * (iY * (iY * C8 + C9) + C10) + C11)) + (iY * (iY * (iY * C12 + C13) + C14) + C15); write_imagef(upscale, (int2)(xOff+j, yOff+i), val); } } }