Разделите версии библиотеки на стороне хоста и на стороне устройства CUDA.

У меня есть библиотека с некоторыми функциями __host__ __device__. У меня также есть гаджет #ifdef __CUDACC__, который гарантирует, что обычный компилятор C++ не увидит __host__ __device__ и, таким образом, сможет скомпилировать эти функции.

Теперь я хочу использовать скомпилированную версию функции моей библиотеки на стороне хоста в простом файле статической библиотеки C++ (.a в Linux) — и я даже хотел бы, чтобы эта библиотека была компилируемой, когда CUDA недоступен; и я хочу, чтобы скомпилированные версии на стороне устройства находились в отдельной статической библиотеке.

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

Что я делаю неправильно?


  • my_lib.hpp (Заголовок библиотеки):
#ifdef __CUDACC__
__host__ __device__
#endif
void foo(int*x, int* y);
int bar();
  • my_lib.cu (источник библиотеки):
#include "my_lib.hpp"

#ifdef __CUDACC__
__host__ __device__
#endif
void foo(int*x, int* y)  { *x = *y; }

int bar() { return 5; }
  • main.cu (тестовая программа):
#include "my_lib.hpp"

__global__ void my_kernel() {
  int z { 78 };
  int w { 90 };
  foo(&z,&w);
}

int main() {
  int z { 123 };
  int w { 456 };
  foo(&z,&w);
  my_kernel<<<1,1>>>();
  cudaDeviceSynchronize();
  cudaDeviceReset();
}

Мои команды сборки:

c++ -c -x c++ -o my_lib-noncuda.o my_lib.cu
ar qc my_lib-noncuda.a my_lib-noncuda.o
ranlib my_lib-noncuda.a
nvcc -dc -o my_lib-cuda.o my_lib.cu
ar qc my_lib-cuda.a my_lib-cuda.o
ranlib my_lib-cuda.a
nvcc -dc -o main.rdc.o main.cu
nvcc -dlink -o main.o main.rdc.o my_lib-cuda.a
c++ -o main main.o my_lib-noncuda.a -lcudart

И ошибки, которые я получаю - по последней, связывающей, команде:

/usr/bin/ld: main.o: in function `__cudaRegisterLinkedBinary_39_tmpxft_00003f88_00000000_6_main_cpp1_ii_e7ab3416':
link.stub:(.text+0x5a): undefined reference to `__fatbinwrap_39_tmpxft_00003f88_00000000_6_main_cpp1_ii_e7ab3416'
/usr/bin/ld: main.o: in function `__cudaRegisterLinkedBinary_41_tmpxft_00003f69_00000000_6_my_lib_cpp1_ii_ab44b3f6':
link.stub:(.text+0xaa): undefined reference to `__fatbinwrap_41_tmpxft_00003f69_00000000_6_my_lib_cpp1_ii_ab44b3f6'
collect2: error: ld returned 1 exit status

Примечания:

  • Я использую CUDA 10.1 и g++ 9.2.1 на Devuan GNU/Linux.
  • Это продолжение удаленного вопроса; @talonmies прокомментировал, что мне лучше показать, что именно я сделал; и это несколько изменило вопрос.
  • Несколько связанный вопрос: этот.

person einpoklum    schedule 16.12.2019    source источник
comment
@talonmies: Теперь я покажу вам, что именно я делаю.   -  person einpoklum    schedule 16.12.2019
comment
Этот пример не работает за пределами ошибок, которые вы описываете. Нигде в этой последовательности сборки не должно быть основного, и также должна быть ошибка main not found, если только я не прочитал это неправильно.   -  person talonmies    schedule 17.12.2019
comment
@talonmies: Ну, ошибки такие, какие они есть, хотя я понимаю, что вы имеете в виду насчет main() и -dc. Дело в том, что если я отбрасываю -dc, я получаю ошибку компиляции: nvcc -o main.o main.cu приводит к Unresolved extern function '_Z3fooPiS_'.   -  person einpoklum    schedule 17.12.2019
comment
Я смущен тем, что вы хотите сделать. Я думал, вы хотите связать обычный файл статической библиотеки C++ с программой C++. Но, насколько я понимаю, вы хотели бы использовать компилятор C++ на последнем шаге только для компоновки и при этом иметь все материалы CUDA. Возможно, это то, что вы ищете: devblogs.nvidia.com/ — раздел Расширенное использование: использование другого компоновщика   -  person BlameTheBits    schedule 17.12.2019
comment
@Shadow: Не совсем так. Я хочу иметь разделенную библиотеку: один файл .a с версиями функций на стороне хоста, который я мог бы использовать в обычной компоновке C++ независимо от CUDA; и еще один файл .a с версиями функций на стороне устройства, которые я могу связать с моим кодом CUDA, который вызывает эти функции из ядра. Тестовая программа иллюстрирует второй вид использования.   -  person einpoklum    schedule 17.12.2019


Ответы (2)


Вот как вы можете создать две библиотеки, одна из которых содержит только функции CUDA-устройства, а другая — только функции хоста. Вы можете опустить «сложный» #if и #ifndef охранник. Но тогда в вашей библиотеке my_lib-cuda.a будет и "не-CUDA-код".

По другим вопросам смотрите вики-ответ сообщества @talonmies или ссылайтесь на ссылку, которую я уже разместил в комментариях: https://devblogs.nvidia.com/separate-compilation-linking-cuda-device-code/ — раздел «Расширенное использование: использование другого компоновщика».

my_lib.cu

#include "my_lib.hpp"

#ifdef __CUDA_ARCH__
__device__
#endif
#if (defined __CUDA_ARCH__) || (not defined __CUDACC__)
void foo(int*x, int* y)  { *x = *y; }
#endif

#ifndef __CUDACC__
int bar() { return 5; }
#endif

Процесс сборки библиотек остается прежним: (только изменены ar qc на ar rc для замены существующих файлов, чтобы вы не получили ошибку при перестроении без предварительного удаления библиотеки)

c++ -c -x c++ -o my_lib-noncuda.o my_lib.cu
ar rc my_lib-noncuda.a my_lib-noncuda.o
ranlib my_lib-noncuda.a
nvcc -dc -o my_lib-cuda.o my_lib.cu
ar rc my_lib-cuda.a my_lib-cuda.o 
ranlib my_lib-cuda.a 

Создание программы CUDA: (упрощено за счет использования только nvcc, а не c++, в качестве альтернативы посмотрите вики-ответ сообщества @talonmies)

nvcc -dc main.cu -o main.o
nvcc main.o my_lib-cuda.a my_lib-noncuda.a -o main

Ссылку на my_lib-noncuda.a можно опустить, если вы также опустите #if и #ifndef в my_lib.cu, как описано выше.

