CUDA atomicAdd для ошибки определения двойников

В предыдущих версиях CUDA, то atomicAdd не был реализован в парном разряде, поэтому она является общей для реализации этого, как здесь. С новым CUDA 8 RC я сталкиваюсь с проблемами, когда пытаюсь скомпилировать свой код, который включает такую функцию. Я предполагаю, что это связано с тем, что с Pascal и Compute Capability 6.0 была добавлена собственная двойная версия atomicAdd, но почему-то это должным образом не игнорируется для предыдущих вычислительных возможностей.

Приведенный ниже код используется для компиляции и запуска отлично с предыдущими версиями CUDA, но теперь я получаю эту ошибку компиляции: 

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

Но если я удалю свою реализацию, я вместо этого получу эту ошибку: 

test.cu(33): error: no instance of overloaded function "atomicAdd" matches the argument list
            argument types are: (double *, double)
Я должен добавить, что я вижу это только в том случае, если я компилирую с -arch=sm_35 или подобным образом. Если я компилирую с -arch=sm_60, я получаю ожидаемое поведение, т. е. только первую ошибку, и успешную компиляцию во втором случае.

Edit: кроме того, он специфичен для atomicAdd - если я изменю имя, он хорошо работает.

Это действительно похоже на компилятор жук. Может ли кто-то еще подтвердить, что это так?

Пример кода: 

__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: я получил ответ от Nvidia, которые признают эту проблему, и вот что говорят об этом разработчики:

Архитектура sm_60, которая недавно поддерживается в CUDA 8.0, имеет собственная функция fp64 atomicAdd. Из-за ограничений нашего инструментария и языка CUDA, то в декларации эта функция должна присутствуйте даже тогда, когда код не используется специально составлено для sm_60. Это вызывает проблему в коде, так как вы также определяете функция fp64 atomicAdd.

Встроенные функции CUDA, такие как atomicAdd, определяются реализацией и может быть изменен между выпусками CUDA. Пользователи не должны определять функции с теми же именами, что и любые встроенные функции CUDA. Мы бы так и сделали предлагаю вам переименовать вашу функцию atomicAdd в ту, которая не является то же самое, что и любые встроенные функции CUDA.

1 8
cuda atomic nvidia

1 ответ:

Этот аромат atomicAdd - новый метод, введенный для вычислительной способности 6.0. Вы можете сохранить свою предыдущую реализацию других вычислительных возможностей, защищая ее с помощью определения макросов 

#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600
#else
<... place here your own pre-pascal atomicAdd definition ...>
#endif

Этот макрос, названный макросом идентификации архитектуры, задокументирован здесь :

[7]}5.7.4. Макрос Идентификации Виртуальной Архитектуры

Макросу идентификации архитектуры __CUDA_ARCH__ присваивается трехзначная строка значений xy0 (оканчивающаяся на литерал 0) во время каждый этап компиляции nvcc 1, который компилируется для compute_xy.

Этот макрос может быть использован при реализации функций GPU для определения виртуальной архитектуры, для которой он в настоящее время компилируется. Код хоста (код, не относящийся к GPU) не должен зависеть от него.

Я предполагаю, что NVIDIA не поместила его для предыдущего CC, чтобы избежать конфликта для пользователей, определяющих его и не переходящих к вычислительным возможностям >= 6.x. я бы не стал считать это ошибкой, а скорее доставкой релиза практика.

EDIT : macro guard был неполным (фиксированным) - вот полный пример.

#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600
#else
__device__ double atomicAdd(double* a, double b) { return b; }
#endif

__device__ double s_global ;
__global__ void kernel () { atomicAdd (&s_global, 1.0) ; }


int main (int argc, char* argv[])
{
        kernel<<<1,1>>> () ;
        return ::cudaDeviceSynchronize () ;
}

Компиляция с: 

$> nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2016 NVIDIA Corporation
Built on Wed_May__4_21:01:56_CDT_2016
Cuda compilation tools, release 8.0, V8.0.26

Командные строки (обе успешные): 

$> nvcc main.cu -arch=sm_60
$> nvcc main.cu -arch=sm_35

Вы можете узнать, почему он работает с включаемым файлом: sm_60_atomic_functions.h, где метод не объявляется, если __CUDA_ARCH__ меньше 600.

11