Как стать автором
Обновить

Как AMD выкрутилась и научилась воспроизводить операции ядер CUDA на своих видеокартах. История ROCm

Уровень сложностиСредний
Время на прочтение8 мин
Количество просмотров46K
Всего голосов 51: ↑51 и ↓0+71
Комментарии55

Комментарии 55

ЗакрепленныеЗакреплённые комментарии

Остаётся только запуск через компилятор DirectML.

DirectML очень медленный (до 10 раз медленнее rocm) и часто утекает память.

Перед запуском Python (или программы с PyTorch) необходимо установить переменную окружения HSA_OVERRIDE_GFX_VERSION=10.3.0 для корректной работы среды ROCm.

Это только для карточек rx 6600/6700. Для rx 570/580 или для встройки vega уже потребуется:
export HSA_OVERRIDE_GFX_VERSION=8.0.3

Для карточек rx 6800/6900 и всей линейки 7000 это не требуется, всё будет работать из коробки без доп. команд. И это только в линуксе, под виндой это делается по другому.

======

На Windows для ускорения на AMD видеокартах сейчас во всю используют ZLUDA, без сложностей с запуском и wsl.
ZLUDA транслирует вызовы CUDA в ROCm 5.7 для Windows, реализуется это подменой CUDA библиотек, так что софт думает, что работает с CUDA-карточкой.

Согласно табличке из статьи, для 6800+ это заведется само, а для rx 6600 и 6700 нужно взять скомпилированные сообществом ROCmLibs отсюда: https://github.com/brknsoul/ROCmLibs

Варианты использования разные:

  • Можно использовать с pytorch, заменяя библиотеки CUDA вручную по инструкции.

  • Запускать как zluda.exe blender.exe, работает далеко не со всем софтом.

  • Брать готовые решения, которые есть для популярный клиентов нейросетей.

SD:

LLM:

  • koboldcpp-rocm веб-клиент с возможностью запустить OpenAI совместимый сервер, поддерживает все карточки, включая rx 6600/6700 и младше, ничего ставить не нужно, всё запустится из коробки само.

  • ollama поддерживает ROCm на Windows для rx 6800+ карточек из коробки, поддержки 6600/6700 нет.

  • LM Studio, Msty, Jan - GUI клиенты, которые тоже поддерживают только rx 6800+ карточки.

По производительности это работает на уровне rtx карт, в отличии от DirectML тут паритет по скорости, 6600 плюс минус равна 3060, rx 7900 xtx процентов на 20% остает от rtx 4090 в kobold-rocm и так далее.
Для SD будет полезно включить оптимизатор --opt-sdp-no-mem-attention, чтобы не было просадок по памяти.

Ткну пальцем в небо.

AMD не интересно тратить кучу денег на портирование софта под неактуальную для них ось.

Основные потребители для ROCm это, как правило, серьёзные конторы с огромными кластерами, битком набитыми всяческими AMD Instinct со товарищи. Естетственно, там НЕ окошки в качестве оси :)

А через WSL, выходит, можно относительно малой кровью ROCm и на окошках запустить. А, раз так, то можно и сделать в качестве бонуса.

видимо потому что это не для запуска моделей, а для разработчиков, у которых по какой-то причине винда

Вот интересно, в соответствии с приведённой сравнительной таблицей с официального сайта у линейки RDNA3 есть поддержка технологии у самой младшей модели 7600 и у старших 7900, а что разве мидл сегмент не поддерживает это?

Статья интересная, спасибо. Хоть узнал вкратце что это

даже на 6500 xt рабоатет если говорить за работу на линукс. А вот на windows, по всей видимости, решили не тратиться на работу с младшими архитектурами, и реализовали поддержку только самых новых решений и c барского плеча дали поддержку на 6800 и выше.

Ну наконец-то! Очень интересно! А что с производительностью?

Тактично умалчивают.

Всё отлично в моём случае. Но опять же из-за того что невозможно натурально сравнить работу напрямую и через ROCm, то нельзя точно говорить о каких-либо цифрах, неизвестен процент потерь. Я догадываюсь что производительность около 80-95%, в той же stable diffusion моя RX 6600 выдает результаты скорости, относительно сопоставимые с RTX 3060.

Остаётся только запуск через компилятор DirectML.

DirectML очень медленный (до 10 раз медленнее rocm) и часто утекает память.

Перед запуском Python (или программы с PyTorch) необходимо установить переменную окружения HSA_OVERRIDE_GFX_VERSION=10.3.0 для корректной работы среды ROCm.

Это только для карточек rx 6600/6700. Для rx 570/580 или для встройки vega уже потребуется:
export HSA_OVERRIDE_GFX_VERSION=8.0.3

Для карточек rx 6800/6900 и всей линейки 7000 это не требуется, всё будет работать из коробки без доп. команд. И это только в линуксе, под виндой это делается по другому.

======

На Windows для ускорения на AMD видеокартах сейчас во всю используют ZLUDA, без сложностей с запуском и wsl.
ZLUDA транслирует вызовы CUDA в ROCm 5.7 для Windows, реализуется это подменой CUDA библиотек, так что софт думает, что работает с CUDA-карточкой.

Согласно табличке из статьи, для 6800+ это заведется само, а для rx 6600 и 6700 нужно взять скомпилированные сообществом ROCmLibs отсюда: https://github.com/brknsoul/ROCmLibs

Варианты использования разные:

  • Можно использовать с pytorch, заменяя библиотеки CUDA вручную по инструкции.

  • Запускать как zluda.exe blender.exe, работает далеко не со всем софтом.

  • Брать готовые решения, которые есть для популярный клиентов нейросетей.

SD:

LLM:

  • koboldcpp-rocm веб-клиент с возможностью запустить OpenAI совместимый сервер, поддерживает все карточки, включая rx 6600/6700 и младше, ничего ставить не нужно, всё запустится из коробки само.

  • ollama поддерживает ROCm на Windows для rx 6800+ карточек из коробки, поддержки 6600/6700 нет.

  • LM Studio, Msty, Jan - GUI клиенты, которые тоже поддерживают только rx 6800+ карточки.

По производительности это работает на уровне rtx карт, в отличии от DirectML тут паритет по скорости, 6600 плюс минус равна 3060, rx 7900 xtx процентов на 20% остает от rtx 4090 в kobold-rocm и так далее.
Для SD будет полезно включить оптимизатор --opt-sdp-no-mem-attention, чтобы не было просадок по памяти.

Спасибо! По всей видимости, я действительно зря не решился начать изучать ZLUDA.

Что будет лучше по скорости: 7900 XTX или 4070 Ti Super? Они примерно в одну цену.

Если речь идёт о нейронках, то 7900XTX, imho лучше. У неё банально больше видеопамяти (24 vs 12?), а недостаток VRAM - главная гиря на ногах при работе с большими моделями.

P.S. Дома Fedora и 7900XT (20Gb VRAM).

Проблем почти нет - в ComfyUI Flux-модели даже в FP16 режиме обрабатываются достаточно шустро (что-то около 2.2 секунды/итерацию). Единственное, пришлось немного поиграться с настройками.

Ну да, хочется покрутить как ллм так и модельки для создания картинок

И, да, AMDшные видюхи в пингвинах поддерживаются из коробки. То есть, никаких заморочек с kmod'ами при апдейте ядра.

1060, в своё время, у меня немало крови попила на этом поприще.

2.2 сек на итерацию для Flux в ComfyUI это твердый уровень 3060, причем 8 гигабайтной. Даже 4060 уже 1.2-1.4 секунды на шаг. Все что выше это уже меньше секунды. А rtx 3060 дешевле 7900 xt где-то раза в три?

А модель флюксовая сколько весит, если не секрет? А то флюкс, он разный бывает ;)

обычная fp8 flux .safetensors unet - 11Гб, bnb-nf4 у меня на 4060 практически с той же скоростью работает (но там лору не подцепить), а flux1-dev-Q4_1.gguf на 7 гиг выдает где-то 1.8-1.9 сек на итерацию.

А я говорил о flux1-dev FP16 22 гига весом. Плюс, иногда ещё и лору сверху цепляю. Разрешение картинки 1024x1024. Полностью в VRAM сия модель не влазит, поэтому скромные 2 секунды за итерацию с небольшим.

И я буду сильно удивлён, если 4060 выдаст при таких условиях цифры лучше тех же 2 секунд.

Так-то, я пробовал и GGUFы и NF4 - да быстрее, но и глючнее на порядок почему-то.

Flux fp16 замечательно работает и в comfy и в фордже на 16 гигах, но да - медленее и на 4060 где-то около 2 секунд и есть.. А если хоть одну лору задействовать, то будет и 2.5. Непонятно только в чем глубинный смысл запуска fp16 на чем-то ниже 3090 :)

Глубинный смысл в том, что качество при FP16 реально местами выше. Ну и часть кастомных моделей для флюкса есть только в FP16 формате. Возможно, их можно ужать до FP8/GGUF, но нафига, если оно и так работает с приемлемой скоростью :)

Дальнейшее развитие ROCm в 2017 году

Семья поддерживаемых линеек видеокарт пополнилась: Radeon RX Vega и Radeon Instinct. Впервые была достигнута бинарная совместимость кода CUDA с платформой.

Ну это же неправда.

Хотелось бы больше конкретики.

Ну, то есть, вы утверждаете, что я могу взять Вегу, накатить ROCm версии 3, а потом запустить на ней произвольный cubin? И для какого поколения? Pascal (как раз актуальное для того времени), Maxwell, или другого?

Ни у Nvidia, ни у AMD нет бинарной совместимости даже между поколениями своих видях, что уж говорить между разными производителями (на разве что у Ampere с Ada внезапно оставили, а ещё RDNA версий 1 и 2, и то если в бинаре подправить версию процессора).

На Nvidia и на AMD очень разные наборы инструкций, например, на вегах нет LOP3 и cross-lane инструкции абсолютно не такие.

Вега (как и все GCN) имеет 64 кб локальной памяти (aka shared memory в терминологии CUDA), в то время как у Nvidia (цитата из документации): Shared memory capacity per SM is 64KB for GP100 and 96KB for GP104. For comparison, Maxwell provided 96KB and up to 112KB of shared memory, respectively.

Как вы запустите тот же кернел, если у вас, например объёма локальной памяти не хватит?

Вы правы в том, что это не бинарная совместимость в строгом смысле слова - я использовал этот термин как упрощение. ROCm/HIP обеспечивает совместимость на уровне исходного кода через автоматическую конвертацию, а не прямое выполнение CUDA бинарников. При этом суть остаётся той же - разработчику в большинстве случаев не требуется вручную переписывать код, так как HIP автоматически handled большинство различий в архитектурах.

Что касается технических различий (LOP3, размеры shared memory и т.д.) - они действительно есть, и именно поэтому используется промежуточный слой трансляции, а не прямое исполнение бинарного кода. HIP/ROCm абстрагирует эти различия, позволяя достичь практически той же цели - запуска CUDA кода на AMD GPU с минимальными модификациями.

Вот бы уже все просто делали под DirectML, а не под вендерлок.

Но что-то даже сама Microsoft не особо спешит развивать DirectML и решения основанные на нём

Поменять один вендор лок на другой?

Так это софтовый менее проблемный вендерлок. Винду можно поставить куда угодно в отличии от железа нвидии

А из кросс платформенного есть Unity Barracuda, кстати. Но это привязываться к Unity

Ну вообще перспектива платить за лицензии винды чтобы использовать не куду а директмл (скорее всего на тех же картах от нвидии) - такое себе.

Ну вообще дома у многих железо от AMD и Intel, а поверх прям с завода винда. Не таскать же с собой на каждый чих eGPU с Nvidia.

Да и в интерпрайзе есть винда.

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

А ещё кудовские зависимости дико жирные и требуются уникальные версии под разные решения (то есть когда на машине у тебя стоит тенсорфлов, питорч и ещё что-то - для этого всего часто могут требоваться разные версии кудовских зависимостей)

Уже придумано, реализовано и доступно прямо сейчас, кроссплатформенное, открытое, универсальное решение работающее везде (если захотеть): на любой ОС, на телефонах, на расбери пи и так далее - Vulkan.
llama.cpp можно скомпилировать с поддержкой вулкана под любую платформу.
Только все всё равно предпочитают CoreML на маке, CUDA/ROCm на линуксе и винде.

Где уже реализована поддержка Vulkan:

SD:

  • SHARK-Studio - для генерации изображений, поддерживает Stable Diffusion модели. Этот стартап купила AMD, кстати.

  • koboldcpp - помимо запускай текстовый моделей, умеет запускать и модели SD, если надо, умеет работать в паре с SHARK-Studio, для генерации картинок во время генерации текста.

  • stable-diffusion.cpp - может запускать и Flux модели, в том числе и в формате квантованных gguf, снижая требования к видеопамяти.

LLM:

  • llama.cpp vulkan - просто скачать версию с vulkan для windows или скомпилировать для linux, или взять в AUR.

  • koboldcpp через флаг --usevulkan (не путать с koboldcpp-rocm, которая заточена на запуска на windows на любых amd gpu через rocm)

Плюс различные реализации нейросетей через ncnn-vulkan на гитхабе. У pytorch есть делегат для вулкана.

Можно даже запускать на разнородных видеокартах, я пробовал rtx 4090 + rx 6600, это работало и добавляло ускорения gguf моделям выгруженным частично на GPU, а частично на CPU. За счет дополнительных 8гб видеопамяти, в быструю vram влезало больше слоев модели.

По моим тестам, Vulkan был медленнее CUDA примерно на 10% в llama-bench. И в Vulkan пока не реализован полноценно flash attention, а из-за этого не работает квантование кэша, что не позволяет выиграть пару гб vram, но это вопрос популярности или времени.

Круто. Поскорее бы уже все на это слезли с иглы нвидии и по умолчанию это использовали.

А NPU умеет использовать?

А на игровых консолях кто-то пробовал реализовать через vulkan? DirectML(Xbox) и Unity Barracuda (Xbox, PS, Switch) там работают.

А в браузере через webgl или webGPU как поддержку vulkan реализовать? Unity Barracuda просто на пиксельных шейдерах в webgl крутит нейронку. Медленнее конечно, чем на вычислительных, но в webgl же и нету вычислительных насколько помню.

А на серверах, где нвидевские видюхи по умолчанию не прокинуты (и у меня было времени разобраться возможно ли вообще их так прокинуть) как графическое устройство - будет работать?

ROCm и zluda просто чудо на rx 470 ,получилось заставить блендер работать.

Ну откуда эта мода на выделение приблизительно случайных слов в тексте болдом, да ещё и с использованием такого вырвиглазного цвета? Это помогает в понимании материала ровно никак, а в большинстве случаев раздражает и/или мешает.
Люди, которые умеют читать больше одного абзаца без потери концентрации, замечательно поймут и без этой цветомузыки. И даже в структурирование смогут, без подсвечивания подзаголовков. А те, кто не в состоянии, те под кат и так не пойдут.

А как насчёт запуска PhysX NVidia на AMD Radeon?

PhysX по моему и так на AMD работает.
С другой стороны, а так ли нужен этот PhysX в современном мире?.. Большая части вещей из ролика легко реализуема и без PhysX, либо на средствах CPU, либо средствами GPU через шейдеры и вычислительные шейдеры

Не, CPU этот PhysX не тянет, а GPU - да: через GPU и у самой NVidia этот PhysX работает, и PhysX у NVidia - это просто ещё один специализированный интерфейс к GPU на подобие CUDA.

А зачем тогда в панели нвидии модно было переключить чтобы physx использовал только цпу..

На CPU PhysX - тормозит. Это как софтверный режим рендеринга без использования GPU.

Я что то не понял, о какой поддержке rocm для windows при использовании Stable Diffusion в статье речь, если в разделе start locally на сайте PyTorch чётко видно, что "NOTE: ROCm is not available on Windows". Как я находил в интернете, проблема в том что под виндой оно до сих пор не поддерживает MiOpen. Вот когда это реализуют, тогда и поговорим. А пока для запуска Stable Diffusion/Flux/Pony приходится грузиться в Linux. И да, по этому ноги моей в AMD больше не будет. Их карты у малого числа пользователей, и решения для ИИ под эти карты разрабатывать никому не интересно...

Куча ультрабуков с AMD и Intel в качестве GPU. Все домашние консоли с AMD.

Все портативные консоли (x86_64 игровые планшеты) с AMD.

И все они начали ещё и получать NPU.

Возможно благодаря NPU уже наконец-то забудут про CUDA

За потенциально глупый коммент - аймсори

Для чайников - я просто могу взять свою залежавшуюся rx-580 и под линуксами попользоваться разным локальным ИИ инструментарием, который по идее работает с ядрами CUDA?

Я помню пробовал в SD побаловаться на ubuntu, но запустилось соответственно только под CPU (8 ядер 16 потоков) и под какие-то там средние настройки, генерацию одной итерации нужно ждать минуты 3. Долговато. Что можно ожидать от rx-580 карточки в среднем?

Дебют на Windows, 2023-2024 год ... взрывной рост интереса к AI

Только вот дебют AMD HIP API for Windows был ещё в 2021 - https://gpuopen.com/blender-cycles-amd-gpu/. В том году именно коллеги из Blender целенаправленно вычистили все баги в стеке. А вот крипто-ии-бро в портировании появившихся на основе HIP ROCm-библиотек не преуспели (де-факто никто не готов делать сборки для amd: automatic1111 - nvidia only, comfyui - nvidia only и т. д.). Факторы, которые мешают портированию (не считая отсутствие рук):

