Corrected the syntax for the reqd_work_group_size attribute

platforms/opencl/src/OpenCLContext.cpp

@@ -70,7 +70,7 @@ OpenCLContext::OpenCLContext(int numParticles, int deviceIndex) : time(0.0), ste
         this->deviceIndex = deviceIndex;
         if (device.getInfo<CL_DEVICE_MAX_WORK_GROUP_SIZE>() < minThreadBlockSize)
             throw OpenMMException("The specified OpenCL device is not compatible with OpenMM");
-        compilationOptions = "-cl-fast-relaxed-math";
+        compilationOptions = "-cl-fast-relaxed-math -DWORK_GROUP_SIZE="+OpenCLExpressionUtilities::intToString(ThreadBlockSize);
         string vendor = device.getInfo<CL_DEVICE_VENDOR>();
         if (vendor.size() >= 6 && vendor.substr(0, 6) == "NVIDIA") {
             compilationOptions += " -DWARPS_ARE_ATOMIC";

platforms/opencl/src/kernels/customGBEnergyN2_default.cl

@@ -6,8 +6,8 @@
  * 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,
+__kernel __attribute__((reqd_work_group_size(WORK_GROUP_SIZE, 1, 1)))
+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,
 #ifdef USE_CUTOFF

platforms/opencl/src/kernels/customGBEnergyN2_nvidia.cl

@@ -6,8 +6,8 @@
  * 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,
+__kernel __attribute__((reqd_work_group_size(WORK_GROUP_SIZE, 1, 1)))
+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,
 #ifdef USE_CUTOFF

platforms/opencl/src/kernels/customGBValueN2_default.cl

@@ -4,8 +4,8 @@
  * 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,
+__kernel __attribute__((reqd_work_group_size(WORK_GROUP_SIZE, 1, 1)))
+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,
 #ifdef USE_CUTOFF

platforms/opencl/src/kernels/customGBValueN2_nvidia.cl

@@ -4,8 +4,8 @@
  * 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,
+__kernel __attribute__((reqd_work_group_size(WORK_GROUP_SIZE, 1, 1)))
+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,
 #ifdef USE_CUTOFF

platforms/opencl/src/kernels/gbsaObc_default.cl

@@ -14,8 +14,8 @@ 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,
+__kernel __attribute__((reqd_work_group_size(WORK_GROUP_SIZE, 1, 1)))
+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) {
 #else
@@ -193,8 +193,8 @@ __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,
+__kernel __attribute__((reqd_work_group_size(WORK_GROUP_SIZE, 1, 1)))
+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,
 #ifdef USE_CUTOFF

platforms/opencl/src/kernels/gbsaObc_nvidia.cl

@@ -14,8 +14,8 @@ 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,
+__kernel __attribute__((reqd_work_group_size(WORK_GROUP_SIZE, 1, 1)))
+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) {
 #else
@@ -253,8 +253,8 @@ __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,
+__kernel __attribute__((reqd_work_group_size(WORK_GROUP_SIZE, 1, 1)))
+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,
 #ifdef USE_CUTOFF

platforms/opencl/src/kernels/nonbonded_default.cl

@@ -11,8 +11,8 @@ 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,
+__kernel __attribute__((reqd_work_group_size(WORK_GROUP_SIZE, 1, 1)))
+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
         __global unsigned int* interactionFlags, __global unsigned int* interactionCount

platforms/opencl/src/kernels/nonbonded_nvidia.cl

@@ -11,8 +11,8 @@ 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,
+__kernel __attribute__((reqd_work_group_size(WORK_GROUP_SIZE, 1, 1)))
+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
         __global unsigned int* interactionFlags, __global unsigned int* interactionCount