Uncoalesced float2 CUDA kernel

StackOverflow https://stackoverflow.com/questions/13408968

  •  29-11-2021
  •  | 
  •  

Question

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

Was it helpful?

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

Licensed under: CC-BY-SA with attribution
Not affiliated with StackOverflow
scroll top