Writing to a buffer is slow in opencl

Question

I have the following buffer:

    int nbytes = 256*256*4;
    uint8_t buffer_window[nbytes;

I allocate it on host like above. Now I'm using the following for creating it

  a_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, nbytes, (void*)&buffer_window[0], &_err));

and I enqeue a 2D Kernel normally on 2D framebuffer:

  global_size[0] = 256;
  global_size[1] = 256;
  

  auto time_start = std::chrono::high_resolution_clock::now();
  CL_CHECK(clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL, global_size, NULL, 0, NULL, NULL));
  CL_CHECK(clFinish(commandQueue));

In the kernel whenever I write data to that buffer, it executes very slow, but when I remove any writing to that buffer, the kernel executes very fast.

Follow up Kernel

__kernel void sendImageToPBO(__global uchar4* dst_buffer, __global struct Triangle_* triangles, int triCount)

{
    
size_t px = get_global_id(0); // triCount
size_t py = get_global_id(1); // triCount
int width = 800;
int height = 800;

float3 v0Raster = (float3)(triangles[px].v[0].pos[0], triangles[px].v[0].pos[1], triangles[px].v[0].pos[2]);
float3 v1Raster = (float3)(triangles[px].v[1].pos[0], triangles[px].v[1].pos[1], triangles[px].v[1].pos[2]);
float3 v2Raster = (float3)(triangles[px].v[2].pos[0], triangles[px].v[2].pos[1], triangles[px].v[2].pos[2]);
float xmin = min3(v0Raster.x, v1Raster.x, v2Raster.x);
float ymin = min3(v0Raster.y, v1Raster.y, v2Raster.y);
float xmax = max3(v0Raster.x, v1Raster.x, v2Raster.x);
float ymax = max3(v0Raster.y, v1Raster.y, v2Raster.y);
float slope = (ymax - ymin) / (xmax - xmin);
int dp, y;
bool discard_;
float ratio;
for (int x = round(xmin); x <= round(xmax); x++) {
    y = slope * (x - round(xmin) + ymin);
    ratio = (x - round(xmin) / (round(xmax) - round(xmin)));
    discard_ = false;

    int flatIdx = width - x + (height - y) * width;
    if (y < 0 || y > height || x < 0 || x > width) {
        discard_ = true;
    }
    if (!discard_) {


        fragments[flatIdx].col[0] = 1.0f;
        fragments[flatIdx].col[1] = 0;
        fragments[flatIdx].col[2] = 0;

    }
}

ProjectPhysX · Answer

If the kernel does not write any data/results to a buffer in global memory, the compiler throws out all the code and you essentially get an empty kernel with zero execution time.

Global range / buffer size in your case is very small. Typically, you want at least several million threads to get full saturation and good performance. Otherwise, initial kernel compile time, PCIe data transfer and and kernel call latency might dominate, rather than kernel execution time.

Writing to a buffer is slow in opencl

1 Answers