I’m working to improve the perfomance of a 2D FFT on images . After a benchmark ; I regconize the transpose function is the reason make the program slow , so I replace it with a more optimized one .
But after that ; I received a return code of all the functions which work fine before
CL_INVALID_KERNEL_NAME
. Except the transpose function and
clSetKernelArg
in the host code ; I don’t change anything else . So I’m out of idea. Hope you guys help me out
Here is the kernel file
The transpose function before :
__kernel void transpose(__global float2 *dst, __global float2* src, int n)
{
unsigned int xgid = get_global_id(0);
unsigned int ygid = get_global_id(1);
unsigned int iid = ygid * n + xgid;
unsigned int oid = xgid * n + ygid;
dst[oid] = src[iid];
}
and the new one :
#define BLOCK_DIM 16
__kernel void transpose(
__global float2* dataout,
__global float2* datain,
int width, int height)
// width = N (signal length)
// height = batch_size (number of signals in a batch)
{
// read the matrix tile into shared memory
__local float2 block[BLOCK_DIM * (BLOCK_DIM + 1)]
unsigned int xIndex = get_global_id(0);
unsigned int yIndex = get_global_id(1);
if((xIndex < width) && (yIndex < height))
{
unsigned int index_in = yIndex * width + xIndex;
int Idin = get_local_id(1)*(BLOCK_DIM+1)+get_local_id(0);
block[Idin]= datain[index_in];
}
barrier(CLK_LOCAL_MEM_FENCE);
// write the transposed matrix tile to global memory
xIndex = get_group_id(1) * BLOCK_DIM + get_local_id(0);
yIndex = get_group_id(0) * BLOCK_DIM + get_local_id(1);
if((xIndex < height) && (yIndex < width))
{
unsigned int index_out = yIndex * height + xIndex;
int Idout = get_local_id(0)*(BLOCK_DIM+1)+get_local_id(1);
dataout[index_out] = block[Idout];
}
}