error: function "atomicAdd(double *, double)" has already been defined

Question

I get this error while trying to compile the caffe derivative DeepLab_v2 on Ubuntu 14.04.5 with Cuda 8.0.

Does anyone know how to solve this?

DeepLab_v2 compiles fine on another computer that has Cuda 7.5, but since in my current computer I have a Pascal Titan X, I probably need to use Cuda 8.0.

CUDA 8.0 provides a definition of `atomicAdd` on `double` quantities that was not present in previous CUDA toolkits. The code you are working with also apparently provides its own definition/implementation, and this is the source of the error message. The correct fix is to make source code changes to the software in question to make it compatible with CUDA 8. [This question](http://stackoverflow.com/questions/37566987/cuda-atomicadd-for-doubles-definition-error) may be of interest. — Robert Crovella, Sep 01 '16 at 14:58
Hmm that answer is marked as correct, however it's not crystal clear to me, what needs to be done. Do I have to insert the `#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600 #else __device__ double atomicAdd(double* a, double b) { return b; } #endif` on top of every `.cpp` file in caffe that uses the `atomicAdd` function and then recompile Caffe? Do you have experience with that? — mcExchange, Sep 01 '16 at 15:57
Apparantly it does work, I will post the complete modification as an answer. Thanks for your help! — mcExchange, Sep 02 '16 at 08:11

score 36 · Accepted Answer · edited Jul 17 '22 at 08:54

I finally got it working with the help of @Robert Crovella's comment. I had to modify the file common.cuh from the DeepLab_v2 master branch in the following way:

#ifndef CAFFE_COMMON_CUH_
#define CAFFE_COMMON_CUH_

  #include <cuda.h>

  #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 600
  static __inline__ __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;
    if (val==0.0)
      return __longlong_as_double(old);
    do {
      assumed = old;
      old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val +__longlong_as_double(assumed)));
    } while (assumed != old);
    return __longlong_as_double(old);
  }
  #endif

#endif

Interestingly I didn't even have the `common.cuh` file, I had to create it, was this the case for you as well? — jlcv, Nov 25 '16 at 01:09

error: function "atomicAdd(double *, double)" has already been defined

1 Answers1