Nas versões anteriores do CUDA, o atomicAdd não era implementado para duplos; portanto, é comum implementá-lo comoaqui. Com o novo CUDA 8 RC, encontro problemas ao tentar compilar meu código, que inclui essa função. Eu acho que isso se deve ao fato de que, com o Pascal e o Compute Capability 6.0, uma versão dupla nativa do atomicAdd foi adicionada, mas de alguma forma isso não é ignorado adequadamente nos recursos de computação anteriores.

O código abaixo usado para compilar e executar bem com as versões CUDA anteriores, mas agora recebo este erro de compilação:

test.cu(3): error: function "atomicAdd(double *, double)" has already been defined

Mas se eu remover minha implementação, em vez disso, recebo este erro:

test.cu(33): error: no instance of overloaded function "atomicAdd" matches the argument list
            argument types are: (double *, double)

Devo acrescentar que só vejo isso se compilar com-arch=sm_35 ou similar. Se eu compilar com-arch=sm_60 Eu recebo o comportamento esperado, ou seja, apenas o primeiro erro e a compilação bem-sucedida no segundo caso.

Editar: Além disso, é específico paraatomicAdd - se eu mudar o nome, funciona bem.

Realmente parece um bug do compilador. Alguém pode confirmar que esse é o caso?

Código de exemplo:

__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)));
    } while (assumed != old);
    return __longlong_as_double(old);
}

__global__ void kernel(double *a)
{
    double b=1.3;
    atomicAdd(a,b);
}

int main(int argc, char **argv)
{
    double *a;
    cudaMalloc(&a,sizeof(double));

    kernel<<<1,1>>>(a);

    cudaFree(a);
    return 0;
}

Edit: Eu recebi uma resposta da Nvidia que reconhece esse problema, e aqui está o que os desenvolvedores dizem sobre isso:

A arquitetura sm_60, recentemente suportada no CUDA 8.0, possui a função nativa fp64 atomicAdd. Devido às limitações de nossa cadeia de ferramentas e linguagem CUDA, a declaração dessa função precisa estar presente mesmo quando o código não está sendo compilado especificamente para sm_60. Isso causa um problema no seu código porque você também define uma função atomicAdd do fp64.

Funções internas do CUDA, como atomicAdd, são definidas pela implementação e podem ser alteradas entre as versões do CUDA. Os usuários não devem definir funções com o mesmo nome que qualquer função interna do CUDA. Sugerimos que você renomeie sua função atomicAdd para uma que não seja a mesma que qualquer função interna do CUDA.