Skip to content

GitLab

Commit 43c8e814 authored by peastman's avatar peastman
Browse files

Minor code cleanup

parent e962bb4e
Show whitespace changes
Inline Side-by-side
Showing with 22 additions and 27 deletions
platforms/cuda/src/kernels/customGBEnergyN2.cu
View file @ 43c8e814
...@@ -234,7 +234,7 @@ extern "C" __global__ void computeN2Energy(unsigned long long* __restrict__ forc ...@@ -234,7 +234,7 @@ extern "C" __global__ void computeN2Energy(unsigned long long* __restrict__ forc
LOAD_ATOM1_PARAMETERS LOAD_ATOM1_PARAMETERS
const unsigned int localAtomIndex = threadIdx.x; const unsigned int localAtomIndex = threadIdx.x;
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
......
platforms/cuda/src/kernels/customGBValueN2.cu
View file @ 43c8e814
...@@ -212,7 +212,7 @@ extern "C" __global__ void computeN2Value(const real4* __restrict__ posq, const ...@@ -212,7 +212,7 @@ extern "C" __global__ void computeN2Value(const real4* __restrict__ posq, const
LOAD_ATOM1_PARAMETERS LOAD_ATOM1_PARAMETERS
const unsigned int localAtomIndex = threadIdx.x; const unsigned int localAtomIndex = threadIdx.x;
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
......
platforms/cuda/src/kernels/gbsaObc1.cu
View file @ 43c8e814
...@@ -264,7 +264,7 @@ extern "C" __global__ void computeBornSum(unsigned long long* __restrict__ globa ...@@ -264,7 +264,7 @@ extern "C" __global__ void computeBornSum(unsigned long long* __restrict__ globa
real4 posq1 = posq[atom1]; real4 posq1 = posq[atom1];
float2 params1 = global_params[atom1]; float2 params1 = global_params[atom1];
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
...@@ -617,7 +617,7 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo ...@@ -617,7 +617,7 @@ extern "C" __global__ void computeGBSAForce1(unsigned long long* __restrict__ fo
real4 posq1 = posq[atom1]; real4 posq1 = posq[atom1];
real bornRadius1 = global_bornRadii[atom1]; real bornRadius1 = global_bornRadii[atom1];
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
......
platforms/cuda/src/kernels/nonbonded.cu
View file @ 43c8e814
...@@ -379,7 +379,7 @@ extern "C" __global__ void computeNonbonded( ...@@ -379,7 +379,7 @@ extern "C" __global__ void computeNonbonded(
LOAD_ATOM1_PARAMETERS LOAD_ATOM1_PARAMETERS
//const unsigned int localAtomIndex = threadIdx.x; //const unsigned int localAtomIndex = threadIdx.x;
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
......
platforms/cuda/src/kernels/sort.cu
View file @ 43c8e814
...@@ -98,7 +98,7 @@ __global__ void computeRange(const DATA_TYPE* __restrict__ data, unsigned int le ...@@ -98,7 +98,7 @@ __global__ void computeRange(const DATA_TYPE* __restrict__ data, unsigned int le
* Assign elements to buckets. * Assign elements to buckets.
*/ */
__global__ void assignElementsToBuckets(const DATA_TYPE* __restrict__ data, unsigned int length, unsigned int numBuckets, const KEY_TYPE* __restrict__ range, __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 minValue = (float) (range[0]);
float maxValue = (float) (range[1]); float maxValue = (float) (range[1]);
float bucketWidth = (maxValue-minValue)/numBuckets; float bucketWidth = (maxValue-minValue)/numBuckets;
......
platforms/opencl/src/kernels/customGBEnergyN2.cl
View file @ 43c8e814
...@@ -250,7 +250,7 @@ __kernel void computeN2Energy( ...@@ -250,7 +250,7 @@ __kernel void computeN2Energy(
LOAD_ATOM1_PARAMETERS LOAD_ATOM1_PARAMETERS
const unsigned int localAtomIndex = get_local_id(0); const unsigned int localAtomIndex = get_local_id(0);
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
......
platforms/opencl/src/kernels/customGBEnergyN2_cpu.cl
View file @ 43c8e814
...@@ -255,7 +255,7 @@ __kernel void computeN2Energy( ...@@ -255,7 +255,7 @@ __kernel void computeN2Energy(
for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) { for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) {
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE+localAtomIndex; unsigned int j = y*TILE_SIZE+localAtomIndex;
#endif #endif
......
platforms/opencl/src/kernels/customGBValueN2.cl
View file @ 43c8e814
...@@ -229,7 +229,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4* ...@@ -229,7 +229,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
LOAD_ATOM1_PARAMETERS LOAD_ATOM1_PARAMETERS
const unsigned int localAtomIndex = get_local_id(0); const unsigned int localAtomIndex = get_local_id(0);
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
......
platforms/opencl/src/kernels/customGBValueN2_cpu.cl
View file @ 43c8e814
...@@ -228,7 +228,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4* ...@@ -228,7 +228,7 @@ __kernel void computeN2Value(__global const real4* restrict posq, __local real4*
for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) { for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) {
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE+localAtomIndex; unsigned int j = y*TILE_SIZE+localAtomIndex;
#endif #endif
......
platforms/opencl/src/kernels/gbsaObc.cl
View file @ 43c8e814
...@@ -232,7 +232,7 @@ __kernel void computeBornSum( ...@@ -232,7 +232,7 @@ __kernel void computeBornSum(
real4 posq1 = posq[atom1]; real4 posq1 = posq[atom1];
float2 params1 = global_params[atom1]; float2 params1 = global_params[atom1];
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
...@@ -617,7 +617,7 @@ __kernel void computeGBSAForce1( ...@@ -617,7 +617,7 @@ __kernel void computeGBSAForce1(
real4 posq1 = posq[atom1]; real4 posq1 = posq[atom1];
real bornRadius1 = global_bornRadii[atom1]; real bornRadius1 = global_bornRadii[atom1];
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
......
platforms/opencl/src/kernels/gbsaObc_cpu.cl
View file @ 43c8e814
...@@ -228,7 +228,7 @@ __kernel void computeBornSum( ...@@ -228,7 +228,7 @@ __kernel void computeBornSum(
for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) { for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) {
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE+localAtomIndex; unsigned int j = y*TILE_SIZE+localAtomIndex;
#endif #endif
...@@ -641,7 +641,7 @@ __kernel void computeGBSAForce1( ...@@ -641,7 +641,7 @@ __kernel void computeGBSAForce1(
for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) { for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) {
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE+localAtomIndex; unsigned int j = y*TILE_SIZE+localAtomIndex;
#endif #endif
......
platforms/opencl/src/kernels/nonbonded.cl
View file @ 43c8e814
...@@ -269,7 +269,7 @@ __kernel void computeNonbonded( ...@@ -269,7 +269,7 @@ __kernel void computeNonbonded(
LOAD_ATOM1_PARAMETERS LOAD_ATOM1_PARAMETERS
const unsigned int localAtomIndex = get_local_id(0); const unsigned int localAtomIndex = get_local_id(0);
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
......
platforms/opencl/src/kernels/nonbonded_cpu.cl
View file @ 43c8e814
...@@ -269,7 +269,7 @@ __kernel void computeNonbonded( ...@@ -269,7 +269,7 @@ __kernel void computeNonbonded(
for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) { for (int localAtomIndex = 0; localAtomIndex < TILE_SIZE; localAtomIndex++) {
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE+localAtomIndex; unsigned int j = y*TILE_SIZE+localAtomIndex;
#endif #endif
......
platforms/opencl/src/kernels/sort.cl
View file @ 43c8e814
...@@ -8,7 +8,7 @@ KEY_TYPE getValue(DATA_TYPE value) { ...@@ -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 * Sort a list that is short enough to entirely fit in local memory. This is executed as
* a single thread block. * 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. // Load the data into local memory.
for (int index = get_local_id(0); index < length; index += get_local_size(0)) 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, ...@@ -96,7 +96,7 @@ __kernel void computeRange(__global const DATA_TYPE* restrict data, uint length,
* Assign elements to buckets. * Assign elements to buckets.
*/ */
__kernel void assignElementsToBuckets(__global const DATA_TYPE* restrict data, uint length, uint numBuckets, __global const KEY_TYPE* restrict range, __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 #ifdef AMD_ATOMIC_WORK_AROUND
// Do a byte write to force all memory accesses to interactionCount to use the complete path. // 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. // 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 @ 43c8e814
...@@ -545,7 +545,7 @@ extern "C" __global__ void computeElectrostatics( ...@@ -545,7 +545,7 @@ extern "C" __global__ void computeElectrostatics(
data.force = make_real3(0); data.force = make_real3(0);
data.torque = make_real3(0); data.torque = make_real3(0);
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
......
plugins/amoeba/platforms/cuda/src/kernels/multipoleFixedField.cu
View file @ 43c8e814
...@@ -655,7 +655,7 @@ extern "C" __global__ void computeFixedField( ...@@ -655,7 +655,7 @@ extern "C" __global__ void computeFixedField(
data.bornRadius = bornRadii[atom1]; data.bornRadius = bornRadii[atom1];
#endif #endif
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
......
plugins/amoeba/platforms/cuda/src/kernels/multipoleInducedField.cu
View file @ 43c8e814
...@@ -514,7 +514,7 @@ extern "C" __global__ void computeInducedField( ...@@ -514,7 +514,7 @@ extern "C" __global__ void computeInducedField(
loadAtomData(data, atom1, posq, inducedDipole, inducedDipolePolar, dampingAndThole); loadAtomData(data, atom1, posq, inducedDipole, inducedDipolePolar, dampingAndThole);
#endif #endif
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #endif
...@@ -607,11 +607,6 @@ extern "C" __global__ void recordInducedDipolesForDIIS(const long long* __restri ...@@ -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, 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) { real* __restrict__ prevDipoles, real* __restrict__ prevDipolesPolar, real* __restrict__ prevErrors, int iteration, bool recordPrevErrors, real* __restrict__ matrix) {
extern __shared__ real2 buffer[]; 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; const real fieldScale = 1/(real) 0x100000000;
real sumErrors = 0; real sumErrors = 0;
real sumPolarErrors = 0; real sumPolarErrors = 0;
......
plugins/amoeba/platforms/cuda/src/kernels/pmeMultipoleElectrostatics.cu
View file @ 43c8e814
...@@ -612,7 +612,7 @@ extern "C" __global__ void computeElectrostatics( ...@@ -612,7 +612,7 @@ extern "C" __global__ void computeElectrostatics(
data.force = make_real3(0); data.force = make_real3(0);
data.torque = make_real3(0); data.torque = make_real3(0);
#ifdef USE_CUTOFF #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 #else
unsigned int j = y*TILE_SIZE + tgx; unsigned int j = y*TILE_SIZE + tgx;
#endif #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