Skip to content

GitLab

Commit be73754b authored by Mark Friedrichs's avatar Mark Friedrichs
Browse files

Added cudaMemcpy for initialization of induced dipoles

parent 6bad9d44
Show whitespace changes
Inline Side-by-side
Showing with 15 additions and 26 deletions
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaMutualInducedAndGkFields.cu
View file @ be73754b
...@@ -249,8 +249,6 @@ void kInitializeMutualInducedAndGkField_kernel( ...@@ -249,8 +249,6 @@ void kInitializeMutualInducedAndGkField_kernel(
float* fixedEFieldPolar, float* fixedEFieldPolar,
float* fixedGkField, float* fixedGkField,
float* polarizability, float* polarizability,
float* inducedDipole,
float* inducedDipolePolar,
float* inducedDipoleS, float* inducedDipoleS,
float* inducedDipolePolarS ) float* inducedDipolePolarS )
{ {
...@@ -260,12 +258,9 @@ void kInitializeMutualInducedAndGkField_kernel( ...@@ -260,12 +258,9 @@ void kInitializeMutualInducedAndGkField_kernel(
{ {
fixedEField[pos] *= polarizability[pos]; fixedEField[pos] *= polarizability[pos];
inducedDipole[pos] = fixedEField[pos];
fixedEFieldPolar[pos] *= polarizability[pos]; fixedEFieldPolar[pos] *= polarizability[pos];
inducedDipolePolar[pos] = fixedEFieldPolar[pos];
fixedGkField[pos] *= polarizability[pos]; fixedGkField[pos] *= polarizability[pos];
inducedDipoleS[pos] = fixedEField[pos] + fixedGkField[pos]; inducedDipoleS[pos] = fixedEField[pos] + fixedGkField[pos];
inducedDipolePolarS[pos] = fixedEFieldPolar[pos] + fixedGkField[pos]; inducedDipolePolarS[pos] = fixedEFieldPolar[pos] + fixedGkField[pos];
...@@ -683,12 +678,13 @@ static void cudaComputeAmoebaMutualInducedAndGkFieldBySOR( amoebaGpuContext amoe ...@@ -683,12 +678,13 @@ static void cudaComputeAmoebaMutualInducedAndGkFieldBySOR( amoebaGpuContext amoe
amoebaGpu->psE_FieldPolar->_pDevData, amoebaGpu->psE_FieldPolar->_pDevData,
amoebaGpu->psGk_Field->_pDevData, amoebaGpu->psGk_Field->_pDevData,
amoebaGpu->psPolarizability->_pDevData, amoebaGpu->psPolarizability->_pDevData,
amoebaGpu->psInducedDipole->_pDevData,
amoebaGpu->psInducedDipolePolar->_pDevData,
amoebaGpu->psInducedDipoleS->_pDevData, amoebaGpu->psInducedDipoleS->_pDevData,
amoebaGpu->psInducedDipolePolarS->_pDevData ); amoebaGpu->psInducedDipolePolarS->_pDevData );
LAUNCHERROR("kInitializeMutualInducedAndGkField"); LAUNCHERROR("kInitializeMutualInducedAndGkField");
cudaMemcpy( amoebaGpu->psInducedDipole->_pDevData, amoebaGpu->psE_Field->_pDevData, 3*gpu->sim.paddedNumberOfAtoms*sizeof( float ), cudaMemcpyDeviceToDevice );
cudaMemcpy( amoebaGpu->psInducedDipolePolar->_pDevData, amoebaGpu->psE_FieldPolar->_pDevData, 3*gpu->sim.paddedNumberOfAtoms*sizeof( float ), cudaMemcpyDeviceToDevice );
#ifdef AMOEBA_DEBUG #ifdef AMOEBA_DEBUG
if( amoebaGpu->log ){ if( amoebaGpu->log ){
......
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaMutualInducedField.cu
View file @ be73754b
...@@ -115,9 +115,7 @@ void kInitializeMutualInducedField_kernel( ...@@ -115,9 +115,7 @@ void kInitializeMutualInducedField_kernel(
int numberOfAtoms, int numberOfAtoms,
float* fixedEField, float* fixedEField,
float* fixedEFieldPolar, float* fixedEFieldPolar,
float* polarizability, float* polarizability )
float* inducedDipole,
float* inducedDipolePolar )
{ {
int pos = __mul24(blockIdx.x,blockDim.x) + threadIdx.x; int pos = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
...@@ -125,10 +123,7 @@ void kInitializeMutualInducedField_kernel( ...@@ -125,10 +123,7 @@ void kInitializeMutualInducedField_kernel(
{ {
fixedEField[pos] *= polarizability[pos]; fixedEField[pos] *= polarizability[pos];
inducedDipole[pos] = fixedEField[pos];
fixedEFieldPolar[pos] *= polarizability[pos]; fixedEFieldPolar[pos] *= polarizability[pos];
inducedDipolePolar[pos] = fixedEFieldPolar[pos];
pos += blockDim.x*gridDim.x; pos += blockDim.x*gridDim.x;
} }
...@@ -488,11 +483,12 @@ static void cudaComputeAmoebaMutualInducedFieldBySOR( amoebaGpuContext amoebaGpu ...@@ -488,11 +483,12 @@ static void cudaComputeAmoebaMutualInducedFieldBySOR( amoebaGpuContext amoebaGpu
gpu->natoms, gpu->natoms,
amoebaGpu->psE_Field->_pDevData, amoebaGpu->psE_Field->_pDevData,
amoebaGpu->psE_FieldPolar->_pDevData, amoebaGpu->psE_FieldPolar->_pDevData,
amoebaGpu->psPolarizability->_pDevData, amoebaGpu->psPolarizability->_pDevData );
amoebaGpu->psInducedDipole->_pDevData,
amoebaGpu->psInducedDipolePolar->_pDevData );
LAUNCHERROR("AmoebaMutualInducedFieldSetup"); LAUNCHERROR("AmoebaMutualInducedFieldSetup");
cudaMemcpy( amoebaGpu->psInducedDipole->_pDevData, amoebaGpu->psE_Field->_pDevData, 3*gpu->sim.paddedNumberOfAtoms*sizeof( float ), cudaMemcpyDeviceToDevice );
cudaMemcpy( amoebaGpu->psInducedDipolePolar->_pDevData, amoebaGpu->psE_FieldPolar->_pDevData, 3*gpu->sim.paddedNumberOfAtoms*sizeof( float ), cudaMemcpyDeviceToDevice );
#ifdef AMOEBA_DEBUG #ifdef AMOEBA_DEBUG
if( amoebaGpu->log ){ if( amoebaGpu->log ){
......
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaPmeMutualInducedField.cu
View file @ be73754b
...@@ -237,19 +237,15 @@ static void kInitializeMutualInducedField_kernel( ...@@ -237,19 +237,15 @@ static void kInitializeMutualInducedField_kernel(
int numberOfAtoms, int numberOfAtoms,
float* fixedEField, float* fixedEField,
float* fixedEFieldPolar, float* fixedEFieldPolar,
float* polarizability, float* polarizability )
float* inducedDipole,
float* inducedDipolePolar )
{ {
int pos = __mul24(blockIdx.x,blockDim.x) + threadIdx.x; int pos = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
while( pos < 3*cSim.atoms ) while( pos < 3*cSim.atoms )
{ {
fixedEField[pos] *= polarizability[pos]; fixedEField[pos] *= polarizability[pos];
inducedDipole[pos] = fixedEField[pos];
fixedEFieldPolar[pos] *= polarizability[pos]; fixedEFieldPolar[pos] *= polarizability[pos];
inducedDipolePolar[pos] = fixedEFieldPolar[pos];
pos += blockDim.x*gridDim.x; pos += blockDim.x*gridDim.x;
} }
...@@ -555,11 +551,12 @@ static void cudaComputeAmoebaPmeMutualInducedFieldBySOR( amoebaGpuContext amoeba ...@@ -555,11 +551,12 @@ static void cudaComputeAmoebaPmeMutualInducedFieldBySOR( amoebaGpuContext amoeba
gpu->natoms, gpu->natoms,
amoebaGpu->psE_Field->_pDevData, amoebaGpu->psE_Field->_pDevData,
amoebaGpu->psE_FieldPolar->_pDevData, amoebaGpu->psE_FieldPolar->_pDevData,
amoebaGpu->psPolarizability->_pDevData, amoebaGpu->psPolarizability->_pDevData );
amoebaGpu->psInducedDipole->_pDevData,
amoebaGpu->psInducedDipolePolar->_pDevData );
LAUNCHERROR("AmoebaPmeMutualInducedFieldSetup"); LAUNCHERROR("AmoebaPmeMutualInducedFieldSetup");
cudaMemcpy( amoebaGpu->psInducedDipole->_pDevData, amoebaGpu->psE_Field->_pDevData, 3*gpu->sim.paddedNumberOfAtoms*sizeof( float ), cudaMemcpyDeviceToDevice );
cudaMemcpy( amoebaGpu->psInducedDipolePolar->_pDevData, amoebaGpu->psE_FieldPolar->_pDevData, 3*gpu->sim.paddedNumberOfAtoms*sizeof( float ), cudaMemcpyDeviceToDevice );
#ifdef AMOEBA_DEBUG #ifdef AMOEBA_DEBUG
if( amoebaGpu->log ){ if( amoebaGpu->log ){
......
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