SHARE
TWEET

louiswu

a guest Feb 7th, 2010 163 Never
  1. __kernel void LinearMetricAddition(__read_only image3d_t target,__read_only image3d_t source,
  2.                                    __global int* _bins,
  3.                                     sampler_t volumeSampler,
  4.                                     __global float* matrix,
  5.                                     )
  6. {
  7.     int i = get_global_id(0); // multiple of 64
  8.     int j = get_global_id(1); // multiple of 64
  9.     int z = get_global_id(2); // multiple of 64
  10.  
  11.     int4 coords = (int4)(i,j,z,1);
  12.  
  13.           // Perform the matrix transformation on the current pixel
  14.                   float _x = matrix[0+0] * i + matrix[0+1] * j + matrix[0+2] * z + matrix[0+3];
  15.                   float _y = matrix[4+0] * i + matrix[4+1] * j + matrix[4+2] * z + matrix[4+3];
  16.                   float _z = matrix[8+0] * i + matrix[8+1] * j + matrix[8+2] * z + matrix[8+3];
  17.  
  18.             float4 transCoords = (float4)(_x,_y,_z,1.0f);
  19.             int res = read_imagei(source,volumeSampler,transCoords).x;
  20.            
  21.             //Block size is 64,1,1 ensuring that all operations to _bins will be coalesced
  22.             _bins[i + j*64 + z*64*64] = read_imagei(target,volumeSampler,coords).x;
  23. }
RAW Paste Data
Top