Just notifying you of a potential compiler conformance issue. When compiling pointer arithmetic such as:
__global *p = q + mad(a, b, c);
with a, b and c all uint, I get the error "Expression must have integral type". This is cleared up by casting the mad result to an int. AFAIK I don't believe there is anything in the C99 or OpenCL standard stopping uint from being used in pointer arithmetic...
Hmm, it seems that mad(a,b,c) is not defined in the spec for int or uint. So the issue here is one of conversions, where int is being converted to float to execute the mad operation...
My mistake, Doh!
While I am at it, AMD OpenCL compiler allows:
int lid = get_local_id(0);
Where (if I am interpreting this right), section 6.2.6 OpenCL spec (rule 4) indicates that uint is higher rank than int, meaning this conversion shouldn't be allowed implicitly.