@@ -324,11 +324,11 @@ bool get_expanded_box(const double rc, const Box& box, DFTD3::ExpandedBox& ebox)
324324}
325325
326326static __device__ void apply_mic_small_box (
327- const Box& box, const DFTD3::ExpandedBox& ebox, double & x12, double & y12, double & z12)
327+ const Box& box, const DFTD3::ExpandedBox& ebox, float & x12, float & y12, float & z12)
328328{
329- double sx12 = ebox.h [9 ] * x12 + ebox.h [10 ] * y12 + ebox.h [11 ] * z12;
330- double sy12 = ebox.h [12 ] * x12 + ebox.h [13 ] * y12 + ebox.h [14 ] * z12;
331- double sz12 = ebox.h [15 ] * x12 + ebox.h [16 ] * y12 + ebox.h [17 ] * z12;
329+ float sx12 = ebox.h [9 ] * x12 + ebox.h [10 ] * y12 + ebox.h [11 ] * z12;
330+ float sy12 = ebox.h [12 ] * x12 + ebox.h [13 ] * y12 + ebox.h [14 ] * z12;
331+ float sz12 = ebox.h [15 ] * x12 + ebox.h [16 ] * y12 + ebox.h [17 ] * z12;
332332 if (box.pbc_x == 1 )
333333 sx12 -= nearbyint (sx12);
334334 if (box.pbc_y == 1 )
@@ -362,9 +362,9 @@ static __global__ void find_neighbor_list_small_box(
362362{
363363 int n1 = blockIdx .x * blockDim .x + threadIdx .x ;
364364 if (n1 < N) {
365- double x1 = g_x[n1];
366- double y1 = g_y[n1];
367- double z1 = g_z[n1];
365+ float x1 = g_x[n1];
366+ float y1 = g_y[n1];
367+ float z1 = g_z[n1];
368368 int count_radial = 0 ;
369369 int count_angular = 0 ;
370370 for (int n2 = 0 ; n2 < N; ++n2) {
@@ -375,30 +375,30 @@ static __global__ void find_neighbor_list_small_box(
375375 continue ; // exclude self
376376 }
377377
378- double delta[3 ];
378+ float delta[3 ];
379379 delta[0 ] = box.cpu_h [0 ] * ia + box.cpu_h [1 ] * ib + box.cpu_h [2 ] * ic;
380380 delta[1 ] = box.cpu_h [3 ] * ia + box.cpu_h [4 ] * ib + box.cpu_h [5 ] * ic;
381381 delta[2 ] = box.cpu_h [6 ] * ia + box.cpu_h [7 ] * ib + box.cpu_h [8 ] * ic;
382382
383- double x12 = g_x[n2] + delta[0 ] - x1;
384- double y12 = g_y[n2] + delta[1 ] - y1;
385- double z12 = g_z[n2] + delta[2 ] - z1;
383+ float x12 = g_x[n2] + delta[0 ] - x1;
384+ float y12 = g_y[n2] + delta[1 ] - y1;
385+ float z12 = g_z[n2] + delta[2 ] - z1;
386386
387387 apply_mic_small_box (box, ebox, x12, y12, z12);
388388
389- float distance_square = float ( x12 * x12 + y12 * y12 + z12 * z12) ;
389+ float distance_square = x12 * x12 + y12 * y12 + z12 * z12;
390390 if (distance_square < rc_radial_sq) {
391391 g_NL_radial[count_radial * N + n1] = n2;
392- g_x12_radial[count_radial * N + n1] = float ( x12) ;
393- g_y12_radial[count_radial * N + n1] = float ( y12) ;
394- g_z12_radial[count_radial * N + n1] = float ( z12) ;
392+ g_x12_radial[count_radial * N + n1] = x12;
393+ g_y12_radial[count_radial * N + n1] = y12;
394+ g_z12_radial[count_radial * N + n1] = z12;
395395 count_radial++;
396396 }
397397 if (distance_square < rc_angular_sq) {
398398 g_NL_angular[count_angular * N + n1] = n2;
399- g_x12_angular[count_angular * N + n1] = float ( x12) ;
400- g_y12_angular[count_angular * N + n1] = float ( y12) ;
401- g_z12_angular[count_angular * N + n1] = float ( z12) ;
399+ g_x12_angular[count_angular * N + n1] = x12;
400+ g_y12_angular[count_angular * N + n1] = y12;
401+ g_z12_angular[count_angular * N + n1] = z12;
402402 count_angular++;
403403 }
404404 }
@@ -484,15 +484,15 @@ __device__ int find_neighbor_cell(
484484 int neighbor_cell = cell_id + zz * nx * ny + yy * nx + xx;
485485 if (cell_id_x + xx < 0 )
486486 neighbor_cell += nx;
487- if (cell_id_x + xx >= nx)
487+ else if (cell_id_x + xx >= nx)
488488 neighbor_cell -= nx;
489489 if (cell_id_y + yy < 0 )
490490 neighbor_cell += ny * nx;
491- if (cell_id_y + yy >= ny)
491+ else if (cell_id_y + yy >= ny)
492492 neighbor_cell -= ny * nx;
493493 if (cell_id_z + zz < 0 )
494494 neighbor_cell += nz * ny * nx;
495- if (cell_id_z + zz >= nz)
495+ else if (cell_id_z + zz >= nz)
496496 neighbor_cell -= nz * ny * nx;
497497
498498 return neighbor_cell;
@@ -547,11 +547,11 @@ __global__ void find_dftd3_coordination_number_large_box(
547547 if (n1 == n2) {
548548 continue ;
549549 }
550- double x12double = g_x[n2] - x1;
551- double y12double = g_y[n2] - y1;
552- double z12double = g_z[n2] - z1;
553- apply_mic (box, x12double, y12double, z12double );
554- float r12[3 ] = {float (x12double), float (y12double), float (z12double) };
550+ float x12 = g_x[n2] - x1;
551+ float y12 = g_y[n2] - y1;
552+ float z12 = g_z[n2] - z1;
553+ apply_mic (box, x12, y12, z12 );
554+ float r12[3 ] = {x12, y12, z12 };
555555 float d12_2 = r12[0 ] * r12[0 ] + r12[1 ] * r12[1 ] + r12[2 ] * r12[2 ];
556556 if (d12_2 < rc * rc) {
557557 int atomic_number_2 = dftd3_para.atomic_number [g_type[n2]];
@@ -637,11 +637,11 @@ __global__ void find_dftd3_force_large_box(
637637 if (n1 == n2) {
638638 continue ;
639639 }
640- double x12double = g_x[n2] - x1;
641- double y12double = g_y[n2] - y1;
642- double z12double = g_z[n2] - z1;
643- apply_mic (box, x12double, y12double, z12double );
644- float r12[3 ] = {float (x12double), float (y12double), float (z12double) };
640+ float x12 = g_x[n2] - x1;
641+ float y12 = g_y[n2] - y1;
642+ float z12 = g_z[n2] - z1;
643+ apply_mic (box, x12, y12, z12 );
644+ float r12[3 ] = {x12, y12, z12 };
645645 float d12_2 = r12[0 ] * r12[0 ] + r12[1 ] * r12[1 ] + r12[2 ] * r12[2 ];
646646 if (d12_2 < rc * rc) {
647647 int atomic_number_2 = dftd3_para.atomic_number [g_type[n2]];
@@ -805,11 +805,11 @@ __global__ void find_dftd3_force_extra_large_box(
805805 if (n1 == n2) {
806806 continue ;
807807 }
808- double x12double = g_x[n2] - x1;
809- double y12double = g_y[n2] - y1;
810- double z12double = g_z[n2] - z1;
811- apply_mic (box, x12double, y12double, z12double );
812- float r12[3 ] = {float (x12double), float (y12double), float (z12double) };
808+ float x12 = g_x[n2] - x1;
809+ float y12 = g_y[n2] - y1;
810+ float z12 = g_z[n2] - z1;
811+ apply_mic (box, x12, y12, z12 );
812+ float r12[3 ] = {x12, y12, z12 };
813813 float d12_2 = r12[0 ] * r12[0 ] + r12[1 ] * r12[1 ] + r12[2 ] * r12[2 ];
814814 if (d12_2 < rc * rc) {
815815 int atomic_number_2 = dftd3_para.atomic_number [g_type[n2]];
0 commit comments