CUDA atomicAdd para erro de definição dupla
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.