static const char* globalAtomicsKernelString= \
"\n"
"\n"
"\n"
"\n"
"//OpenCL 1.1 has atomic_inc build-in (no extension needed)\n"
"//see http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/atomic_inc.html\n"
"__kernel void  globalAtomicKernelOpenCL1_1( volatile __global int* counter)\n"
"{\n"
"	atomic_inc(counter);\n"
"}\n"
"\n"
"//OpenCL 1.1 atomic device counters extension, usually faster on current AMD hardware\n"
"//http://www.khronos.org/registry/cl/extensions/ext/cl_ext_atomic_counters_32.txt\n"
"#pragma OPENCL EXTENSION cl_ext_atomic_counters_32 : enable\n"
"__kernel void  counterAtomicKernelExt( counter32_t counter)\n"
"{\n"
"	atomic_inc(counter);\n"
"}\n"
"\n"
"\n"
"//OpenCL 1.0 optional extension, using atom_inc\n"
"//see http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/cl_khr_global_int32_base_atomics.html\n"
"#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable //atomic_inc\n"
"__kernel void  globalAtomicKernelExt( __global int* counter)\n"
"{\n"
"	atom_inc(counter);\n"
"}\n"
"\n"
"\n"
"__kernel void  globalAtomicKernelCounters32Broken( __global int* counter)\n"
"{\n"
"	(*counter)++;\n"
"}\n"
"\n"
;