workarounds/fixes for VIV OpenCL

- query allowed workgroup size for each kernel

- use 2D global workgroup to work around 1<<16 dimension limit
  - possibly better alternative is to use clEnqueueNDRangeKernel in a loop with an offset
  - but how to query this limit?!