Working with texture


Press the button to start the texture read/write test.

To write a 2×2 texture (image) enqueueWriteImage function is used

  var format = {channelOrder:WebCL.CL_R, channelDataType:WebCL.CL_FLOAT};
  var imgIn  = ctx.createImage2D(WebCL.CL_MEM_READ_ONLY, format, 2, 2, 0);
  ...
  cmdQueue.enqueueWriteImage(imgIn, true, [0,0,0], [2,2,1], 0, 0, data, []);

Textures vs. buffer arrays

In difference schemes for differencial equations every point is used many times by the "nearest" kernels (see below pressure calculations for 2D droplets)
constant int n = 512, m = 511, sh = 9;
__kernel void kP(global float* P, global float* G, global float* P1 ){
  int x = get_global_id(0),  y = get_global_id(1),  ny = y<<sh;
  P1[x + ny] = (P[((x + m)&m) + ny] + P[((x + 1)&m) + ny] +
    P[x + (((y + 1)&m)<<sh)] + P[x + (((y + m)&m)<<sh)] - G[x + ny]) *.25f;
}
On new GPU there are L1 and L2 caches common for textures and buffers. In any case I don't see difference between them on AMD E-350.

Global memory bandwidth profiling

Global memory bandwidth tests: copying 1024×1024 float4 matrices and 2048×2048 float matrices.

Operation   float write   float copy   float4 copy   (GB/s)
E-350 3.3 5.5 5.9
HD 4870 3.6   7.2   58
GTX 560   33.8   104.6

For GTX 560 the maximum bandwidth is 128 and for E-350 - 9 GB/s.


WebCL Demos     updated 4 Nov 2011