Этот документ является продолжением усовершенствований WebAssembly и WebGPU для более быстрого веб-ИИ, часть 1 . Прежде чем продолжить, мы рекомендуем вам прочитать этот пост или посмотреть выступление на IO 24 .
Веб-графический процессор
WebGPU предоставляет веб-приложениям доступ к оборудованию графического процессора клиента для выполнения эффективных высокопараллельных вычислений. С момента запуска WebGPU в Chrome мы видели в сети невероятные демонстрации искусственного интеллекта (ИИ) и машинного обучения (МО).
Например, Web Stable Diffusion продемонстрировала, что можно использовать ИИ для создания изображений из текста прямо в браузере. Ранее в этом году собственная команда Google Mediapipe опубликовала экспериментальную поддержку вывода больших языковых моделей .
На следующей анимации показана Gemma , модель большого языка Google (LLM) с открытым исходным кодом, работающая полностью на устройстве в Chrome в режиме реального времени.
Следующая демонстрация Hugging Face модели Segment Anything Model компании Meta создает высококачественные маски объектов полностью на клиенте.
Это всего лишь пара замечательных проектов, демонстрирующих возможности WebGPU для искусственного интеллекта и машинного обучения. WebGPU позволяет этим и другим моделям работать значительно быстрее, чем на процессоре.
Тест Hugging Face WebGPU для встраивания текста демонстрирует огромное ускорение по сравнению с реализацией той же модели на ЦП. На ноутбуке Apple M1 Max WebGPU работал более чем в 30 раз быстрее. Другие сообщают, что WebGPU ускоряет тест более чем в 120 раз .
Улучшение функций WebGPU для искусственного интеллекта и машинного обучения
WebGPU отлично подходит для моделей искусственного интеллекта и машинного обучения, которые могут иметь миллиарды параметров благодаря поддержке вычислительных шейдеров . Вычислительные шейдеры выполняются на графическом процессоре и помогают выполнять параллельные операции с массивами над большими объемами данных.
Среди многочисленных улучшений WebGPU за последний год мы продолжили добавлять новые возможности для повышения производительности машинного обучения и искусственного интеллекта в Интернете. Недавно мы запустили две новые функции: 16-битные числа с плавающей запятой и упакованные целочисленные скалярные произведения.
16-битная с плавающей запятой
Помните, что рабочие нагрузки машинного обучения не требуют точности . shader-f16 — это функция, которая позволяет использовать тип f16 в языке шейдеров WebGPU. Этот тип с плавающей запятой занимает 16 бит вместо обычных 32 бит. f16 имеет меньший диапазон и менее точен, но для многих моделей ML этого достаточно.
Эта функция повышает эффективность несколькими способами:
Уменьшение памяти : тензоры с элементами f16 занимают половину места, что сокращает использование памяти вдвое. Вычисления на графическом процессоре часто ограничивают пропускную способность памяти, поэтому использование половины памяти часто означает, что шейдеры работают в два раза быстрее. Технически вам не нужна клавиша f16 для экономии пропускной способности памяти. Можно хранить данные в формате низкой точности, а затем расширять их до полного f32 в шейдере для вычислений. Но графический процессор тратит дополнительную вычислительную мощность на упаковку и распаковку данных.
Сокращение преобразования данных : f16 использует меньше вычислений за счет минимизации преобразования данных. Данные низкой точности можно сохранять и затем использовать напрямую, без преобразования.
Повышенный параллелизм . Современные графические процессоры способны одновременно помещать больше значений в исполнительные блоки графического процессора, что позволяет ему выполнять большее количество параллельных вычислений. Например, графический процессор, поддерживающий до 5 триллионов операций с плавающей запятой f32 в секунду, может поддерживать 10 триллионов операций с плавающей запятой f16 в секунду.
shader-f16 тест Hugging Face WebGPU для встраивания текста выполняет тест в 3 раза быстрее, чем f32 на ноутбуке Apple M1 Max.WebLLM — это проект, который может запускать несколько больших языковых моделей. Он использует Apache TVM , среду компилятора машинного обучения с открытым исходным кодом.
Я попросил WebLLM спланировать поездку в Париж, используя модель Llama 3 с восемью миллиардами параметров. Результаты показывают, что на этапе предварительного заполнения модели f16 работает в 2,1 раза быстрее, чем f32. На этапе декодирования скорость более чем в 1,3 раза выше.
Приложения должны сначала подтвердить, что адаптер графического процессора поддерживает f16, и, если он доступен, явно включить его при запросе устройства графического процессора. Если f16 не поддерживается, вы не можете запросить его в массиве requiredFeatures .
// main.js
const adapter = await navigator.gpu.requestAdapter();
const supportsF16 = adapter.features.has('shader-f16');
if (supportsF16) {
// Use f16.
const device = await adapter.requestDevice({
requiredFeatures: ['shader-f16'],
});
initApp(device);
}
Затем в шейдерах WebGPU вы должны явно включить f16 вверху. После этого вы можете использовать его в шейдере, как и любой другой тип данных с плавающей запятой.
// my-shader.wgsl
enable f16;
struct Data {
values : array<vec4<f16>>
}
@group(0) @binding(0) var<storage, read> data : Data;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid : vec3u) {
let value : vec4<f16> = data.values[gid.x];
...
}
Упакованные целочисленные скалярные произведения
Многие модели по-прежнему хорошо работают с точностью всего 8 бит (половина f16). Это популярно среди LLM и моделей изображений для сегментации и распознавания объектов. Тем не менее, качество вывода моделей ухудшается с меньшей точностью, поэтому 8-битное квантование подходит не для каждого приложения.
Относительно немногие графические процессоры изначально поддерживают 8-битные значения. Именно здесь на помощь приходят упакованные целочисленные скалярные произведения. Мы включили DP4a в Chrome 123 .
Современные графические процессоры имеют специальные инструкции для того, чтобы взять два 32-битных целых числа, интерпретировать каждое из них как 4 последовательно упакованных 8-битных целых числа и вычислить скалярное произведение между их компонентами.
Это особенно полезно для искусственного интеллекта и машинного обучения, поскольку ядра матричного умножения состоят из множества скалярных произведений.
Например, давайте умножим матрицу 4 x 8 на вектор 8 x 1. Для этого необходимо взять 4 скалярных произведения для вычисления каждого значения выходного вектора; А, Б, С и Д.
Процесс вычисления каждого из этих выходных данных одинаков; мы рассмотрим шаги, необходимые для вычисления одного из них. Перед любыми вычислениями нам сначала нужно преобразовать 8-битные целочисленные данные в тип, с которым мы можем выполнять арифметические действия, например f16. Затем мы выполняем поэлементное умножение и, наконец, складываем все произведения вместе. В общей сложности для всего умножения матрицы на вектор мы выполняем 40 преобразований целых чисел в числа с плавающей запятой для распаковки данных, 32 умножения с плавающей запятой и 28 операций сложения с плавающей запятой.
Для более крупных матриц с большим количеством операций упакованные целочисленные скалярные произведения могут помочь уменьшить объем работы.
Для каждого из выходных данных вектора результатов мы выполняем две операции упакованного скалярного произведения, используя встроенный язык шейдинга WebGPU dot4U8Packed , а затем суммируем результаты. Итого, для всего умножения матрицы на вектор мы не производим никакого преобразования данных. Мы выполняем 8 упакованных скалярных произведений и 4 сложения целых чисел.
Мы протестировали упакованные целочисленные скалярные произведения с 8-битными данными на различных потребительских графических процессорах. По сравнению с 16-битной плавающей запятой мы видим, что 8-битная операция работает в 1,6–2,8 раза быстрее. Когда мы дополнительно используем упакованные целочисленные скалярные произведения, производительность становится еще лучше. Это в 1,7–2,9 раза быстрее.
Проверьте поддержку браузера с помощью свойства wgslLanguageFeatures . Если графический процессор изначально не поддерживает упакованные скалярные произведения, браузер полифилизирует свою собственную реализацию.
// main.js
if (navigator.gpu.wgslLanguageFeatures.has('packed_4x8_integer_dot_product')) {
// Use dot4U8Packed, dot4I8Packed builtin
// functions in the shaders.
}
Следующий фрагмент кода diff (разница), подчеркивающий изменения, необходимые для использования упакованных целочисленных продуктов в шейдере WebGPU.
Раньше — шейдер WebGPU, который накапливает частичные скалярные произведения в переменную sum. В конце цикла сумма сохраняет полное скалярное произведение между вектором и одной строкой входной матрицы.
// my-dot-product.wgsl @compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid : vec3u) { var sum : f16; let start = gid.x * uniforms.dim; for (var i = 0u; i < uniforms.dim; i++) { let v1 : vec4<f16> = vector.values[i]; let v2 : vec4<f16> = matrix.values[start + i]; sum += dot(v1, v2); } }
After — шейдер WebGPU, написанный для использования скалярных произведений упакованных целых чисел. Основное отличие состоит в том, что вместо загрузки 4 значений с плавающей запятой из вектора и матрицы этот шейдер загружает одно 32-битное целое число. Это 32-битное целое число содержит данные четырех 8-битных целочисленных значений. Затем мы вызываем dot4U8Packed для вычисления скалярного произведения двух значений.
// my-dot-product.wgsl
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid : vec3u) {
var sum : f32;
let start = gid.x * uniforms.dim;
for (var i = 0u; i < uniforms.dim; i++) {
let v1 : u32 = vector.values[i];
let v2 : u32 = matrix.values[start + i];
sum += dot4U8Packed(v1, v2);
}
}
Как 16-битные скалярные произведения с плавающей запятой, так и упакованные целочисленные скалярные произведения — это встроенные функции Chrome, которые ускоряют работу искусственного интеллекта и машинного обучения. 16-битные числа с плавающей запятой доступны, если аппаратное обеспечение их поддерживает, а Chrome реализует упакованные целочисленные скалярные произведения на всех устройствах.
Вы можете использовать эти функции в Chrome Stable уже сегодня, чтобы повысить производительность.
Предлагаемые функции
Забегая вперед, мы исследуем еще две функции: подгруппы и кооперативное матричное умножение.
Функция подгрупп обеспечивает параллелизм на уровне SIMD для взаимодействия или выполнения коллективных математических операций, таких как суммирование более 16 чисел. Это обеспечивает эффективный обмен данными между потоками. Подгруппы поддерживаются современными API-интерфейсами графических процессоров с разными именами и в несколько разных формах.
Мы превратили общий набор в предложение, которое подали в группу стандартизации WebGPU. Мы создали прототипы подгрупп в Chrome под экспериментальным флагом и представили на обсуждение наши первоначальные результаты. Основной вопрос заключается в том, как обеспечить переносимость.
Кооперативное матричное умножение — это более позднее дополнение к графическим процессорам. Большое матричное умножение можно разбить на несколько меньших матричных умножений. Кооперативное матричное умножение выполняет умножение этих меньших блоков фиксированного размера за один логический шаг. На этом этапе группа потоков эффективно взаимодействует для вычисления результата.
Мы изучили поддержку базовых API-интерфейсов графических процессоров и планируем представить предложение группе стандартизации WebGPU. Как и в случае с подгруппами, мы ожидаем, что большая часть обсуждений будет сосредоточена вокруг переносимости.
Чтобы оценить производительность операций с подгруппами, в реальном приложении мы интегрировали экспериментальную поддержку подгрупп в MediaPipe и протестировали ее с прототипом Chrome для операций с подгруппами.
Мы использовали подгруппы в ядрах графического процессора на этапе предварительного заполнения большой языковой модели, поэтому я сообщаю об ускорении только на этапе предварительного заполнения. На графическом процессоре Intel мы видим, что подгруппы работают в два с половиной раза быстрее, чем базовый уровень. Однако эти улучшения не одинаковы для разных графических процессоров.
На следующей диаграмме показаны результаты применения подгрупп для оптимизации микротеста умножения матриц на нескольких потребительских графических процессорах. Умножение матриц — одна из самых сложных операций в больших языковых моделях. Данные показывают, что на многих графических процессорах подгруппы увеличивают скорость в два, пять и даже тринадцать раз по сравнению с базовым уровнем. Однако обратите внимание, что на первом графическом процессоре подгруппы ничуть не лучше.
Оптимизация графического процессора сложна
В конечном счете, лучший способ оптимизировать ваш графический процессор зависит от того, какой графический процессор предлагает клиент. Использование необычных новых функций графического процессора не всегда окупается так, как вы ожидаете, поскольку здесь может быть задействовано множество сложных факторов. Лучшая стратегия оптимизации для одного графического процессора может не быть лучшей стратегией для другого графического процессора.
Вы хотите минимизировать пропускную способность памяти, полностью используя вычислительные потоки графического процессора.
Шаблоны доступа к памяти также могут быть очень важны. Графические процессоры, как правило, работают намного лучше, когда вычислительные потоки обращаются к памяти по схеме, оптимальной для оборудования. Важно: на разном оборудовании графического процессора следует ожидать разных характеристик производительности. Возможно, вам придется запустить различные оптимизации в зависимости от графического процессора.
На следующей диаграмме мы взяли тот же алгоритм умножения матриц, но добавили еще одно измерение, чтобы дополнительно продемонстрировать влияние различных стратегий оптимизации, а также сложность и различия между разными графическими процессорами. Мы представили здесь новую технику, которую назовем «Swizzle». Swizzle оптимизирует шаблоны доступа к памяти, чтобы сделать их более оптимальными для оборудования.
Вы можете видеть, что изменение памяти имеет значительное влияние; иногда это даже более эффективно, чем подгруппы. На графическом процессоре 6 swizzle обеспечивает ускорение в 12 раз, а подгруппы — в 13 раз. В совокупности они дают невероятное ускорение в 26 раз. Для других графических процессоров иногда объединение swizzle и подгрупп работает лучше, чем каждый из них по отдельности. А на других графических процессорах лучше всего работает использование swizzle.
Настройка и оптимизация алгоритмов графического процессора для хорошей работы на любом оборудовании может потребовать большого опыта. Но, к счастью, ведется огромная талантливая работа над инфраструктурами библиотек более высокого уровня, такими как Mediapipe , Transformers.js , Apache TVM , ONNX Runtime Web и другие.
Библиотеки и платформы имеют хорошие возможности для решения сложных задач управления различными архитектурами графических процессоров и создания кода для конкретной платформы, который будет хорошо работать на клиенте.
Вынос
Команда Chrome продолжает способствовать развитию стандартов WebAssembly и WebGPU, чтобы улучшить веб-платформу для рабочих нагрузок машинного обучения. Мы инвестируем в более быстрые вычислительные примитивы, лучшее взаимодействие веб-стандартов и обеспечиваем эффективную работу как больших, так и малых моделей на разных устройствах.
Наша цель — максимально расширить возможности платформы, сохранив при этом лучшее от Интернета: охват, удобство использования и мобильность. И мы делаем это не в одиночку. Мы работаем в сотрудничестве с другими поставщиками браузеров в W3C и многими партнерами по разработке.
Мы надеемся, что при работе с WebAssembly и WebGPU вы помните следующее:
- Выводы ИИ теперь доступны в Интернете на всех устройствах. Это дает преимущества работы на клиентских устройствах, такие как снижение затрат на сервер, низкая задержка и повышение конфиденциальности.
- Хотя многие обсуждаемые функции актуальны в первую очередь для авторов платформы, ваши приложения могут получить от этого выгоду без особых затрат.
- Веб-стандарты изменчивы и развиваются, и мы всегда ищем обратную связь. Поделитесь своими для WebAssembly и WebGPU .
Благодарности
Мы хотели бы поблагодарить команду Intel по веб-графике, которая сыграла важную роль в разработке WebGPU f16 и функций целочисленного скалярного произведения. Мы хотели бы поблагодарить других членов рабочих групп WebAssembly и WebGPU в W3C, включая других поставщиков браузеров.
Спасибо командам искусственного интеллекта и машинного обучения как в Google, так и в сообществе разработчиков программного обеспечения с открытым исходным кодом за то, что вы стали замечательными партнерами. И, конечно же, всех наших товарищей по команде, которые делают все это возможным.
,Этот документ является продолжением усовершенствований WebAssembly и WebGPU для более быстрого веб-ИИ, часть 1 . Прежде чем продолжить, мы рекомендуем вам прочитать этот пост или посмотреть выступление на IO 24 .
Веб-графический процессор
WebGPU предоставляет веб-приложениям доступ к оборудованию графического процессора клиента для выполнения эффективных высокопараллельных вычислений. С момента запуска WebGPU в Chrome мы видели в сети невероятные демонстрации искусственного интеллекта (ИИ) и машинного обучения (МО).
Например, Web Stable Diffusion продемонстрировала, что можно использовать ИИ для создания изображений из текста прямо в браузере. Ранее в этом году собственная команда Google Mediapipe опубликовала экспериментальную поддержку вывода больших языковых моделей .
На следующей анимации показана Gemma , модель большого языка Google (LLM) с открытым исходным кодом, работающая полностью на устройстве в Chrome в режиме реального времени.
Следующая демонстрация Hugging Face модели Segment Anything Model компании Meta создает высококачественные маски объектов полностью на клиенте.
Это всего лишь пара замечательных проектов, демонстрирующих возможности WebGPU для искусственного интеллекта и машинного обучения. WebGPU позволяет этим и другим моделям работать значительно быстрее, чем на процессоре.
Тест Hugging Face WebGPU для встраивания текста демонстрирует огромное ускорение по сравнению с реализацией той же модели на ЦП. На ноутбуке Apple M1 Max WebGPU работал более чем в 30 раз быстрее. Другие сообщают, что WebGPU ускоряет тест более чем в 120 раз .
Улучшение функций WebGPU для искусственного интеллекта и машинного обучения
WebGPU отлично подходит для моделей искусственного интеллекта и машинного обучения, которые могут иметь миллиарды параметров благодаря поддержке вычислительных шейдеров . Вычислительные шейдеры выполняются на графическом процессоре и помогают выполнять параллельные операции с массивами над большими объемами данных.
Среди многочисленных улучшений WebGPU за последний год мы продолжили добавлять новые возможности для повышения производительности машинного обучения и искусственного интеллекта в Интернете. Недавно мы запустили две новые функции: 16-битные числа с плавающей запятой и упакованные целочисленные скалярные произведения.
16-битная с плавающей запятой
Помните, что рабочие нагрузки машинного обучения не требуют точности . shader-f16 — это функция, которая позволяет использовать тип f16 в языке шейдеров WebGPU. Этот тип с плавающей запятой занимает 16 бит вместо обычных 32 бит. f16 имеет меньший диапазон и менее точен, но для многих моделей ML этого достаточно.
Эта функция повышает эффективность несколькими способами:
Уменьшение памяти : тензоры с элементами f16 занимают половину места, что сокращает использование памяти вдвое. Вычисления на графическом процессоре часто ограничивают пропускную способность памяти, поэтому использование половины памяти часто означает, что шейдеры работают в два раза быстрее. Технически вам не нужна клавиша f16 для экономии пропускной способности памяти. Можно хранить данные в формате низкой точности, а затем расширять их до полного f32 в шейдере для вычислений. Но графический процессор тратит дополнительную вычислительную мощность на упаковку и распаковку данных.
Сокращение преобразования данных : f16 использует меньше вычислений за счет минимизации преобразования данных. Данные низкой точности можно сохранять и затем использовать напрямую, без преобразования.
Повышенный параллелизм . Современные графические процессоры способны одновременно помещать больше значений в исполнительные блоки графического процессора, что позволяет ему выполнять большее количество параллельных вычислений. Например, графический процессор, поддерживающий до 5 триллионов операций с плавающей запятой f32 в секунду, может поддерживать 10 триллионов операций с плавающей запятой f16 в секунду.
shader-f16 тест Hugging Face WebGPU для встраивания текста выполняет тест в 3 раза быстрее, чем f32 на ноутбуке Apple M1 Max.WebLLM — это проект, который может запускать несколько больших языковых моделей. Он использует Apache TVM , среду компилятора машинного обучения с открытым исходным кодом.
Я попросил WebLLM спланировать поездку в Париж, используя модель Llama 3 с восемью миллиардами параметров. Результаты показывают, что на этапе предварительного заполнения модели f16 работает в 2,1 раза быстрее, чем f32. На этапе декодирования скорость более чем в 1,3 раза выше.
Приложения должны сначала подтвердить, что адаптер графического процессора поддерживает f16, и, если он доступен, явно включить его при запросе устройства графического процессора. Если f16 не поддерживается, вы не можете запросить его в массиве requiredFeatures .
// main.js
const adapter = await navigator.gpu.requestAdapter();
const supportsF16 = adapter.features.has('shader-f16');
if (supportsF16) {
// Use f16.
const device = await adapter.requestDevice({
requiredFeatures: ['shader-f16'],
});
initApp(device);
}
Затем в шейдерах WebGPU вы должны явно включить f16 вверху. После этого вы можете использовать его в шейдере, как и любой другой тип данных с плавающей запятой.
// my-shader.wgsl
enable f16;
struct Data {
values : array<vec4<f16>>
}
@group(0) @binding(0) var<storage, read> data : Data;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid : vec3u) {
let value : vec4<f16> = data.values[gid.x];
...
}
Упакованные целочисленные скалярные произведения
Многие модели по-прежнему хорошо работают с точностью всего 8 бит (половина f16). Это популярно среди LLM и моделей изображений для сегментации и распознавания объектов. Тем не менее, качество вывода моделей ухудшается с меньшей точностью, поэтому 8-битное квантование подходит не для каждого приложения.
Относительно немногие графические процессоры изначально поддерживают 8-битные значения. Именно здесь на помощь приходят упакованные целочисленные скалярные произведения. Мы включили DP4a в Chrome 123 .
Современные графические процессоры имеют специальные инструкции для того, чтобы взять два 32-битных целых числа, интерпретировать каждое из них как 4 последовательно упакованных 8-битных целых числа и вычислить скалярное произведение между их компонентами.
Это особенно полезно для искусственного интеллекта и машинного обучения, поскольку ядра матричного умножения состоят из множества скалярных произведений.
Например, давайте умножим матрицу 4 x 8 на вектор 8 x 1. Для этого необходимо взять 4 скалярных произведения для вычисления каждого значения выходного вектора; А, Б, С и Д.
Процесс вычисления каждого из этих выходных данных одинаков; мы рассмотрим шаги, необходимые для вычисления одного из них. Перед любыми вычислениями нам сначала нужно преобразовать 8-битные целочисленные данные в тип, с которым мы можем выполнять арифметические действия, например f16. Затем мы выполняем поэлементное умножение и, наконец, складываем все произведения вместе. В общей сложности для всего умножения матрицы на вектор мы выполняем 40 преобразований целых чисел в числа с плавающей запятой для распаковки данных, 32 умножения с плавающей запятой и 28 операций сложения с плавающей запятой.
Для более крупных матриц с большим количеством операций упакованные целочисленные скалярные произведения могут помочь уменьшить объем работы.
Для каждого из выходных данных в векторе результатов мы выполняем две операции упакованного скалярного произведения, используя встроенный язык шейдинга WebGPU dot4U8Packed , а затем суммируем результаты. Итого, для всего умножения матрицы на вектор мы не производим никакого преобразования данных. Мы выполняем 8 упакованных скалярных произведений и 4 сложения целых чисел.
Мы протестировали упакованные целочисленные скалярные произведения с 8-битными данными на различных потребительских графических процессорах. По сравнению с 16-битной плавающей запятой мы видим, что 8-битная операция работает в 1,6–2,8 раза быстрее. Когда мы дополнительно используем упакованные целочисленные скалярные произведения, производительность становится еще лучше. Это в 1,7–2,9 раза быстрее.
Проверьте поддержку браузера с помощью свойства wgslLanguageFeatures . Если графический процессор изначально не поддерживает упакованные скалярные произведения, браузер полифилизирует свою собственную реализацию.
// main.js
if (navigator.gpu.wgslLanguageFeatures.has('packed_4x8_integer_dot_product')) {
// Use dot4U8Packed, dot4I8Packed builtin
// functions in the shaders.
}
Следующий фрагмент кода diff (разница), подчеркивающий изменения, необходимые для использования упакованных целочисленных продуктов в шейдере WebGPU.
Раньше — шейдер WebGPU, который накапливает частичные скалярные произведения в переменную sum. В конце цикла сумма сохраняет полное скалярное произведение между вектором и одной строкой входной матрицы.
// my-dot-product.wgsl @compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid : vec3u) { var sum : f16; let start = gid.x * uniforms.dim; for (var i = 0u; i < uniforms.dim; i++) { let v1 : vec4<f16> = vector.values[i]; let v2 : vec4<f16> = matrix.values[start + i]; sum += dot(v1, v2); } }
After — шейдер WebGPU, написанный для использования скалярных произведений упакованных целых чисел. Основное отличие состоит в том, что вместо загрузки 4 значений с плавающей запятой из вектора и матрицы этот шейдер загружает одно 32-битное целое число. Это 32-битное целое число содержит данные четырех 8-битных целочисленных значений. Затем мы вызываем dot4U8Packed для вычисления скалярного произведения двух значений.
// my-dot-product.wgsl
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid : vec3u) {
var sum : f32;
let start = gid.x * uniforms.dim;
for (var i = 0u; i < uniforms.dim; i++) {
let v1 : u32 = vector.values[i];
let v2 : u32 = matrix.values[start + i];
sum += dot4U8Packed(v1, v2);
}
}
Как 16-битные скалярные произведения с плавающей запятой, так и упакованные целочисленные скалярные произведения — это встроенные функции Chrome, которые ускоряют работу искусственного интеллекта и машинного обучения. 16-битные числа с плавающей запятой доступны, если аппаратное обеспечение их поддерживает, а Chrome реализует упакованные целочисленные скалярные произведения на всех устройствах.
Вы можете использовать эти функции в Chrome Stable уже сегодня, чтобы повысить производительность.
Предлагаемые функции
Забегая вперед, мы исследуем еще две функции: подгруппы и кооперативное матричное умножение.
Функция подгрупп обеспечивает параллелизм на уровне SIMD для взаимодействия или выполнения коллективных математических операций, таких как суммирование более 16 чисел. Это обеспечивает эффективный обмен данными между потоками. Подгруппы поддерживаются современными API-интерфейсами графических процессоров с разными именами и в несколько разных формах.
Мы превратили общий набор в предложение, которое подали в группу стандартизации WebGPU. Мы создали прототипы подгрупп в Chrome под экспериментальным флагом и представили на обсуждение наши первоначальные результаты. Основной вопрос заключается в том, как обеспечить переносимость.
Кооперативное матричное умножение — это более позднее дополнение к графическим процессорам. Большое матричное умножение можно разбить на несколько меньших матричных умножений. Кооперативное матричное умножение выполняет умножение этих меньших блоков фиксированного размера за один логический шаг. На этом этапе группа потоков эффективно взаимодействует для вычисления результата.
Мы изучили поддержку базовых API-интерфейсов графических процессоров и планируем представить предложение группе стандартизации WebGPU. Как и в случае с подгруппами, мы ожидаем, что большая часть обсуждений будет сосредоточена вокруг переносимости.
Чтобы оценить производительность операций с подгруппами, в реальном приложении мы интегрировали экспериментальную поддержку подгрупп в MediaPipe и протестировали ее с прототипом Chrome для операций с подгруппами.
Мы использовали подгруппы в ядрах графического процессора на этапе предварительного заполнения большой языковой модели, поэтому я сообщаю об ускорении только на этапе предварительного заполнения. На графическом процессоре Intel мы видим, что подгруппы работают в два с половиной раза быстрее, чем базовый уровень. Однако эти улучшения не одинаковы для разных графических процессоров.
На следующей диаграмме показаны результаты применения подгрупп для оптимизации микротеста умножения матриц на нескольких потребительских графических процессорах. Умножение матриц — одна из самых сложных операций в больших языковых моделях. Данные показывают, что на многих графических процессорах подгруппы увеличивают скорость в два, пять и даже тринадцать раз по сравнению с базовым уровнем. Однако обратите внимание, что на первом графическом процессоре подгруппы ничуть не лучше.
Оптимизация графического процессора сложна
В конечном счете, лучший способ оптимизировать ваш графический процессор зависит от того, какой графический процессор предлагает клиент. Использование необычных новых функций графического процессора не всегда окупается так, как вы ожидаете, поскольку здесь может быть задействовано множество сложных факторов. Лучшая стратегия оптимизации для одного графического процессора может не быть лучшей стратегией для другого графического процессора.
Вы хотите минимизировать пропускную способность памяти, полностью используя вычислительные потоки графического процессора.
Шаблоны доступа к памяти также могут быть очень важны. Графические процессоры, как правило, работают намного лучше, когда вычислительные потоки обращаются к памяти по схеме, оптимальной для оборудования. Важно: на разном оборудовании графического процессора следует ожидать разных характеристик производительности. Возможно, вам придется запустить различные оптимизации в зависимости от графического процессора.
На следующей диаграмме мы взяли тот же алгоритм умножения матриц, но добавили еще одно измерение, чтобы дополнительно продемонстрировать влияние различных стратегий оптимизации, а также сложность и различия между разными графическими процессорами. Мы представили здесь новую технику, которую назовем «Swizzle». Swizzle оптимизирует шаблоны доступа к памяти, чтобы сделать их более оптимальными для оборудования.
Вы можете видеть, что изменение памяти имеет значительное влияние; иногда это даже более эффективно, чем подгруппы. На графическом процессоре 6 swizzle обеспечивает ускорение в 12 раз, а подгруппы — в 13 раз. В совокупности они дают невероятное ускорение в 26 раз. Для других графических процессоров иногда объединение swizzle и подгрупп работает лучше, чем каждый из них по отдельности. А на других графических процессорах лучше всего работает использование swizzle.
Настройка и оптимизация алгоритмов графического процессора для хорошей работы на любом оборудовании может потребовать большого опыта. Но, к счастью, ведется огромная талантливая работа над инфраструктурами библиотек более высокого уровня, такими как Mediapipe , Transformers.js , Apache TVM , ONNX Runtime Web и другие.
Библиотеки и платформы имеют хорошие возможности для решения сложных задач управления различными архитектурами графических процессоров и создания кода для конкретной платформы, который будет хорошо работать на клиенте.
Вынос
Команда Chrome продолжает способствовать развитию стандартов WebAssembly и WebGPU, чтобы улучшить веб-платформу для рабочих нагрузок машинного обучения. Мы инвестируем в более быстрые вычислительные примитивы, лучшее взаимодействие между веб-стандартами и обеспечиваем эффективную работу как больших, так и малых моделей на разных устройствах.
Наша цель — максимизировать возможности платформы, сохранив при этом лучшее от Интернета: охват, удобство использования и мобильность. И мы делаем это не в одиночку. Мы работаем в сотрудничестве с другими поставщиками браузеров в W3C и многими партнерами по разработке.
Мы надеемся, что при работе с WebAssembly и WebGPU вы помните следующее:
- Выводы ИИ теперь доступны в Интернете на всех устройствах. Это дает преимущества работы на клиентских устройствах, такие как снижение затрат на сервер, низкая задержка и повышение конфиденциальности.
- Хотя многие обсуждаемые функции актуальны в первую очередь для авторов платформы, ваши приложения могут получить от этого выгоду без особых затрат.
- Веб-стандарты изменчивы и развиваются, и мы всегда ищем обратную связь. Поделитесь своими для WebAssembly и WebGPU .
Благодарности
Мы хотели бы поблагодарить команду Intel по веб-графике, которая сыграла важную роль в разработке WebGPU f16 и функций целочисленного скалярного произведения. Мы хотели бы поблагодарить других членов рабочих групп WebAssembly и WebGPU в W3C, включая других поставщиков браузеров.
Спасибо командам искусственного интеллекта и машинного обучения как в Google, так и в сообществе разработчиков программного обеспечения с открытым исходным кодом за то, что вы стали замечательными партнерами. И, конечно же, всех наших товарищей по команде, которые делают все это возможным.
,Этот документ является продолжением усовершенствований WebAssembly и WebGPU для более быстрого веб-ИИ, часть 1 . Прежде чем продолжить, мы рекомендуем вам прочитать этот пост или посмотреть выступление на IO 24 .
Веб-графический процессор
WebGPU предоставляет веб-приложениям доступ к оборудованию графического процессора клиента для выполнения эффективных высокопараллельных вычислений. С момента запуска WebGPU в Chrome мы видели в сети невероятные демонстрации искусственного интеллекта (ИИ) и машинного обучения (МО).
Например, Web Stable Diffusion продемонстрировала, что можно использовать ИИ для создания изображений из текста прямо в браузере. Ранее в этом году собственная команда Google Mediapipe опубликовала экспериментальную поддержку вывода больших языковых моделей .
На следующей анимации показана Gemma , модель большого языка Google (LLM) с открытым исходным кодом, работающая полностью на устройстве в Chrome в режиме реального времени.
Следующая демонстрация Hugging Face модели Segment Anything Model компании Meta создает высококачественные маски объектов полностью на клиенте.
Это всего лишь пара замечательных проектов, демонстрирующих возможности WebGPU для искусственного интеллекта и машинного обучения. WebGPU позволяет этим и другим моделям работать значительно быстрее, чем на процессоре.
Тест Hugging Face WebGPU для встраивания текста демонстрирует огромное ускорение по сравнению с реализацией той же модели на ЦП. На ноутбуке Apple M1 Max WebGPU работал более чем в 30 раз быстрее. Другие сообщают, что WebGPU ускоряет тест более чем в 120 раз .
Улучшение функций WebGPU для искусственного интеллекта и машинного обучения
WebGPU отлично подходит для моделей искусственного интеллекта и машинного обучения, которые могут иметь миллиарды параметров благодаря поддержке вычислительных шейдеров . Вычислительные шейдеры выполняются на графическом процессоре и помогают выполнять параллельные операции с массивами над большими объемами данных.
Среди многочисленных улучшений WebGPU за последний год мы продолжили добавлять новые возможности для повышения производительности машинного обучения и искусственного интеллекта в Интернете. Недавно мы запустили две новые функции: 16-битные числа с плавающей запятой и упакованные целочисленные скалярные произведения.
16-битная с плавающей запятой
Помните, что рабочие нагрузки машинного обучения не требуют точности . shader-f16 — это функция, которая позволяет использовать тип f16 в языке шейдеров WebGPU. Этот тип с плавающей запятой занимает 16 бит вместо обычных 32 бит. f16 имеет меньший диапазон и менее точен, но для многих моделей ML этого достаточно.
Эта функция повышает эффективность несколькими способами:
Уменьшение памяти : тензоры с элементами f16 занимают половину места, что сокращает использование памяти вдвое. Вычисления на графическом процессоре часто ограничивают пропускную способность памяти, поэтому использование половины памяти часто означает, что шейдеры работают в два раза быстрее. Технически вам не нужна клавиша f16 для экономии пропускной способности памяти. Можно хранить данные в формате низкой точности, а затем расширять их до полного f32 в шейдере для вычислений. Но графический процессор тратит дополнительную вычислительную мощность на упаковку и распаковку данных.
Сокращение преобразования данных : f16 использует меньше вычислений за счет минимизации преобразования данных. Данные низкой точности можно сохранять и затем использовать напрямую, без преобразования.
Повышенный параллелизм . Современные графические процессоры способны одновременно помещать больше значений в исполнительные блоки графического процессора, что позволяет ему выполнять большее количество параллельных вычислений. Например, графический процессор, поддерживающий до 5 триллионов операций с плавающей запятой f32 в секунду, может поддерживать 10 триллионов операций с плавающей запятой f16 в секунду.
shader-f16 , эталонный эталон WebGPU Hugging Face для текстового встроенного эталона запускает эталон в 3 раза быстрее, чем F32 на ноутбуке Apple M1 Max.Webllm - это проект, который может запускать несколько крупных языковых моделей. Он использует Apache TVM , фреймворк компилятора с открытым исходным кодом.
Я попросил Webllm спланировать поездку в Париж, используя модель параметров Llama 3. Результаты показывают, что на фазе предварительной заполнения модели F16 в 2,1 раза быстрее, чем F32. Во время фазы декодирования это более чем в 1,3 раза быстрее.
Приложения должны сначала подтвердить, что адаптер графического процессора поддерживает F16, и, если он доступен, явно включает его при запросе устройства GPU. Если F16 не поддерживается, вы не можете запросить его в массиве requiredFeatures .
// main.js
const adapter = await navigator.gpu.requestAdapter();
const supportsF16 = adapter.features.has('shader-f16');
if (supportsF16) {
// Use f16.
const device = await adapter.requestDevice({
requiredFeatures: ['shader-f16'],
});
initApp(device);
}
Затем в ваших шейдерах WebGPU вы должны явно включить F16 вверху. После этого вы можете использовать его в шейдере, как любой другой тип данных.
// my-shader.wgsl
enable f16;
struct Data {
values : array<vec4<f16>>
}
@group(0) @binding(0) var<storage, read> data : Data;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid : vec3u) {
let value : vec4<f16> = data.values[gid.x];
...
}
Упакованные целочисленные точки продукты
Многие модели по -прежнему хорошо работают с 8 битами точности (половина F16). Это популярно среди LLMS и моделей изображений для сегментации и распознавания объектов. Тем не менее, качество выходных данных для моделей ухудшается с меньшей точностью, поэтому 8-битное квантование не подходит для каждого приложения.
Относительно мало графических процессоров изначально поддерживают 8-битные значения. Именно здесь появляются упакованные целочисленные точки. Мы отправили DP4A в Chrome 123 .
Современные графические процессоры имеют специальные инструкции, чтобы взять два 32-разрядных целых числа, интерпретировать их как 4 8-битных целых числа последовательно упакованных, и вычислить точечный продукт между их компонентами.
Это особенно полезно для ИИ и машинного обучения, потому что ядра умножения матрицы состоит из многих, многих точечных продуктов.
Например, давайте умножим матрицу 4 x 8 с вектором 8 x 1. Вычисление этого включает в себя принятие 4 точечных продуктов для расчета каждого из значений в выходном векторе; A, B, C и D.
Процесс вычисления каждого из этих выходов одинаковы; Мы посмотрим на шаги, связанные с вычислением одного из них. Перед каким-либо вычислением нам сначала нужно преобразовать 8-битные целочисленные данные в тип, с которым мы можем выполнить арифметику, например, F16. Затем мы запускаем умножение по элементу и, наконец, добавляем все продукты вместе. В общей сложности для всего умножения матрицы-вектора мы выполняем 40 целых числа переоборудования, чтобы распаковывать данные, 32 умножения плавания и 28 дополнений поплавок.
Для более крупных матриц с большим количеством операций упакованные целочисленные точки продукты могут помочь уменьшить объем работы.
Для каждого из выходов в векторе результатов мы выполняем две упакованные операции продукта Dot, используя dot4U8Packed язык затенения WebGPU, а затем добавляем результаты вместе. В общей сложности для всего умножения матрицы-вектора мы не выполняем какое-либо преобразование данных. Мы выполняем 8 упакованных точечных продуктов и 4 целочисленных дополнения.
Мы протестировали упакованные целочисленные точечные продукты с 8-битными данными о различных потребительских графических процессорах. По сравнению с 16-битной плавающей запятой, мы видим, что 8-битный составляет в 1,6-2,8 раза быстрее. Когда мы дополнительно используем упакованные целочисленные точки точки, производительность еще лучше. Это в 1,7 до 2,9 раза быстрее.
Проверьте поддержку браузера с помощью собственности wgslLanguageFeatures . Если графический процессор не поддерживает упакованные точечные продукты, то браузер полифиляет свою собственную реализацию.
// main.js
if (navigator.gpu.wgslLanguageFeatures.has('packed_4x8_integer_dot_product')) {
// Use dot4U8Packed, dot4I8Packed builtin
// functions in the shaders.
}
Следующий фрагмент кода Diff (разница) подчеркивает изменения, необходимые для использования упакованных целочисленных продуктов в шейдере WebGPU.
До - шейдер WebGPU, который накапливает частичные точечные продукты в переменную `sum '. В конце цикла «Sum» удерживает продукт полной точки между вектором и одной строкой входной матрицы.
// my-dot-product.wgsl @compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid : vec3u) { var sum : f16; let start = gid.x * uniforms.dim; for (var i = 0u; i < uniforms.dim; i++) { let v1 : vec4<f16> = vector.values[i]; let v2 : vec4<f16> = matrix.values[start + i]; sum += dot(v1, v2); } }
После - шейдер WebGPU, записанный для использования упакованных продуктов с целочисленными точками. Основное отличие состоит в том, что вместо загрузки 4 значений поплавок из вектора и матрицы этот шейдер загружает одно 32-разрядное целое число. Это 32-разрядное целое число содержит данные четырех 8-битных целочисленных значений. Затем мы называем dot4U8Packed для вычисления точечного продукта двух значений.
// my-dot-product.wgsl
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid : vec3u) {
var sum : f32;
let start = gid.x * uniforms.dim;
for (var i = 0u; i < uniforms.dim; i++) {
let v1 : u32 = vector.values[i];
let v2 : u32 = matrix.values[start + i];
sum += dot4U8Packed(v1, v2);
}
}
Как 16-битная плавающая точка, так и упакованные целочисленные точечные продукты представляют собой отправленные функции в Chrome, которые ускоряют AI и ML. 16-битная плавающая точка доступна, когда оборудование поддерживает его, а Chrome реализует упакованные целочисленные точки продукты на всех устройствах.
Вы можете использовать эти функции в Chrome Stable сегодня для достижения лучшей производительности.
Предлагаемые функции
С нетерпением ждем, мы исследуем еще две функции: подгруппы и кооперативная матрица умножны.
Функция подгрупп позволяет параллелизму уровня SIMD для общения или выполнения коллективных математических операций, таких как сумма для более чем 16 чисел. Это позволяет эффективно обмен данными. Подгруппы поддерживаются на современных API -интерфейсах графических процессоров, с различными именами и в немного разных формах.
Мы перевели общий набор в предложение, которое мы приняли в группу стандартизации WebGPU. И мы прототипировали подгруппы в Chrome за экспериментальным флагом и принесли наши первоначальные результаты в обсуждение. Основная проблема заключается в том, как обеспечить портативное поведение.
Кооперативная матрица Умножение является более недавним дополнением к графическим процессорам. Большая матрица может быть разбита на несколько меньших размножений матрицы. Умножение кооперативной матрицы выполняет умножение на этих меньших блоках с фиксированным размером на одном логическом шаге. В пределах этого шага группа потоков эффективно сотрудничает, чтобы вычислить результат.
Мы обследовали поддержку в базовых API -интерфейсах GPU и планируем представить предложение в группу стандартизации WebGPU. Как и в случае подгрупп, мы ожидаем, что большая часть обсуждения будет сосредоточена вокруг портативности.
Чтобы оценить эффективность операций подгрупп, в реальном приложении мы интегрировали экспериментальную поддержку подгрупп в MediaPipe и протестировали его с прототипом Chrome для операций подгрупп.
Мы использовали подгруппы в ядрах графических процессоров на фазе префиллы большой языковой модели, поэтому я сообщаю только о ускорении для фазы префилля. На графическом процессоре Intel мы видим, что подгруппы выполняют два с половиной раза быстрее, чем базовый. Тем не менее, эти улучшения не соответствуют различным графическим процессорам.
Следующая диаграмма показывает результаты применения подгрупп для оптимизации матрицы, умножающего микробанк, на нескольких графических процессорах потребителя. Умножение матрицы является одной из более тяжелых операций в моделях крупных языков. Данные показывают, что на многих графических процессорах подгруппы увеличивают скорость два, пять и даже тринадцать раз больше базовой линии. Тем не менее, обратите внимание, что в первом графическом процессоре подгруппы вообще не намного лучше.
Оптимизация графического процессора сложно
В конечном счете, лучший способ оптимизации вашего GPU зависит от того, какой графический процессор предлагает клиент. Использование новых функций GPU не всегда окупается, как вы можете ожидать, потому что может быть много сложных факторов. Лучшая стратегия оптимизации на одном графическом процессоре может быть не лучшей стратегией в другом графическом процессоре.
Вы хотите минимизировать пропускную способность памяти, полностью используя вычислительные потоки GPU.
Паттерны доступа к памяти тоже могут быть действительно важными. Графические процессоры имеют тенденцию работать намного лучше, когда вычислительные потоки получают доступ к памяти в шаблоне, который является оптимальным для аппаратного обеспечения. ВАЖНО: Вы должны ожидать различных характеристик производительности на различном оборудовании GPU. Вам может потребоваться запускать различные оптимизации в зависимости от графического процессора.
В следующей диаграмме мы приняли тот же алгоритм умножения матрицы, но добавили еще одно измерение, чтобы дополнительно продемонстрировать влияние различных стратегий оптимизации, а также сложность и дисперсию на разные графические процессоры. Мы представили здесь новую технику, которую мы будем называть «Swizzle». Swizzle оптимизирует шаблоны доступа к памяти, чтобы они были более оптимальными для оборудования.
Вы можете видеть, что Swizzle памяти оказывает значительное влияние; Иногда это даже более эффективно, чем подгруппы. На GPU 6 Swizzle обеспечивает 12 -кратное ускорение, в то время как подгруппы обеспечивают 13 -кратное ускорение. В совокупности у них невероятное 26 -кратное ускорение. Для других графических процессоров иногда подгруппы и подгруппы совместно работают лучше, чем один один. А на других графических процессорах, исключительно с помощью Swizzle выполняет лучшее.
Настройка и оптимизация алгоритмов графических процессоров для хорошо работы на каждом кусочке оборудования может потребовать много знаний. Но, к счастью, существует огромное количество талантливых работ, занимающихся библиотеками более высокого уровня, таких как MediaPipe , Transformers.js , Apache TVM , Onnx Runtime Web и многое другое.
Библиотеки и рамки хорошо расположены, чтобы справиться с сложностью управления различными архитектурами графических процессоров и генерации кода, специфичного для платформы, который будет хорошо работать на клиенте.
Вынос
Команда Chrome по -прежнему помогает развивать стандарты Webassembly и WebGPU для улучшения веб -платформы для рабочих нагрузок машинного обучения. Мы инвестируем в более быстрые вычислительные примитивы, лучшие взаимосвязи между веб -стандартами и следим за тем, чтобы модели как крупные, так и малые способны эффективно работать на разных устройствах.
Наша цель состоит в том, чтобы максимизировать возможности платформы, сохраняя при этом лучшее из Интернета: это охват, удобство использования и мобильность. И мы не делаем это в одиночку. Мы работаем в сотрудничестве с другими поставщиками браузеров в W3C и многими партнерами по разработке.
Мы надеемся, что вы помните следующее, работая с Webassembly и Webgpu:
- Вывод AI теперь доступен в Интернете, по всему устройствам. Это приносит преимущество работы на клиентских устройствах, таких как снижение стоимости сервера, низкая задержка и повышенная конфиденциальность.
- В то время как многие обсуждаемые функции имеют отношение к авторам фреймворка, ваши приложения могут выиграть без особых накладных расходов.
- Веб -стандарты плавны и развиваются, и мы всегда ищем обратную связь. Поделитесь своими за Webassembly и Webgpu .
Благодарности
Мы хотели бы поблагодарить команду Intel Web Graphics, которая сыграла важную роль в управлении WebGPU F16 и упакованными функциями продукта Integer Dot. Мы хотели бы поблагодарить других членов рабочих групп Webassembly и Webgpu в W3C, включая других поставщиков браузеров.
Спасибо командам ИИ и ML как в Google, так и в сообществе с открытым исходным кодом за то, что они были невероятными партнерами. И, конечно, все наши товарищи по команде, которые делают все это возможным.
,Этот документ является продолжением усовершенствований Webassembly и WebGPU для более быстрого веб -ИИ, часть 1 . Мы рекомендуем вам прочитать этот пост или посмотреть разговор в IO 24, прежде чем продолжить .
Webgpu
WebGPU предоставляет веб-приложения доступ к аппаратному обеспечению GPU клиента для выполнения эффективных, высокопараллельных вычислений. После запуска WebGPU в Chrome мы видели невероятные демонстрации искусственного интеллекта (ИИ) и машинного обучения (ML) в Интернете.
Например, веб -стабильная диффузия продемонстрировала, что можно было использовать ИИ для генерации изображений из текста, непосредственно в браузере. Ранее в этом году собственная команда MediaPipe от Google опубликовала экспериментальную поддержку для вывода с большой языком .
На следующей анимации показана Gemma , Google с открытым исходным кодом (LLM), полностью работающая в Chrome в Chrome в режиме реального времени.
Следующая демонстрация сегмента Meta's Segret Metue Model Ashipe создает высококачественные маски объектов полностью на клиенте.
Это всего лишь пара удивительных проектов, которые демонстрируют силу WebGPU для AI и ML. WebGPU позволяет этим моделям и другим работать значительно быстрее, чем они могли бы на процессоре.
Очередный эталон WebGPU обнимающего лица для встраивания текста демонстрирует огромное ускорение по сравнению с реализацией ЦП той же модели. На ноутбуке Apple M1 Max WebGPU был в 30 раз быстрее. Другие сообщили, что WebGPU ускоряет эталон более 120 раз .
Улучшение функций WebGPU для ИИ и ML
WebGPU отлично подходит для моделей искусственного интеллекта и ML, которые могут иметь миллиарды параметров, благодаря поддержке вычислительных шейдеров . Вычислить шейдеры, работающие на графическом процессоре и помогают запустить параллельные операции массива на больших объемах данных.
Среди многочисленных улучшений в WebGPU в прошлом году мы продолжали добавлять больше возможностей для повышения производительности ML и ИИ в Интернете. Недавно мы запустили две новые функции: 16-битная плавающая точка и упакованные продукты Dot Dot.
16-битная плавающая точка
Помните, что рабочие нагрузки ML не требуют точности . shader-f16 -это функция, которая позволяет использовать тип F16 на языке затенения WebGPU. Этот тип с плавающей запятой занимает 16 бит вместо обычных 32 бит. F16 имеет меньший диапазон и менее точен, но для многих моделей ML этого достаточно.
Эта функция повышает эффективность несколькими способами:
Уменьшенная память : тензоры с элементами F16 занимают половину пространства, которое сокращает использование памяти пополам. Расчеты графических процессоров часто бывают узкими местами на пропускной способности памяти, поэтому половина памяти часто может означать, что шейдеры работают вдвое быстрее. Технически вам не нужно F16, чтобы сохранить пропускную способность памяти. Можно сохранить данные в формате низкой конкретной решения, а затем расширить их до полного F32 в шейдере для вычислений. Но графический процессор тратит дополнительную вычислительную мощность, чтобы упаковать и распаковывать данные.
Снижение преобразования данных : F16 использует меньше вычислительных вычислений, минимизируя преобразование данных. Данные с низкой точностью могут храниться, а затем использовать непосредственно без преобразования.
Повышенная параллелизм : современные графические процессоры способны одновременно впитывать больше значений в единицах выполнения GPU, что позволяет ему выполнять большее количество параллельных вычислений. Например, графический процессор, который поддерживает до 5 триллионов F32 с плавающей точкой в секунду, может поддерживать 10 триллиона F16 с плавающей запятой в секунду.
shader-f16 , эталонный эталон WebGPU Hugging Face для текстового встроенного эталона запускает эталон в 3 раза быстрее, чем F32 на ноутбуке Apple M1 Max.Webllm - это проект, который может запускать несколько крупных языковых моделей. Он использует Apache TVM , фреймворк компилятора с открытым исходным кодом.
Я попросил Webllm спланировать поездку в Париж, используя модель параметров Llama 3. Результаты показывают, что на фазе предварительной заполнения модели F16 в 2,1 раза быстрее, чем F32. Во время фазы декодирования это более чем в 1,3 раза быстрее.
Приложения должны сначала подтвердить, что адаптер графического процессора поддерживает F16, и, если он доступен, явно включает его при запросе устройства GPU. Если F16 не поддерживается, вы не можете запросить его в массиве requiredFeatures .
// main.js
const adapter = await navigator.gpu.requestAdapter();
const supportsF16 = adapter.features.has('shader-f16');
if (supportsF16) {
// Use f16.
const device = await adapter.requestDevice({
requiredFeatures: ['shader-f16'],
});
initApp(device);
}
Затем в ваших шейдерах WebGPU вы должны явно включить F16 вверху. После этого вы можете использовать его в шейдере, как любой другой тип данных.
// my-shader.wgsl
enable f16;
struct Data {
values : array<vec4<f16>>
}
@group(0) @binding(0) var<storage, read> data : Data;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid : vec3u) {
let value : vec4<f16> = data.values[gid.x];
...
}
Упакованные целочисленные точки продукты
Многие модели по -прежнему хорошо работают с 8 битами точности (половина F16). Это популярно среди LLMS и моделей изображений для сегментации и распознавания объектов. Тем не менее, качество выходных данных для моделей ухудшается с меньшей точностью, поэтому 8-битное квантование не подходит для каждого приложения.
Относительно мало графических процессоров изначально поддерживают 8-битные значения. Именно здесь появляются упакованные целочисленные точки. Мы отправили DP4A в Chrome 123 .
Современные графические процессоры имеют специальные инструкции, чтобы взять два 32-разрядных целых числа, интерпретировать их как 4 8-битных целых числа последовательно упакованных, и вычислить точечный продукт между их компонентами.
Это особенно полезно для ИИ и машинного обучения, потому что ядра умножения матрицы состоит из многих, многих точечных продуктов.
Например, давайте умножим матрицу 4 x 8 с вектором 8 x 1. Вычисление этого включает в себя принятие 4 точечных продуктов для расчета каждого из значений в выходном векторе; A, B, C и D.
Процесс вычисления каждого из этих выходов одинаковы; Мы посмотрим на шаги, связанные с вычислением одного из них. Перед каким-либо вычислением нам сначала нужно преобразовать 8-битные целочисленные данные в тип, с которым мы можем выполнить арифметику, например, F16. Затем мы запускаем умножение по элементу и, наконец, добавляем все продукты вместе. В общей сложности для всего умножения матрицы-вектора мы выполняем 40 целых числа переоборудования, чтобы распаковывать данные, 32 умножения плавания и 28 дополнений поплавок.
Для более крупных матриц с большим количеством операций упакованные целочисленные точки продукты могут помочь уменьшить объем работы.
Для каждого из выходов в векторе результатов мы выполняем две упакованные операции продукта Dot, используя dot4U8Packed язык затенения WebGPU, а затем добавляем результаты вместе. В общей сложности для всего умножения матрицы-вектора мы не выполняем какое-либо преобразование данных. Мы выполняем 8 упакованных точечных продуктов и 4 целочисленных дополнения.
Мы протестировали упакованные целочисленные точечные продукты с 8-битными данными о различных потребительских графических процессорах. По сравнению с 16-битной плавающей запятой, мы видим, что 8-битный составляет в 1,6-2,8 раза быстрее. Когда мы дополнительно используем упакованные целочисленные точки точки, производительность еще лучше. Это в 1,7 до 2,9 раза быстрее.
Проверьте поддержку браузера с помощью собственности wgslLanguageFeatures . Если графический процессор не поддерживает упакованные точечные продукты, то браузер полифиляет свою собственную реализацию.
// main.js
if (navigator.gpu.wgslLanguageFeatures.has('packed_4x8_integer_dot_product')) {
// Use dot4U8Packed, dot4I8Packed builtin
// functions in the shaders.
}
Следующий фрагмент кода Diff (разница) подчеркивает изменения, необходимые для использования упакованных целочисленных продуктов в шейдере WebGPU.
До - шейдер WebGPU, который накапливает частичные точечные продукты в переменную `sum '. В конце цикла «Sum» удерживает продукт полной точки между вектором и одной строкой входной матрицы.
// my-dot-product.wgsl @compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid : vec3u) { var sum : f16; let start = gid.x * uniforms.dim; for (var i = 0u; i < uniforms.dim; i++) { let v1 : vec4<f16> = vector.values[i]; let v2 : vec4<f16> = matrix.values[start + i]; sum += dot(v1, v2); } }
После - шейдер WebGPU, записанный для использования упакованных продуктов с целочисленными точками. Основное отличие состоит в том, что вместо загрузки 4 значений поплавок из вектора и матрицы этот шейдер загружает одно 32-разрядное целое число. Это 32-разрядное целое число содержит данные четырех 8-битных целочисленных значений. Затем мы называем dot4U8Packed для вычисления точечного продукта двух значений.
// my-dot-product.wgsl
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid : vec3u) {
var sum : f32;
let start = gid.x * uniforms.dim;
for (var i = 0u; i < uniforms.dim; i++) {
let v1 : u32 = vector.values[i];
let v2 : u32 = matrix.values[start + i];
sum += dot4U8Packed(v1, v2);
}
}
Как 16-битная плавающая точка, так и упакованные целочисленные точечные продукты представляют собой отправленные функции в Chrome, которые ускоряют AI и ML. 16-битная плавающая точка доступна, когда оборудование поддерживает его, а Chrome реализует упакованные целочисленные точки продукты на всех устройствах.
Вы можете использовать эти функции в Chrome Stable сегодня для достижения лучшей производительности.
Предлагаемые функции
С нетерпением ждем, мы исследуем еще две функции: подгруппы и кооперативная матрица умножны.
Функция подгрупп позволяет параллелизму уровня SIMD для общения или выполнения коллективных математических операций, таких как сумма для более чем 16 чисел. Это позволяет эффективно обмен данными. Подгруппы поддерживаются на современных API -интерфейсах графических процессоров, с различными именами и в немного разных формах.
Мы перевели общий набор в предложение, которое мы приняли в группу стандартизации WebGPU. И мы прототипировали подгруппы в Chrome за экспериментальным флагом и принесли наши первоначальные результаты в обсуждение. Основная проблема заключается в том, как обеспечить портативное поведение.
Кооперативная матрица Умножение является более недавним дополнением к графическим процессорам. Большая матрица может быть разбита на несколько меньших размножений матрицы. Умножение кооперативной матрицы выполняет умножение на этих меньших блоках с фиксированным размером на одном логическом шаге. В пределах этого шага группа потоков эффективно сотрудничает, чтобы вычислить результат.
Мы обследовали поддержку в базовых API -интерфейсах GPU и планируем представить предложение в группу стандартизации WebGPU. Как и в случае подгрупп, мы ожидаем, что большая часть обсуждения будет сосредоточена вокруг портативности.
Чтобы оценить эффективность операций подгрупп, в реальном приложении мы интегрировали экспериментальную поддержку подгрупп в MediaPipe и протестировали его с прототипом Chrome для операций подгрупп.
Мы использовали подгруппы в ядрах графических процессоров на фазе префиллы большой языковой модели, поэтому я сообщаю только о ускорении для фазы префилля. На графическом процессоре Intel мы видим, что подгруппы выполняют два с половиной раза быстрее, чем базовый. Тем не менее, эти улучшения не соответствуют различным графическим процессорам.
Следующая диаграмма показывает результаты применения подгрупп для оптимизации матрицы, умножающего микробанк, на нескольких графических процессорах потребителя. Умножение матрицы является одной из более тяжелых операций в моделях крупных языков. Данные показывают, что на многих графических процессорах подгруппы увеличивают скорость два, пять и даже тринадцать раз больше базовой линии. Тем не менее, обратите внимание, что в первом графическом процессоре подгруппы вообще не намного лучше.
Оптимизация графического процессора сложно
В конечном счете, лучший способ оптимизации вашего GPU зависит от того, какой графический процессор предлагает клиент. Использование новых функций GPU не всегда окупается, как вы можете ожидать, потому что может быть много сложных факторов. Лучшая стратегия оптимизации на одном графическом процессоре может быть не лучшей стратегией в другом графическом процессоре.
Вы хотите минимизировать пропускную способность памяти, полностью используя вычислительные потоки GPU.
Паттерны доступа к памяти тоже могут быть действительно важными. Графические процессоры имеют тенденцию работать намного лучше, когда вычислительные потоки получают доступ к памяти в шаблоне, который является оптимальным для аппаратного обеспечения. ВАЖНО: Вы должны ожидать различных характеристик производительности на различном оборудовании GPU. Вам может потребоваться запускать различные оптимизации в зависимости от графического процессора.
В следующей диаграмме мы приняли тот же алгоритм умножения матрицы, но добавили еще одно измерение, чтобы дополнительно продемонстрировать влияние различных стратегий оптимизации, а также сложность и дисперсию на разные графические процессоры. Мы представили здесь новую технику, которую мы будем называть «Swizzle». Swizzle оптимизирует шаблоны доступа к памяти, чтобы они были более оптимальными для оборудования.
Вы можете видеть, что Swizzle памяти оказывает значительное влияние; Иногда это даже более эффективно, чем подгруппы. На GPU 6 Swizzle обеспечивает 12 -кратное ускорение, в то время как подгруппы обеспечивают 13 -кратное ускорение. В совокупности у них невероятное 26 -кратное ускорение. Для других графических процессоров иногда подгруппы и подгруппы совместно работают лучше, чем один один. А на других графических процессорах, исключительно с помощью Swizzle выполняет лучшее.
Настройка и оптимизация алгоритмов графических процессоров для хорошо работы на каждом кусочке оборудования может потребовать много знаний. Но, к счастью, существует огромное количество талантливых работ, занимающихся библиотеками более высокого уровня, таких как MediaPipe , Transformers.js , Apache TVM , Onnx Runtime Web и многое другое.
Библиотеки и рамки хорошо расположены, чтобы справиться с сложностью управления различными архитектурами графических процессоров и генерации кода, специфичного для платформы, который будет хорошо работать на клиенте.
Вынос
Команда Chrome по -прежнему помогает развивать стандарты Webassembly и WebGPU для улучшения веб -платформы для рабочих нагрузок машинного обучения. Мы инвестируем в более быстрые вычислительные примитивы, лучшие взаимосвязи между веб -стандартами и следим за тем, чтобы модели как крупные, так и малые способны эффективно работать на разных устройствах.
Наша цель состоит в том, чтобы максимизировать возможности платформы, сохраняя при этом лучшее из Интернета: это охват, удобство использования и мобильность. И мы не делаем это в одиночку. Мы работаем в сотрудничестве с другими поставщиками браузеров в W3C и многими партнерами по разработке.
Мы надеемся, что вы помните следующее, работая с Webassembly и Webgpu:
- Вывод AI теперь доступен в Интернете, по всему устройствам. Это приносит преимущество работы на клиентских устройствах, таких как снижение стоимости сервера, низкая задержка и повышенная конфиденциальность.
- В то время как многие обсуждаемые функции имеют отношение к авторам фреймворка, ваши приложения могут выиграть без особых накладных расходов.
- Веб -стандарты плавны и развиваются, и мы всегда ищем обратную связь. Поделитесь своими за Webassembly и Webgpu .
Благодарности
Мы хотели бы поблагодарить команду Intel Web Graphics, которая сыграла важную роль в управлении WebGPU F16 и упакованными функциями продукта Integer Dot. Мы хотели бы поблагодарить других членов рабочих групп Webassembly и Webgpu в W3C, включая других поставщиков браузеров.
Спасибо командам ИИ и ML как в Google, так и в сообществе с открытым исходным кодом за то, что они были невероятными партнерами. И, конечно, все наши товарищи по команде, которые делают все это возможным.