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, []);
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.
| 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 |