21 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
25 __global ulong* restrict squareOfPixels,
26 __global ulong* restrict sumOfPixels)
33 int i = get_global_id(0);
37 ushort8 pixelShort = convert_ushort8(vload8(i, imagePixels));
39 ushort8 newSquareShort = pixelShort * pixelShort;
46 ulong8 pixelLong = convert_ulong8(pixelShort);
47 ulong8 newSquareLong = convert_ulong8(newSquareShort);
53 ulong4 sumLongPixels1 = pixelLong.hi + pixelLong.lo;
54 ulong2 sumLongPixels2 = sumLongPixels1.hi + sumLongPixels1.lo;
55 ulong sumLongPixels3 = sumLongPixels2.hi + sumLongPixels2.lo;
57 ulong4 sumLongSquares1 = newSquareLong.hi + newSquareLong.lo;
58 ulong2 sumLongSquares2 = sumLongSquares1.hi + sumLongSquares1.lo;
59 ulong sumLongSquares3 = sumLongSquares2.hi + sumLongSquares2.lo;
71 atom_add(sumOfPixels, sumLongPixels3);
72 atom_add(squareOfPixels, sumLongSquares3);