1) подвязанность рантайм-библиотек на формат ELF (например, https://github.com/ROCm/ROCR-Runtime/blob/rocm-6.2.0/src/core/inc/amd_elf_image.hpp).

2) зоопарк в имплементациях: rocblas - hipblaslt - miopen - rocfft - у каждой из этих библиотек свой способ сборки (кто-то в elf, кто-то в fatbinary-elf, кто-то байткод в sqlite, кто-то JIT)

3) понимание того, что собранный проект будет занимать порядка 3ГБ (см. размер torch-rocm для linux)

Видел вот такую новость

>С выходом CUDA 11.6 компания NVIDIA обновила лицензионное соглашение, предупредив разработчиков о недопустимости применять эту технологию на графических процессорах сторонних разработчиков с использованием «уровней трансляции», преобразующих один код в другой. Производитель подчеркнул, что методы обратного проектирования, декомпиляции или дизассемблирования SDK с его переводом на платформы, отличные от NVIDIA, запрещены.

Это касается описанного в статье или нет?

так никто и не использует "cuda" на radeon GPU, там используется ROCm.

Прямо в статье написано что это платформа для обеспечения совместимости с CUDA

Есть код, написанный под cuda. Написан код не сотрудниками nvidia. Его транслируют в совместимый с радеонами код. Запускают этот код на радеонах. О каком использовании cuda речь?

>методы обратного проектирования, декомпиляции или дизассемблирования SDK с его переводом на платформы, отличные от NVIDIA, запрещены

Из этого следует, что трансляция в совместимый с Радеонами код запрещена лицензионным соглашением CUDA.

Если я чего-то не понимаю и технически это под данный запрет не попадает - то интересно почему именно и что именно.

Есть выражение add(a,b) - вам не нужно декомпелировать исходники, чтобы посчитать результат на rocm.

А где реверс пригодился бы, если вы хотите ускорить add(a,b) для огромных матриц, но.... не знаете как сделать...и лезете подсматривать к нвидеа.

для Куда есть свой язык(?) с оптимизациями. Вот его наверное не получится портировать, а мейнстримные либы которые и на cpu работают - только подсунуть нужную обёртку.

Если я правильно понял из статьи - первый rocm получился именно реверс инженирингом(так быстрее, но в среднем - тормозить будет), а потом уже стали(?) норм разрабатывать и оптимизировать под амд.

Вообще странно, почему, когда игра на триллион долларов на кону, энтузиасты еще где-то должны смотреть таблицы, выискивать крупицы.

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

Так AMD наверняка и тратит миллионы. Только не на то, чтобы окучивать тебя, а на то, чтобы окучивать компании, которые набивают серверные железом на миллиарды долларов.

Вот не верю, что оно так работает здесь. Кто бы из мегакорпораций железо не покупал, и кому б в итоге не сдавал в аренду, но там будет крутится код, написанный инженерами, которые хотят решать интересные задачи и иметь прогресс в своей работе, а не сутки напролет писать костыли для обхода проблем совместимости. Да, можно попробовать сломать инженеров через колено, поставив перед фактом, что они отныне будут писать код для железа с сомнительной совместимостью с передовыми библиотеками и постоянным ожиданием странных багов лишь потому, что CFO заключил удачную сделку, и, да, не все так страшно, ибо им всегда поможет 24х7 premium enterprise support... Но, кажется, что технологические прорывы делаются как-то не так

А как обстоят дела с Tensor Cores ? Тоже код с их применением портируется без просадок? Стоит ли пробовать решения AMD для сверточных нейросетей?

В прошлом году каким-то чудом смог заставить свою RX 6600 работать с pytorch на arch linux, через какое-то время снова стал пытаться запускать старый код, а уже не работает... Не помогает ни HSA_OVERRIDE_GFX_VERSION=10.3.0, ни переустановка rocm и торча. На винде HIP SDK, как отмечено в статье, даже не поддерживатеся. Я пробовал ставить zluda с пропатченным ROCm, и оно даже запускается, но поддерживаются только самые простые операции, по типу сложения и умножения тензоров, а уже RNN блоки крашатся на torch._cudnn_rnn_flatten_weight с CUDNN_STATUS_INTERNAL_ERROR.

Короче, жаль, что я 3 года назад выбрал не ту сторону и купил карту от AMD и столько с ней настрадался (причем не только с глубоким обучением, но и с драйверами для видеоигр), веры в их продукцию у меня уже нет.

Зарегистрируйтесь на Хабре, чтобы оставить комментарий