Skip to content

Commit 8c3e915

Browse files
committed
Merge branch 'release/v2.2.4'
2 parents b922a5e + 18c4d6d commit 8c3e915

12 files changed

+63
-11
lines changed

Bassetti_Erskine.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
# This file is part of the code:
99
#
1010
#
11-
# PyPIC Version 2.2.3
11+
# PyPIC Version 2.2.4
1212
#
1313
#
1414
# Author and contact: Giovanni IADAROLA

FFT_OpenBoundary.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
# This file is part of the code:
99
#
1010
#
11-
# PyPIC Version 2.2.3
11+
# PyPIC Version 2.2.4
1212
#
1313
#
1414
# Author and contact: Giovanni IADAROLA

FFT_OpenBoundary_SquareGrid.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
# This file is part of the code:
99
#
1010
#
11-
# PyPIC Version 2.2.3
11+
# PyPIC Version 2.2.4
1212
#
1313
#
1414
# Author and contact: Giovanni IADAROLA

FFT_PEC_Boundary_SquareGrid.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
# This file is part of the code:
99
#
1010
#
11-
# PyPIC Version 2.2.3
11+
# PyPIC Version 2.2.4
1212
#
1313
#
1414
# Author and contact: Giovanni IADAROLA

FiniteDifferences_ShortleyWeller_SquareGrid.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
# This file is part of the code:
99
#
1010
#
11-
# PyPIC Version 2.2.3
11+
# PyPIC Version 2.2.4
1212
#
1313
#
1414
# Author and contact: Giovanni IADAROLA

FiniteDifferences_ShortleyWeller_SquareGrid_extrapolation.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
# This file is part of the code:
99
#
1010
#
11-
# PyPIC Version 2.2.3
11+
# PyPIC Version 2.2.4
1212
#
1313
#
1414
# Author and contact: Giovanni IADAROLA

FiniteDifferences_Staircase_SquareGrid.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
# This file is part of the code:
99
#
1010
#
11-
# PyPIC Version 2.2.3
11+
# PyPIC Version 2.2.4
1212
#
1313
#
1414
# Author and contact: Giovanni IADAROLA

GPU/p2m/p2m_kernels.cu

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010
// 2017-05-22 edit based on:
1111
// http://stackoverflow.com/questions/39274472/error-function-atomicadddouble-double-has-already-been-defined
1212

13+
/*
1314
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600
1415
1516
#else
@@ -33,6 +34,31 @@ static __inline__ __device__ double atomicAdd(double* address, double val)
3334
return __longlong_as_double(old);
3435
}
3536
#endif
37+
*/
38+
39+
#ifdef __CUDA_ARCH__
40+
#if __CUDA_ARCH__ < 600
41+
static __inline__ __device__ double atomicAdd(double* address, double val)
42+
{
43+
unsigned long long int* address_as_ull =
44+
(unsigned long long int*)address;
45+
unsigned long long int old = *address_as_ull, assumed;
46+
47+
if (val==0.0)
48+
return __longlong_as_double(old);
49+
do {
50+
assumed = old;
51+
old = atomicCAS(address_as_ull, assumed,
52+
__double_as_longlong(val +
53+
__longlong_as_double(assumed)));
54+
55+
// Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
56+
} while (assumed != old);
57+
58+
return __longlong_as_double(old);
59+
}
60+
#endif
61+
#endif
3662

3763

3864
extern "C" {

GPU/p2m/p2m_kernels_inclmeshing.cu

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010
// 2017-05-22 edit based on:
1111
// http://stackoverflow.com/questions/39274472/error-function-atomicadddouble-double-has-already-been-defined
1212

13+
/*
1314
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600
1415
1516
#else
@@ -33,6 +34,31 @@ static __inline__ __device__ double atomicAdd(double* address, double val)
3334
return __longlong_as_double(old);
3435
}
3536
#endif
37+
*/
38+
39+
#ifdef __CUDA_ARCH__
40+
#if __CUDA_ARCH__ < 600
41+
static __inline__ __device__ double atomicAdd(double* address, double val)
42+
{
43+
unsigned long long int* address_as_ull =
44+
(unsigned long long int*)address;
45+
unsigned long long int old = *address_as_ull, assumed;
46+
47+
if (val==0.0)
48+
return __longlong_as_double(old);
49+
do {
50+
assumed = old;
51+
old = atomicCAS(address_as_ull, assumed,
52+
__double_as_longlong(val +
53+
__longlong_as_double(assumed)));
54+
55+
// Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
56+
} while (assumed != old);
57+
58+
return __longlong_as_double(old);
59+
}
60+
#endif
61+
#endif
3662

3763

3864
extern "C" {

PyPIC_Scatter_Gather.py

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
# This file is part of the code:
99
#
1010
#
11-
# PyPIC Version 2.2.3
11+
# PyPIC Version 2.2.4
1212
#
1313
#
1414
# Author and contact: Giovanni IADAROLA
@@ -66,7 +66,7 @@ class PyPIC_Scatter_Gather(object):
6666
def __init__(self, x_aper=None, y_aper=None, dx=None, dy=None, xg=None, yg=None,
6767
x_min=None, x_max=None, y_min=None, y_max=None, *args, **kwargs):
6868

69-
print 'PyPIC Version 2.2.3'
69+
print 'PyPIC Version 2.2.4'
7070

7171
if xg!=None and yg!=None:
7272
assert(x_aper==None and y_aper==None and dx==None and dy==None)

0 commit comments

Comments
 (0)