One can specify work group sizes with kernel attribute reqd_work_group_size and/or with the local_work_size argument of clEnqueueNDRangeKernel(). If reqd_work_group_size is specified, then local_work_size must match it. I wonder what the rationale is behind it. Why can not local_work_size be NULL?