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
atomicAdd
ondouble
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 may be of interest.#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 theatomicAdd
function and then recompile Caffe? Do you have experience with that?