Skip to content

GitLab

Commit 10b51d25 authored by peastman's avatar peastman Committed by GitHub
Browse files

Merge pull request #1579 from peastman/cleanup 

Minor code cleanup
parents 9e34d39e 43c8e814
Hide whitespace changes
Inline Side-by-side
Showing with 22 additions and 27 deletions
platforms/cuda/src/kernels/customGBEnergyN2.cu
View file @ 10b51d25
......@@ -236,7 +236,7 @@ extern "C" __global__ void computeN2Energy(unsigned long long* __restrict__ forc
LOAD_ATOM1_PARAMETERS
const unsigned int localAtomIndex = threadIdx.x;
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......
platforms/cuda/src/kernels/customGBValueN2.cu
View file @ 10b51d25
......@@ -212,7 +212,7 @@ extern "C" __global__ void computeN2Value(const real4* __restrict__ posq, const
LOAD_ATOM1_PARAMETERS
const unsigned int localAtomIndex = threadIdx.x;
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......
platforms/cuda/src/kernels/gbsaObc1.cu
View file @ 10b51d25
......@@ -264,7 +264,7 @@ extern "C" __global__ void computeBornSum(unsigned long long* __restrict__ globa
real4 posq1 = posq[atom1];
float2 params1 = global_params[atom1];
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......@@ -619,7 +619,7 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo
real4 posq1 = posq[atom1];
real bornRadius1 = global_bornRadii[atom1];
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......
platforms/cuda/src/kernels/nonbonded.cu
View file @ 10b51d25
......@@ -379,7 +379,7 @@ extern "C" __global__ void computeNonbonded(
LOAD_ATOM1_PARAMETERS
//const unsigned int localAtomIndex = threadIdx.x;
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......
platforms/cuda/src/kernels/sort.cu
View file @ 10b51d25
......@@ -98,7 +98,7 @@ __global__ void computeRange(const DATA_TYPE* __restrict__ data, unsigned int le
* Assign elements to buckets.
*/
__global__ void assignElementsToBuckets(const DATA_TYPE* __restrict__ data, unsigned int length, unsigned int numBuckets, const KEY_TYPE* __restrict__ range,
unsigned int* bucketOffset, unsigned int* __restrict__ bucketOfElement, unsigned int* __restrict__ offsetInBucket) {
unsigned int* __restrict__ bucketOffset, unsigned int* __restrict__ bucketOfElement, unsigned int* __restrict__ offsetInBucket) {
float minValue = (float) (range[0]);
float maxValue = (float) (range[1]);
float bucketWidth = (maxValue-minValue)/numBuckets;
......
platforms/opencl/src/kernels/customGBEnergyN2.cl
View file @ 10b51d25
......@@ -252,7 +252,7 @@ __kernel void computeN2Energy(
LOAD_ATOM1_PARAMETERS
const unsigned int localAtomIndex = get_local_id(0);
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......
platforms/opencl/src/kernels/customGBEnergyN2_cpu.cl
View file @ 10b51d25
......@@ -255,7 +255,7 @@ __kernel void computeN2Energy(
for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) {
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+localAtomIndex] : y*TILE_SIZE+localAtomIndex);
unsigned int j = interactingAtoms[pos*TILE_SIZE+localAtomIndex];
#else
unsigned int j = y*TILE_SIZE+localAtomIndex;
#endif
......
platforms/opencl/src/kernels/customGBValueN2.cl
View file @ 10b51d25
......@@ -229,7 +229,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
LOAD_ATOM1_PARAMETERS
const unsigned int localAtomIndex = get_local_id(0);
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......
platforms/opencl/src/kernels/customGBValueN2_cpu.cl
View file @ 10b51d25
......@@ -228,7 +228,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) {
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+localAtomIndex] : y*TILE_SIZE+localAtomIndex);
unsigned int j = interactingAtoms[pos*TILE_SIZE+localAtomIndex];
#else
unsigned int j = y*TILE_SIZE+localAtomIndex;
#endif
......
platforms/opencl/src/kernels/gbsaObc.cl
View file @ 10b51d25
......@@ -232,7 +232,7 @@ __kernel void computeBornSum(
real4 posq1 = posq[atom1];
float2 params1 = global_params[atom1];
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......@@ -619,7 +619,7 @@ __kernel void computeGBSAForce1(
real4 posq1 = posq[atom1];
real bornRadius1 = global_bornRadii[atom1];
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......
platforms/opencl/src/kernels/gbsaObc_cpu.cl
View file @ 10b51d25
......@@ -228,7 +228,7 @@ __kernel void computeBornSum(
for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) {
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+localAtomIndex] : y*TILE_SIZE+localAtomIndex);
unsigned int j = interactingAtoms[pos*TILE_SIZE+localAtomIndex];
#else
unsigned int j = y*TILE_SIZE+localAtomIndex;
#endif
......@@ -641,7 +641,7 @@ __kernel void computeGBSAForce1(
for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) {
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+localAtomIndex] : y*TILE_SIZE+localAtomIndex);
unsigned int j = interactingAtoms[pos*TILE_SIZE+localAtomIndex];
#else
unsigned int j = y*TILE_SIZE+localAtomIndex;
#endif
......
platforms/opencl/src/kernels/nonbonded.cl
View file @ 10b51d25
......@@ -269,7 +269,7 @@ __kernel void computeNonbonded(
LOAD_ATOM1_PARAMETERS
const unsigned int localAtomIndex = get_local_id(0);
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......
platforms/opencl/src/kernels/nonbonded_cpu.cl
View file @ 10b51d25
......@@ -269,7 +269,7 @@ __kernel void computeNonbonded(
for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) {
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+localAtomIndex] : y*TILE_SIZE+localAtomIndex);
unsigned int j = interactingAtoms[pos*TILE_SIZE+localAtomIndex];
#else
unsigned int j = y*TILE_SIZE+localAtomIndex;
#endif
......
platforms/opencl/src/kernels/sort.cl
View file @ 10b51d25
......@@ -8,7 +8,7 @@ KEY_TYPE getValue(DATA_TYPE value) {
* Sort a list that is short enough to entirely fit in local memory. This is executed as
* a single thread block.
*/
__kernel void sortShortList(__global DATA_TYPE* __restrict__ data, uint length, __local DATA_TYPE* dataBuffer) {
__kernel void sortShortList(__global DATA_TYPE* restrict data, uint length, __local DATA_TYPE* dataBuffer) {
// Load the data into local memory.
for (int index = get_local_id(0); index < length; index += get_local_size(0))
......@@ -96,7 +96,7 @@ __kernel void computeRange(__global const DATA_TYPE* restrict data, uint length,
* Assign elements to buckets.
*/
__kernel void assignElementsToBuckets(__global const DATA_TYPE* restrict data, uint length, uint numBuckets, __global const KEY_TYPE* restrict range,
__global uint* bucketOffset, __global uint* restrict bucketOfElement, __global uint* restrict offsetInBucket) {
__global uint* restrict bucketOffset, __global uint* restrict bucketOfElement, __global uint* restrict offsetInBucket) {
#ifdef AMD_ATOMIC_WORK_AROUND
// Do a byte write to force all memory accesses to interactionCount to use the complete path.
// This avoids the atomic access from causing all word accesses to other buffers from using the slow complete path.
......
plugins/amoeba/platforms/cuda/src/kernels/multipoleElectrostatics.cu
View file @ 10b51d25
......@@ -545,7 +545,7 @@ extern "C" __global__ void computeElectrostatics(
data.force = make_real3(0);
data.torque = make_real3(0);
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......
plugins/amoeba/platforms/cuda/src/kernels/multipoleFixedField.cu
View file @ 10b51d25
......@@ -655,7 +655,7 @@ extern "C" __global__ void computeFixedField(
data.bornRadius = bornRadii[atom1];
#endif
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......
plugins/amoeba/platforms/cuda/src/kernels/multipoleInducedField.cu
View file @ 10b51d25
......@@ -514,7 +514,7 @@ extern "C" __global__ void computeInducedField(
loadAtomData(data, atom1, posq, inducedDipole, inducedDipolePolar, dampingAndThole);
#endif
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......@@ -607,11 +607,6 @@ extern "C" __global__ void recordInducedDipolesForDIIS(const long long* __restri
const real* __restrict__ inducedDipole, const real* __restrict__ inducedDipolePolar, const float* __restrict__ polarizability, float2* __restrict__ errors,
real* __restrict__ prevDipoles, real* __restrict__ prevDipolesPolar, real* __restrict__ prevErrors, int iteration, bool recordPrevErrors, real* __restrict__ matrix) {
extern __shared__ real2 buffer[];
#ifdef USE_EWALD
const real ewaldScale = (4/(real) 3)*(EWALD_ALPHA*EWALD_ALPHA*EWALD_ALPHA)/SQRT_PI;
#else
const real ewaldScale = 0;
#endif
const real fieldScale = 1/(real) 0x100000000;
real sumErrors = 0;
real sumPolarErrors = 0;
......
plugins/amoeba/platforms/cuda/src/kernels/pmeMultipoleElectrostatics.cu
View file @ 10b51d25
......@@ -612,7 +612,7 @@ extern "C" __global__ void computeElectrostatics(
data.force = make_real3(0);
data.torque = make_real3(0);
#ifdef USE_CUTOFF
unsigned int j = (numTiles <= maxTiles ? interactingAtoms[pos*TILE_SIZE+tgx] : y*TILE_SIZE + tgx);
unsigned int j = interactingAtoms[pos*TILE_SIZE+tgx];
#else
unsigned int j = y*TILE_SIZE + tgx;
#endif
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment