Cuda виртуальный класс

Я хотел бы выполнить некоторые виртуальные методы в ядре cuda, но вместо создания объекта в том же ядре я хотел бы создать его на хосте и скопировать его в память gpu.

Я успешно создаю объекты в ядре и вызываю виртуальный метод. Проблема возникает при копировании объекта. Это имеет смысл, потому что, очевидно, указатель виртуальной функции является поддельным.
То, что происходит, это просто «запуск сетки Cuda не удался», по крайней мере, так говорит Nsight.
Но при взгляде на SASS происходит сбой при разыменовании указателя виртуальной функции, что имеет смысл.

Я, конечно, использую Cuda 4.2, а также компилирую с «compute_30» на подходящей карте.

Так каков рекомендуемый путь? Или эта функция просто не поддерживается?

У меня была идея запустить сначала другое ядро, которое создает фиктивные объекты и извлекает указатель виртуальной функции, чтобы «исправить» мои объекты перед их копированием. К сожалению, это на самом деле не работает (пока не разобрался), а также было бы уродливым решением.

Постскриптум Это на самом деле повтор этот вопрос, на который, к сожалению, так и не был дан полный ответ.

Редактировать :

Поэтому я нашел способ сделать то, что хотел. Но просто для ясности: это совсем не ответ или решение, ответ уже был предоставлен, это всего лишь взлом, просто для удовольствия.

Итак, сначала давайте посмотрим, что делает Cuda при вызове виртуального метода, ниже отладка SASS

//R0 is the address of our object
LD.CG R0, [R0];
IADD R0, R0, 0x4;
NOP;
MOV R0, R0;
LD.CG R0, [R0];
...
IADD R0, RZ, R9;
MOV R0, R0;
LDC R0, c[0x2][R0];
...
BRX R0 - 0x5478

Таким образом, предполагая, что «c [0x2] [INDEX]» является константой для всех ядер, мы можем просто получить индекс для класса, просто запустив ядро ​​и выполнив это, где obj — это вновь созданный объект класса, который просматривает:

unsigned int index = *(unsigned int*)(*(unsigned int*)obj + 4);

Затем используйте что-то вроде этого:

struct entry
{
unsigned int vfptr;// := &vfref, thats our value to store in an object
int dummy;// := 1234, great for debugging
unsigned int vfref;// := &dummy
unsigned int index;
char ClassName[256];//use it as a key for a dict
};

Сохраните это в хосте, а также в памяти устройства (ячейки памяти — в памяти устройства), и на хосте вы можете использовать ClassName в качестве поиска объекта, который нужно «исправить».

Но опять же: я бы не стал использовать это ни в чем серьезном, потому что с точки зрения производительности, виртуальные функции совсем не велики.

5

Решение

То, что вы пытаетесь сделать, в настоящее время не поддерживается компилятором CUDA и средой выполнения (начиная с CUDA 5.0). Раздел D.2.6.3 Руководства по программированию CUDA C v5.0 гласит:

D.2.6.3 Виртуальные функции

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

Не разрешается передавать в качестве аргумента __global__ функционировать объект класса
с виртуальными функциями.

Таблица виртуальных функций размещается компилятором в глобальной или постоянной памяти.

Я рекомендую вам инкапсулировать данные вашего класса отдельно от функциональных возможностей класса. Например, хранить данные в структуре. Если вы планируете работать с массивами этих объектов, храните данные в структуре массивов (для производительности — за рамками этого вопроса). Выделите структуры данных на хосте, используя cudaMalloc, а затем передать данные ядру в качестве аргументов, а не передавать класс виртуальными методами.

Затем создайте свои объекты с помощью виртуальных методов на устройстве. Конструктор вашего класса с виртуальными методами будет принимать параметры ядра указателя устройства в качестве аргументов. Виртуальный устройство методы могут затем работать с данными устройства.

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

5

Другие решения

Других решений пока нет …

По вопросам рекламы ammmcru@yandex.ru
Adblock
detector