Merge branch 'devel' into olz_documentation

.github/workflows/build.yml

@@ -18,6 +18,9 @@ jobs:
         - {cc: gcc, cxx: g++}
 
         include:
+          - os: 'ubuntu-20.04'
+            cuda: {version: '11.3.1', method: 'network'}
+            compiler: {cc: gcc-9, cxx: g++-9, update-alternatives: True}
           - os: 'ubuntu-20.04'
             cuda: {version: '10.2.89', method: 'local'}
             compiler: {cc: gcc-8, cxx: g++-8, update-alternatives: True}

CMakeLists.txt

@@ -80,11 +80,18 @@ if(${XMIPP_USE_CUDA})
       string(APPEND CMAKE_CUDA_FLAGS " --expt-extended-lambda")
 
       if(CMAKE_VERSION VERSION_GREATER_EQUAL "3.18")
-        set(CMAKE_CUDA_ARCHITECTURES 60 75)
+        list(APPEND CMAKE_CUDA_ARCHITECTURES 60 75)
 
         if(${CUDAToolkit_VERSION} VERSION_GREATER_EQUAL "11")
           list(APPEND CMAKE_CUDA_ARCHITECTURES 86)
         endif()
+      else()
+        string(APPEND CMAKE_CUDA_FLAGS " -gencode arch=compute_60,code=sm_60")
+        string(APPEND CMAKE_CUDA_FLAGS " -gencode arch=compute_75,code=sm_75")
+
+        if(${CUDAToolkit_VERSION} VERSION_GREATER_EQUAL "11")
+          string(APPEND CMAKE_CUDA_FLAGS " -gencode arch=compute_86,code=sm_86")
+        endif()
       endif()
 
     else()

src/xmipp/libraries/reconstruction_cuda/cuda_angular_sph_alignment.cu

@@ -5,6 +5,7 @@
 #ifndef KTT_USED
 #include "cuda_volume_deform_sph_defines.h"
 #include "cuda_angular_sph_alignment.h"
+#include "cuda_compatibility.h"
 #endif
 
 namespace AngularAlignmentGpu {
@@ -21,26 +22,6 @@ using PrecisionType3 = double3;
 #define SIN sin
 #define CUDA_FLOOR floor
 
-#if __CUDA_ARCH__ < 600
-__device__ double atomicAdd(double* address, double val)
-{
-    unsigned long long int* address_as_ull =
-                              (unsigned long long int*)address;
-    unsigned long long int old = *address_as_ull, assumed;
-
-    do {
-        assumed = old;
-        old = atomicCAS(address_as_ull, assumed,
-                        __double_as_longlong(val +
-                               __longlong_as_double(assumed)));
-
-    // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
-    } while (assumed != old);
-
-    return __longlong_as_double(old);
-}
-#endif
-
 #else
 // Types
 using PrecisionType = float;