Фрагментный API WMMA
Загрузка, mma_sync и сохранение фрагментов
«Фрагментный API WMMA» — бесплатный урок CUDA Academy на CoddyKit. Это урок 3 из 4. Ты можешь прочитать весь урок бесплатно ниже — а потом практиковать его прямо в браузере с встроенным редактором кода и ИИ-репетитором 24/7. Это часть пути обучения CUDA Academy, и твой прогресс синхронизируется между веб-версией и приложением CoddyKit. Курс CUDA Academy содержит 4 уроков всего.
API WMMA
Для прямого программирования Tensor Core используется WMMA — API матричного умножения с накоплением на уровне варпа в пространстве имён nvcuda::wmma. 🧩
Весь варп работает сообща
WMMA работает на уровне всего варпа: все 32 потока в варпе совместно обрабатывают один фрагмент. Вы мыслите фрагментами, а не отдельными потоками.
Знакомьтесь: фрагмент
Фрагмент — это тип данных WMMA, содержащий часть фрагмента матрицы, принадлежащую одному варпу. Каждый поток незаметно владеет несколькими его элементами.
Три роли фрагментов
Вы объявляете фрагменты с метками matrix_a, matrix_b или накопитель. Метка сообщает WMMA, как этот фрагмент участвует в умножении с накоплением.
Размеры фрагментов фиксированы
Фрагменты используют заданные размеры, например 16 на 16 на 16. Вы выбираете поддерживаемую форму; аппаратное обеспечение принимает только определённые сочетания.
Шаг 1: загрузка
Сначала вызовите load_matrix_sync, чтобы загрузить фрагмент из памяти во фрагмент WMMA. Эта загрузка автоматически распределяет данные между потоками варпа.
wmma::load_matrix_sync(a_frag, ptr, ldm);Шаг 2: очистка накопителя
Перед сложением обнулите фрагмент результата с помощью fill_fragment. Чистый накопитель означает, что суммы начинаются с известного значения.
wmma::fill_fragment(c_frag, 0.0f);Шаг 3: умножение с накоплением
Основной вызов — mma_sync. Он выполняет D = A умножить на B плюс C для загруженных фрагментов, напрямую используя Tensor Core.
wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);Шаг 4: сохранение
Наконец, store_matrix_sync записывает фрагмент накопителя обратно в память. Это сохранение собирает части всех потоков в один фрагмент.
wmma::store_matrix_sync(out, c_frag, ldm, layout);Синхронизация означает синхронную работу варпа
Суффикс _sync напоминает: каждый вызов WMMA является синхронным для варпа. Все 32 линии должны достичь его одновременно, иначе поведение не определено.
Компоновка и ведущая размерность
Для загрузки и сохранения нужны ведущая размерность, то есть шаг между строками, а также построчная или постолбцовая компоновка для правильного чтения памяти.
Быстрая проверка
Какой вызов WMMA фактически выполняет матричное умножение с накоплением на Tensor Core?
Итоги
Вы прошли весь процесс WMMA: объявить фрагмент, загрузить фрагменты, очистить накопитель, вызвать mma_sync, а затем сохранить результат. Каждый шаг синхронен для варпа. 🙌
Часто задаваемые вопросы
Урок «Фрагментный API WMMA» бесплатный?
Да — полный текст урока «Фрагментный API WMMA» бесплатно доступен здесь в веб-версии. Чтобы практиковать его интерактивно (встроенный редактор кода и ИИ-репетитор 24/7) и разблокировать остальной курс CUDA Academy, подпишись на CoddyKit PRO. Курс CUDA Academy содержит 4 уроков всего.
Чему я научусь в уроке «Фрагментный API WMMA»?
Загрузка, mma_sync и сохранение фрагментов Ты практикуешь CUDA Academy с помощью реального кода, который запускаешь прямо в браузере, и ИИ-репетитор 24/7 отвечает на твои вопросы во время урока.
Нужен ли мне опыт, чтобы начать CUDA Academy?
Предыдущий опыт не требуется. CUDA Academy на CoddyKit структурирован для всех уровней — от новичков до продвинутых, поэтому ты можешь начать отсюда или с самого начала и учиться в своем темпе. Это урок 3 из 4.
Сколько времени занимает урок «Фрагментный API WMMA»?
Большинство уроков CoddyKit занимают около 5–10 минут. Каждый из них компактный и интерактивный, поэтому ты постоянно делаешь прогресс и продолжаешь с того же места в веб-версии и приложении.
Можно ли писать и запускать код в этом уроке CUDA Academy?
Да. Каждый урок CUDA Academy включает встроенный редактор кода, поэтому ты пишешь и запускаешь реальный код прямо в браузере и получаешь моментальную обратную связь от AI — локальная установка не требуется.
Все уроки этого курса
- Что вычисляют тензорные ядра
- Смешанная точность: FP16, BF16, TF32
- Фрагментный API WMMA
- Компромиссы численной устойчивости