CUDA: почему невозможно определить статические глобальные функции-члены?

При компиляции приведенного ниже кода с помощью nvcc (CUDA 5.0) появляется ошибка "недопустимая комбинация квалификаторов памяти", так как, видимо, невозможно иметь глобальные ядра в классе.

class A
{
public:
    __global__ static void kernel();
};

__global__ void A::kernel()
{}

Я могу понять это ограничение при работе с нестатическими членами, но почему ошибка все еще возникает, когда ядро ​​объявлено статическим? Вызов таких членов ничем не отличается от вызова функции, когда она объявлена ​​в пространстве имен (в данном случае A).

A::kernel <<< 1, 1 >>> ();

Есть ли причина, по которой я не понимаю, почему это не реализовано (пока)?

EDIT: Судя по ответам как в ответах, так и в комментариях, я недостаточно ясно выразил свой вопрос. Мой вопрос не в том, почему появляется ошибка. Очевидно, это потому, что он не был реализован. Мой вопрос заключается в том, почему это не реализовано. До сих пор я не мог придумать причину, по которой эта функция не может быть реализована. Я понимаю, что мог забыть об особом случае, который усложнил бы дело, отсюда и вопрос.

Причины, по которым я считаю это разумной функцией, следующие:

  • Статическая функция не имеет указателя this. Таким образом, даже если ядро ​​вызывается для объекта, который находится на хосте, нет конфликта при доступе к его данным, так как эти данные изначально недоступны (данные из какого объекта? ?).
  • Вы можете возразить, что если у класса есть связанные с ним статические данные, живущие на хосте, они в принципе должны быть доступны из статического ядра. Однако статические данные также не поддерживаются, так что опять же нет конфликта.
  • Вызов статического ядра для объекта на хосте (A a; a.staticKernel<<<...,...>>>();)) был бы полностью эквивалентен вызову его вообще без объекта (A::staticKernel<<<...,...>>>();), как мы привыкли в обычном C++.

Что мне не хватает?


c++ member static global cuda
person JorenHeit    schedule 31.08.2013    source источник
comment
ваша цель - иметь класс с _ глобальной _ функцией?   -  person 4pie0    schedule 31.08.2013
comment
Мой вопрос в том, почему это не было реализовано. До сих пор я не мог придумать причину, по которой эта функция не может быть реализована. - есть миллион вещей, которые можно было бы реализовать, но в этом нет необходимости, и они остаются нереализованными   -  person 4pie0    schedule 01.09.2013
comment
Итак, ваш ответ: это просто не так.?   -  person JorenHeit    schedule 01.09.2013
comment
Итак, настоящая цель этого вопроса - обсудить, почему объектная модель CUDA такая, какая она есть? Это не совсем вопрос переполнения стека. Голосуйте за закрытие, поскольку в первую очередь основано на мнении.   -  person talonmies    schedule 01.09.2013
comment
Я хочу отправить запрос функции в NVidia, потому что мне показалось странным, что это не поддерживается. Прежде чем я это сделаю, я хочу убедиться, что нет никакой очевидной причины, по которой это так, как есть. Видимо, нет, и, видимо, это причина, чтобы понизить меня. Это не имеет ничего общего с мнением или дебатами.   -  person JorenHeit    schedule 01.09.2013
comment
Никто из нас здесь не проектировал объектную модель CUDA. Поэтому никто из нас не может сказать, почему он не поддерживается, и поэтому любые ответы в лучшем случае спекулятивны. Если бы я догадывался, я бы сказал, что это потому, что это нарушает модель компиляции - функция __global__ одновременно компилирует как объект хоста, так и объект устройства. Классы и структуры CUDA могут быть созданы только в одном пространстве памяти. Казалось бы, это исключает наличие объектов __constant__ или __global__ внутри структур и классов.   -  person talonmies    schedule 01.09.2013
comment
@talonmies: я не мог знать заранее, что на мой вопрос нет очевидного ответа, поэтому я не мог предвидеть предположения. Пониженные голоса и негативный тон в ответах/комментариях совершенно не нужны, если вы спросите меня. Что касается вашего догадки: ничего не добавляется к классу/структуре (в памяти), когда вы объявляете внутри него статическую функцию. Объявление просто добавляет область действия функции, подобно пространству имен. Я свяжусь с NVidia, чтобы узнать, что они скажут.   -  person JorenHeit    schedule 02.09.2013

Ответы (1)


arrow_upward
3
arrow_downward

К счастью, примерно через 4 года после того, как был задан этот вопрос, clang 4.0 может компилировать язык CUDA. Рассмотрим этот пример:

class A
{
public:
    __global__ static void kernel();
};

__device__ void A::kernel()
{}

int main()
{
    A::kernel <<< 1, 1 >>> ();
};

Когда я пытаюсь скомпилировать его с помощью clang 4.0, я получаю следующую ошибку:

test.cu:7:1: error: kernel function 'kernel' must be a free function or static member function
__global__ void A::kernel()
^
/usr/local/cuda/include/host_defines.h:191:9: note: expanded from macro '__global__'
        __location__(global)
        ^
/usr/local/cuda/include/host_defines.h:88:9: note: expanded from macro '__location__'
        __annotate__(a)
        ^
/usr/local/cuda/include/host_defines.h:86:9: note: expanded from macro '__annotate__'
        __attribute__((a))
        ^
test.cu:7:20: error: __host__ function 'kernel' cannot overload __global__ function 'kernel'
__global__ void A::kernel()
                   ^
test.cu:4:28: note: previous declaration is here
    __global__ static void kernel();
                           ^
2 errors generated.

Чтобы устранить эти ошибки, я встроил определение ядра в объявление класса:

class A
{
public:
    __global__ static void kernel()
    {
        // implementation would go here
    }
};

Тогда clang 4.0 успешно скомпилирует его и сможет запустить без ошибок. Так что это явно не ограничение языка CUDA, а его де-факто стандартный компилятор. Кстати, у nvcc есть много подобных неоправданных ограничений, которых нет у clang.

person Jakub Klinkovský    schedule 14.05.2017
comment
@talonmies Раздел E. Поддержка языка C/C++ руководства по программированию — это просто (частичная) спецификация того, что nvcc поддерживает или не поддерживает. Это не имеет ничего общего со спецификацией языка CUDA, которой, к сожалению, не существует. B. Раздел «Расширения языка C» закрыт, но он смешан с nvcc особенностями и API среды выполнения CUDA. - person Jakub Klinkovský; 15.05.2017
comment
В настоящее время (CUDA 10.2) это все еще не работает, если ядро ​​шаблонно, например. template<typename T> __global__ static void kernel(T something){}. Получение warning: inline qualifier ignored for "__global__" function без встроенного квалификатора, а затем error: illegal combination of memory qualifiers в той же строке с объявлением функции. - person Greg Kramida; 08.06.2020
comment
Могу подтвердить, что CUDA 11.1 все еще имеет этот недостаток в nvcc. То же сообщение об ошибке и все такое. Теперь придется придумывать совершенно новый дизайн. - person Meta.x.gdb; 12.03.2021