atomicAdd com double em CUDA

Olá a todos,

neste tutorial vou mostrar como somar duas variáveis do tipo “double” atomicamente em CUDA. Andei pesquisando na net, mas pelo que vi a NVidia ainda não tem suporte para esta função. A função “atomicAdd” só suporta no máximo variáveis do tipo “float”. Pesquisando nos foruns da NVidia vi que poderia usar esta função abaixo para contornar a situação:

/*
* Funcao para adicionar atomicamente dois valores double em CUDA
*/
__device__ inline void myAtomicAdd(double *address, double value)  //See CUDA official forum
{
    unsigned long long oldval, newval, readback;

    oldval = __double_as_longlong(*address);
    newval = __double_as_longlong(__longlong_as_double(oldval) + value);

    while ((readback=atomicCAS((unsigned long long *)address, oldval, newval)) != oldval){
        oldval = readback;
        newval = __double_as_longlong(__longlong_as_double(oldval) + value);
    }
}

E para compilar ainda é necessário utilizar a opção de compilação -arch compute_20 (isso se o compute capability de sua placa for 2.0)

Até agora ela me foi útil. Enquanto que a NVidia não dá suporte a este tipo de função, nos resta apenas utilizar esta função temporariamente.

Me avisem assim que esta situação mudar.

Até a próxima!

Anúncios

Deixe um comentário

Preencha os seus dados abaixo ou clique em um ícone para log in:

Logotipo do WordPress.com

Você está comentando utilizando sua conta WordPress.com. Sair / Alterar )

Imagem do Twitter

Você está comentando utilizando sua conta Twitter. Sair / Alterar )

Foto do Facebook

Você está comentando utilizando sua conta Facebook. Sair / Alterar )

Foto do Google+

Você está comentando utilizando sua conta Google+. Sair / Alterar )

Conectando a %s