Added reqd_work_group_size attribute to some kernels

platforms/opencl/src/kernels/customGBEnergyN2_default.cl

@@ -6,6 +6,7 @@
  * Compute a force based on pair interactions.
  */
 
+__attribute__((reqd_work_group_size(64, 1, 1)))
 __kernel void computeN2Energy(__global float4* forceBuffers, __global float* energyBuffer, __local float4* local_force,
 	__global float4* posq, __local float4* local_posq, __global unsigned int* exclusions, __global unsigned int* exclusionIndices,
         __local float4* tempForceBuffer, __global unsigned int* tiles,

platforms/opencl/src/kernels/customGBEnergyN2_nvidia.cl

@@ -6,6 +6,7 @@
  * Compute a force based on pair interactions.
  */
 
+__attribute__((reqd_work_group_size(64, 1, 1)))
 __kernel void computeN2Energy(__global float4* forceBuffers, __global float* energyBuffer, __local float4* local_force,
 	__global float4* posq, __local float4* local_posq, __global unsigned int* exclusions, __global unsigned int* exclusionIndices,
         __local float4* tempBuffer, __global unsigned int* tiles,

platforms/opencl/src/kernels/customGBValueN2_default.cl

@@ -4,6 +4,7 @@
  * Compute a value based on pair interactions.
  */
 
+__attribute__((reqd_work_group_size(64, 1, 1)))
 __kernel void computeN2Value(__global float4* posq, __local float4* local_posq, __global unsigned int* exclusions,
         __global unsigned int* exclusionIndices, __global float* global_value, __local float* local_value,
         __local float* tempBuffer, __global unsigned int* tiles,

platforms/opencl/src/kernels/customGBValueN2_nvidia.cl

@@ -4,6 +4,7 @@
  * Compute a value based on pair interactions.
  */
 
+__attribute__((reqd_work_group_size(64, 1, 1)))
 __kernel void computeN2Value(__global float4* posq, __local float4* local_posq, __global unsigned int* exclusions,
         __global unsigned int* exclusionIndices, __global float* global_value, __local float* local_value,
         __local float* tempBuffer, __global unsigned int* tiles,

platforms/opencl/src/kernels/gbsaObc_default.cl

@@ -14,6 +14,7 @@ typedef struct {
  * Compute the Born sum.
  */
 
+__attribute__((reqd_work_group_size(64, 1, 1)))
 __kernel void computeBornSum(__global float* global_bornSum, __global float4* posq, __global float2* global_params, __local AtomData* localData, __local float* tempBuffer, __global unsigned int* tiles,
 #ifdef USE_CUTOFF
         __global unsigned int* interactionFlags, __global unsigned int* interactionCount) {
@@ -192,6 +193,7 @@ __kernel void computeBornSum(__global float* global_bornSum, __global float4* po
  * First part of computing the GBSA interaction.
  */
 
+__attribute__((reqd_work_group_size(64, 1, 1)))
 __kernel void computeGBSAForce1(__global float4* forceBuffers, __global float* energyBuffer,
         __global float4* posq, __global float* global_bornRadii,
         __global float* global_bornForce, __local AtomData* localData, __local float4* tempBuffer, __global unsigned int* tiles,

platforms/opencl/src/kernels/gbsaObc_nvidia.cl

@@ -14,6 +14,7 @@ typedef struct {
  * Compute the Born sum.
  */
 
+__attribute__((reqd_work_group_size(64, 1, 1)))
 __kernel void computeBornSum(__global float* global_bornSum, __global float4* posq, __global float2* global_params, __local AtomData* localData, __local float* tempBuffer, __global unsigned int* tiles,
 #ifdef USE_CUTOFF
         __global unsigned int* interactionFlags, __global unsigned int* interactionCount) {
@@ -252,6 +253,7 @@ __kernel void computeBornSum(__global float* global_bornSum, __global float4* po
  * First part of computing the GBSA interaction.
  */
 
+__attribute__((reqd_work_group_size(64, 1, 1)))
 __kernel void computeGBSAForce1(__global float4* forceBuffers, __global float* energyBuffer,
         __global float4* posq, __global float* global_bornRadii,
         __global float* global_bornForce, __local AtomData* localData, __local float4* tempBuffer, __global unsigned int* tiles,

platforms/opencl/src/kernels/nonbonded_default.cl

@@ -11,6 +11,7 @@ typedef struct {
  * Compute nonbonded interactions.
  */
 
+__attribute__((reqd_work_group_size(64, 1, 1)))
 __kernel void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffer, __global float4* posq, __global unsigned int* exclusions,
         __global unsigned int* exclusionIndices, __local AtomData* localData, __local float4* tempBuffer, __global unsigned int* tiles,
 #ifdef USE_CUTOFF

platforms/opencl/src/kernels/nonbonded_nvidia.cl

@@ -11,6 +11,7 @@ typedef struct {
  * Compute nonbonded interactions.
  */
 
+__attribute__((reqd_work_group_size(64, 1, 1)))
 __kernel void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffer, __global float4* posq, __global unsigned int* exclusions,
         __global unsigned int* exclusionIndices, __local AtomData* localData, __local float4* tempBuffer, __global unsigned int* tiles,
 #ifdef USE_CUTOFF