Uncoalesced float2 CUDA kernel

I am having trouble optimizing the grid and block sizes of the example below. When I do profiling, it appears that the memory write operation in the kernel code is not coalesced.

I found some solutions on the internet but they suggested me to change the structure of c_image to [x1, x2, x3...] [y1, y2, y3...]

However I need the structure as [x1, y1] [x2, y2]... since I am using CUFFT library on the code somewhere else which requires this exact form.

Is there a coalesced way to perform operations in c_image structure?

My code:

void main()
{
    float2 *c_image;  // x1 y1 x2 y2 x3 y3 x4 y4 .. .. .. .. x2048 y2048
    cudamalloc(c_image, 2048*2048*8);

    //warp size = 32
    //max thread count = 1024
    dim3 blocksize(1024, 1);
    dim3 gridsize(2048, 2);
    test<<<gridsize, blocksize>>(C_image);  
}


__global__ void test(float2 *o) 
{
    int x = threadIdx.x + blockIdx.y*1024;
    int y = blockIdx.x;

    int index = x + 2048*y;

        o[index].x = x;
        o[index].y = y;
}

Thanks a lot!

PS: I tried this, but no luck! CUDA float2 coalescing

Solution

Reducing this to a single assignment using a temporary float2 variable should result in a 64-bit write.

_global__ void test(float2 *o) 
{
    int x = threadIdx.x + blockIdx.y * 1024;
    int y = blockIdx.x;
    int index = x + 2048 * y;
    float2 tmp = float2(x, y);
    o[index] = tmp;
}

Additional details can be found at