Создание программы на C++: (учитывая, что вокруг кода CUDA в main.cu имеется #ifdef __CUDACC__ охранников)

c++ -x c++ -c main.cu -o main.o
c++ main.o my_lib-noncuda.a -o main
person BlameTheBits    schedule 17.12.2019
comment
Дело в том, что если я выберу такой подход, у меня будут столкновения. То есть предположим, что мое приложение напрямую использует my_lib на стороне устройства и other_lib на стороне хоста. Теперь other_lib, в свою очередь, использует my_lib на стороне хоста. Когда я попытаюсь связать все вместе, функции на стороне хоста окажутся в двух конфликтующих местах, не так ли? - person einpoklum; 17.12.2019
comment
Если вы сделаете это, как я, my_lib будет состоять только из кода устройства или только хост-кода. Так что ваш пример не имеет смысла. А конфликты — это всегда то, на что нужно обращать внимание, и с тем, что я сделал, дополнительных потенциальных конфликтов нет. - person BlameTheBits; 17.12.2019
comment
Итак, это не удается. То есть вторая из двух строк для сборки программы CUDA завершается ошибкой: nvlink error : Multiple definition of '_Z3fooPiS_' in 'my_lib-cuda.a:my_lib-cuda.rdc.o', first defined in 'my_lib-cuda.a:my_lib-cuda.o' nvlink fatal : merge_elf failed — если ваши изменения не относятся к тому, что написал @talonmies. - person einpoklum; 18.12.2019
comment
Ой. Ты прав. Но только если я повторю некоторые шаги компиляции. На первой компиляции у меня все работает нормально. Но это, кажется, проблема создания библиотеки. При втором запуске my_lib-cuda.a не перезаписывается, а добавляется. - person BlameTheBits; 18.12.2019
comment
Действительно, проблема в том, как я создаю две библиотеки в вопросе. Мне нужно сделать что-то еще... но помните, что суть вопроса заключается в разделении кода хоста и устройства, а не в компиляции примера программы. - person einpoklum; 18.12.2019
comment
Да. Но разве не этим я занимаюсь? Библиотека my_lib-cuda.a содержит только код устройства, а my_lib-noncuda.a содержит только код хоста. - person BlameTheBits; 18.12.2019
comment
Я добавил недостающие части. Но создание библиотек действительно остается прежним. Конечно, вы также можете разделить основной файл, как это сделал @talonmies, но я пропустил это, так как понял ваш вопрос о том, как иметь библиотеку только для устройства и только для хоста. И в вашем примере вы решили не включать ядро ​​в библиотеку устройств, поэтому я подумал, что содержимое и структура основного файла не имеют значения. - person BlameTheBits; 18.12.2019
comment
Хорошо, теперь, когда это работает для меня, +1. Но я не совсем понимаю, почему это работает: (1.) Как переключатели на команды ar могут повлиять на связывание в более поздней команде? Нет необходимости в замене, пока позже. (2.) почему nvcc хочет принять функцию только для __device__ в .cu, когда в заголовке указано __host__ __device__? - person einpoklum; 18.12.2019
comment
(1) Библиотека — это просто набор реализаций. Его не волнует, содержит ли он две реализации одной и той же функции. Но когда вам действительно нужно вызвать функцию, то есть когда вы связываетесь с библиотекой, вы получите ошибку связывания. Поэтому вам нужно убедиться, что вы не просто добавляете файлы вслепую (то есть реализации функций), а добавляете их только в том случае, если они еще не присутствуют. Честно говоря, я никогда не добавляю вариант добавления вместо замены, если он уже присутствует, или добавления, если он отсутствует, когда-либо полезная вещь. - person BlameTheBits; 18.12.2019
comment
(2) На самом деле у вас может быть две совершенно разные реализации функции. (См. мою первую или вторую версию ответа). И, конечно же, всегда нормально иметь некоторые объявления без определений, если вы добавляете определения до того, как начнете их использовать (что делается на этапе компоновки). Подводя итог, можно сказать: версия хоста и версия устройства foo — это две совершенно разные и независимые функции. Для вас может быть просто удобно иметь функцию хоста и устройства с одним и тем же именем (и, возможно, также сохранить несколько строк повторяющегося кода). - person BlameTheBits; 18.12.2019
comment
И не возникнет ли у меня потом проблем с написанием __host__ __device__ функций, которые вызывают bar()? - person einpoklum; 18.12.2019
comment
Конечно, только хост-функция может вызывать bar (которая является хост-функцией). Но это всегда так. Или вы имеете в виду __host__ __device__ функции, которые вызывают foo? Я не вижу в этом проблемы. - person BlameTheBits; 18.12.2019

Давайте изменим ваш пример так, как я думаю, будет ваш фактический вариант использования. Модификация помещает main() в файл .cpp для компиляции g++, а код CUDA в отдельный файл .cu для компиляции nvcc. Это важно для того, чтобы ваша установка с двумя библиотеками работала; и оправдано, потому что «main содержит ядра CUDA, требующие отдельной компиляции и компоновки» — это особый крайний случай для модели компиляции nvcc.

Реструктурированный код:

main.cu:

include "my_lib.hpp"

__global__ void my_kernel() {
  int z { 78 };
  int w { 90 };
  foo(&z,&w);
}

int cudamain()
{
  my_kernel<<<1,1>>>();
  return 0;
}

main.cpp:

#include <cuda_runtime_api.h>
#include "my_lib.hpp"

extern int cudamain();

int main() {
  int z { 123 };
  int w { 456 };
  foo(&z,&w);
  cudamain();
  cudaDeviceSynchronize();
  cudaDeviceReset();
}

все остальные файлы остаются как в вопросе.

Команды, необходимые для сборки программы, теперь следующие:

c++ -c -x c++ -o my_lib-noncuda.o my_lib.cu
ar qc my_lib-noncuda.a my_lib-noncuda.o
ranlib my_lib-noncuda.a

nvcc -std=c++11 -dc -o my_lib-cuda.rdc.o my_lib.cu
ar qc my_lib-cuda.a my_lib-cuda.rdc.o
ranlib my_lib-cuda.a

# Until this line - identical to what you have tried in your question

nvcc -std=c++11 -c -rdc=true main.cu -o main.cu.o 
nvcc -dlink -o main.o main.cu.o my_lib-cuda.a

c++ -std=c++11 -o main main.cpp main.o main.cu.o -I/path/to/cuda/include \
    -L/path/to/cuda/lib64 my_lib-cuda.a my_lib-noncuda.a -lcudart -lcudadevrt

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

person Community    schedule 17.12.2019
comment
Ссылка на my_lib-noncuda.a здесь не нужна, или нет? - person BlameTheBits; 17.12.2019
comment
Да. main вызывает хост-версию функции хоста/устройства - person talonmies; 17.12.2019
comment
Я просто скопировал точный код и команды компиляции и удалил только ссылку на my_lib-noncuda.a. У меня нет ошибки. - person BlameTheBits; 17.12.2019
comment
Честно говоря, я не смотрел на хост-сторону вещей, просто на отсутствующие зависимости и отсутствие основного - person talonmies; 17.12.2019
comment
@talonmies; С библиотеками все еще есть дублирование, что является проблемой в моей собственной попытке сборки в вопросе. Как предполагает @Shadow, my_lib-cuda.a имеет версии foo() как на стороне хоста, так и на стороне устройства, а также версию bar() на стороне хоста. И если вы переключите порядок компоновки в вашей последней команде сборки - сначала -cuda.a, затем -noncuda.a (что вполне вероятно произойдет непреднамеренно с генераторами системы сборки или просто по наивности) - вы получите ошибки: - person einpoklum; 18.12.2019
comment
@talonmies; Кроме того, вы правы насчет моего предполагаемого использования, поэтому модификация в основном приемлема :-) - person einpoklum; 18.12.2019
comment
Ваше редактирование комментария TLDR было неправильным, и я удалил его. Исправление вызвано не разделением main и cuda, а дополнительными объектами, которые необходимо связать. Если вы разделите основную часть, скомпилируете и свяжете, как вы это делали, ошибка все равно будет сохраняться. Моя единственная точка зрения заключалась в том, что вам нужно будет сделать что-то совсем другое, если вы хотите, чтобы ядра связывались в той же единице перевода, что и ваша основная. И я не получаю повторяющиеся ошибки символов, используя этот точный рецепт в этом точном коде - person talonmies; 18.12.2019