cutlass
CUTLASS
CUTLASS 3.6.0 — октябрь 2024 г.
CUTLASS — это набор абстракций шаблонов CUDA C++ для реализации высокопроизводительного матричного умножения (GEMM) и связанных с ним вычислений на всех уровнях и масштабах внутри CUDA. Он включает стратегии иерархической декомпозиции и перемещения данных, аналогичные тем, которые используются для реализации cuBLAS и cuDNN. CUTLASS разлагает эти «подвижные части» на многократно используемые модульные программные компоненты, абстрагированные классами шаблонов C++. Примитивы для разных уровней концептуальной иерархии распараллеливания можно специализировать и настраивать с помощью пользовательских размеров листов, типов данных и других алгоритмических политик. Полученная в результате гибкость упрощает их использование в качестве строительных блоков в пользовательских ядрах и приложениях.
Для поддержки широкого спектра приложений CUTLASS обеспечивает обширную поддержку вычислений смешанной точности, предоставляя специализированные абстракции перемещения данных и умножения-накопления для плавающей запятой половинной точности (FP16), BFloat16 (BF16), Tensor Float 32 (TF32), с плавающей запятой одинарной точности (FP32), эмуляция FP32 с помощью инструкций ядра тензора, типы с плавающей запятой двойной точности (FP64), целочисленные типы данных (4b и 8б) и двоичные типы данных (1б). CUTLASS демонстрирует синхронные операции умножения матриц, ориентированные на программируемые высокопроизводительные тензорные ядра, реализованные на архитектурах NVIDIA Volta, Turing, Ampere и Hopper.
См. Краткое руководство, чтобы быстро приступить к работе.
См. список функций, в котором приведен список операций, поддерживаемых на каждом уровне иерархии модели выполнения.
В CUTLASS 3.0 представлена новая базовая библиотека CuTe для описания и управления тензорами потоков и данных. CuTe — это набор абстракций шаблонов C++ CUDA для определения и работы с иерархически многомерными макетами потоков и данных. CuTe предоставляет объекты Layout и Tensor , которые компактно упаковывают тип, форму, пространство памяти и расположение данных, одновременно выполняя сложную индексацию для пользователя. Это позволяет программистам сосредоточиться на логических описаниях своих алгоритмов, в то время как CuTe выполняет за них механическую бухгалтерию. С помощью этих инструментов мы можем быстро проектировать, реализовывать и изменять все операции плотной линейной алгебры.
Основные абстракции CuTe — это иерархически многомерные макеты, которые можно составлять из массивов данных для представления тензоров. Представление макетов достаточно мощное, чтобы представить почти все, что нам нужно для реализации эффективной плотной линейной алгебры. Макеты также можно комбинировать и манипулировать ими с помощью функциональной композиции, на основе которой мы строим большой набор общих операций, таких как разбиение на плитки и секционирование.
CUTLASS 3.0 и более поздние версии используют CuTe во всей иерархии GEMM в своих шаблонах. Это значительно упрощает проектирование и улучшает компоновку и читаемость кода. Дополнительную документацию по CuTe можно найти в специальном каталоге документации.
В дополнение к GEMM, CUTLASS реализует высокопроизводительную свертку с помощью неявного алгоритма GEMM. Неявный GEMM — это формулировка операции свертки как GEMM, тем самым используя преимущества модульного конвейера GEMM CUTLASS. Это позволяет CUTLASS создавать свертки, повторно используя высокооптимизированные компоненты GEMM.
CUTLASS 3.6.0 — это обновление CUTLASS, в которое добавлены:
Хоппер структурированный разреженный GEMM.
РП16
РП8
INT8
ТФ32
Рефакторинг kernel::ConvUniversal API для приведения его в соответствие с gemm::GemmUniversal . Теперь API свертки 3.x больше не считается бета-версией API.
Улучшенный смешанный ввод GEMM и реализация таблицы поиска для режима только масштабирования INT4 x FP8 .
Узлы EVT для выбора Top-K и пример softmax и GEMM с их использованием.
Программно-зависимый запуск (PDL), который использует новую функцию Hopper для ускорения двух последовательных ядер и соответствующей документации.
Новый инструмент отладки synclog для выгрузки всех событий синхронизации из ядра в файл. Подробности смотрите в документации по synclog.
Новый эпилог с поддержкой TMA для сгруппированного GEMM, который обеспечивает значительное улучшение производительности, а также поддержку EVT.
Эпилог массива указателей с поддержкой SIMT.
Новое расписание ядра Ping-Pong для группового GEMM и некоторые другие оптимизации.
Новая стратегия создания экземпляров для ядер профилировщика CUTLASS, а также улучшенная документация для уровня создания экземпляров в профилировщике CUTLASS.
Новая аппаратная поддержка сравнений и вычислений cutlass::bfloat16_t
Исправлено использование isnan в Windows для half_t .
Минимальные требования:
Архитектура: Вольта
Компилятор: Должен поддерживать как минимум C++17.
Версия набора инструментов CUDA: 11.4.
Начиная с CUTLASS 3.0, в CUTLASS удалена поддержка следующего:
Архитектуры графических процессоров Maxwell и Pascal
Убунту 16.04
КУДА 10.2
Версии языка C++ ниже 17.
Подробный список выпусков и обновлений см. в CHANGELOG.
Примитивы CUTLASS очень эффективны. При использовании для построения ядер GEMM на уровне устройства они демонстрируют пиковую производительность, сравнимую с cuBLAS для скалярных вычислений GEMM. На рисунке выше показано постоянное улучшение производительности CUTLASS на NVIDIA H100 (архитектура NVIDIA Hopper) с момента выхода CUTLASS 3.1. CUTLASS 3.5.1 был скомпилирован с использованием набора инструментов CUDA 12.5u1. Операции Tensor Core реализуются с использованием инструкций CUDA mma и wgmma.
При использовании строительных блоков CUTLASS для создания неявных ядер gemm (Fprop, Dgrad и Wgrad) на уровне устройства производительность CUTLASS также сопоставима с cuDNN при запуске слоев Resnet-50 на NVIDIA A100, как показано на рисунке выше. Операции Tensor Core реализуются с помощью инструкции CUDA mma.
CUTLASS требует хост-компилятора C++17 и лучше всего работает при сборке с помощью набора инструментов CUDA 12.4 . Он также совместим с CUDA 11.4, CUDA 11.5, CUDA 11.6, CUDA 11.7, CUDA 11.8, CUDA 12.0, CUDA 12.1, CUDA 12.2.2, CUDA 12.3.1 и CUDA 12.3.2.
Мы протестировали следующие среды.
| Операционная система | Компилятор |
|---|---|
| Убунту 18.04 | GCC 7.5.0 |
| Убунту 20.04 | GCC 10.3.0 |
| Убунту 22.04 | GCC 11.2.0 |
| Убунту 22.04 | Кланг 10.0.0 |
| Убунту 22.04 | Кланг 14.0.6 |
| Убунту 22.04 | Кланг 17.0.6 |
| Windows 10.0 | Visual Studio 2019 v16.11.27 |
Примечание. В GCC 8.5.0 имеются известные регрессии в отношении выражений свертки и перегруженных операторов. Рекомендуется использовать GCC 7.5.0 или (предпочтительно) GCC >= 9.
CUTLASS успешно работает на следующих графических процессорах NVIDIA и, как ожидается, будет эффективен на графических процессорах NVIDIA на базе архитектур Volta, Turing, Ampere, Ada и Hopper.
| графический процессор | Вычислительные возможности CUDA | Минимальный набор инструментов CUDA, необходимый для CUTLASS-3 |
|---|---|---|
| Графический процессор NVIDIA V100 с тензорным ядром | 7.0 | 11.4 |
| NVIDIA ТитанВ | 7.0 | 11.4 |
| NVIDIA GeForce RTX 2080 TI, 2080, 2070 | 7,5 | 11.4 |
| NVIDIA Т4 | 7,5 | 11.4 |
| Графический процессор NVIDIA A100 с тензорным ядром | 8.0 | 11.4 |
| NVIDIA А10 | 8,6 | 11.4 |
| NVIDIA GeForce RTX 3090 | 8,6 | 11.4 |
| NVIDIA GeForce RTX 4090 | 8,9 | 11,8 |
| NVIDIA L40 | 8,9 | 11,8 |
| Графический процессор NVIDIA H100 с тензорным ядром | 9,0 | 11,8 |
В общем, код PTX, созданный для одной целевой архитектуры, может быть запущен на будущих архитектурах (т. е. он совместим с будущими версиями). Однако в CUDA 12.0 появилась концепция «функций с ускорением архитектуры», PTX которых не имеет гарантий прямой совместимости. Некоторые инструкции Hopper PTX подпадают под эту категорию функций с архитектурным ускорением и, следовательно, требуют целевой архитектуры sm_90a (обратите внимание на добавленную букву «a»). Более подробную информацию об этой и других инструкциях по ускорению архитектуры можно найти в документации CUDA.
Информация о целевой архитектуре передается в CUTLASS через флаг cmake CUTLASS_NVCC_ARCHS . Чтобы максимизировать производительность Hopper GH100, пользователям необходимо собрать CUTLASS с целевой архитектурой 90a . Если пользователь случайно собирает ядро, которое использует функции SM90a (например, инструкции Hopper Tensor Core), используя цель SM90 (обратите внимание на отсутствие «a»), с помощью CUDA Toolkit 12 или 11.8, ожидается, что ядро выйдет из строя во время выполнения. ошибка.
cmake .. -DCUTLASS_NVCC_ARCHS="90a"
Пожалуйста, обратитесь к функциональной документации для получения подробной информации о том, какие ядра требуют каких целевых архитектур.
CUTLASS описан в следующих документах и сопроводительной документации Doxygen.
Краткое руководство — сборка и запуск CUTLASS
Функциональность — суммирует функциональные возможности, доступные в CUTLASS.
Эффективный GEMM в CUDA — описывает, как ядра GEMM могут быть эффективно реализованы в CUDA.
CUTLASS 3.x Design — описывает дизайн CUTLASS 3.x, его преимущества и то, как CuTe позволяет нам писать гораздо больше компонуемых компонентов.
GEMM API 3.x — описывает модель CUTLASS 3.x GEMM и концепции шаблонов C++.
GEMM API 2.x — описывает модель CUTLASS 2.x GEMM и концепции шаблонов C++.
Неявная свертка GEMM — описывает двухмерную и трехмерную свертку в CUTLASS.
Организация кода — описывает организацию и содержание проекта CUTLASS.
Терминология — описывает термины, используемые в коде.
Programming Guidelines — рекомендации по написанию эффективного современного CUDA C++.
Фундаментальные типы — описывают базовые классы C++, используемые в CUTLASS для представления числовых величин и массивов.
Макеты — описывает расположение матриц и тензоров в памяти.
Tile Iterators — описывает концепции C++ для перебора ячеек матриц в памяти.
CUTLASS Profiler — приложение для профилирования, управляемое из командной строки
CUTLASS Utilities — дополнительные шаблоны, используемые для облегчения быстрой разработки.
Запуск зависимого ядра — описывает новую функцию в Hopper, которая позволяет перекрывать зависимые ядра в одном потоке, и то, как она используется в CUTLASS.
Мы также описали структуру эффективного GEMM в нашем выступлении на конференции GPU Technology Conference 2018.
CUTLASS: программные примитивы для плотной линейной алгебры на всех уровнях и масштабах в CUDA
Разработка ядер CUDA для максимально возможного использования тензорных ядер на NVIDIA A100
Ускорение свертки с помощью тензорных ядер в CUTLASS
Ускорение обратного градиента данных за счет увеличения использования тензорного ядра в CUTLASS
CUTLASS: Python API, улучшения и NVIDIA Hopper
CUTLASS — это библиотека шаблонов только для заголовков, и ее не нужно создавать для использования в других проектах. Клиентские приложения должны ориентироваться на каталог include/ CUTLASS в своих путях включения.
Модульные тесты, примеры и утилиты CUTLASS можно создавать с помощью CMake. Минимальная версия CMake указана в кратком руководстве. Убедитесь, что переменная среды CUDACXX указывает на NVCC в наборе инструментов CUDA, установленном в вашей системе.
$ экспорт CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc Создайте каталог сборки в проекте CUTLASS, затем запустите CMake. По умолчанию CUTLASS собирает ядра для версий архитектуры CUDA 5.0, 6.0, 6.1, 7.0, 7.5, 8.0, 8.6, 8.9 и 9.0. Чтобы сократить время компиляции, вы можете указать архитектуры для сборки CUTLASS, изменив параметр конфигурации CMake CUTLASS_NVCC_ARCHS .
$ mkdir build && cd build $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # компилируется для архитектуры NVIDIA Ampere
Из каталога build/ скомпилируйте и запустите модульные тесты CUTLASS, создав целевой test_unit с помощью make.
Модульные тесты организованы в виде нескольких двоичных файлов, отражающих пространства имен верхнего уровня CUTLASS, и их можно выполнять параллельно с помощью аргумента командной строки -j команды make.
$ make test_unit -j ... ... ... [----------] Демонтаж глобальной тестовой среды [==========] Выполнено 946 тестов из 57 тестовых случаев. (всего 10812 мс) [ ПРОЙДЕНО ] 946 тестов.
Все тесты должны проходить на поддерживаемых платформах, хотя точное количество тестов может меняться со временем.
CUTLASS организован как библиотека только заголовков вместе с утилитами, инструментами, примерами и модульными тестами. Документация Doxygen предоставляет полный список файлов, классов и концепций шаблонов, определенных в проекте CUTLASS.
Подробное объяснение организации исходного кода можно найти в документации CUTLASS, но ниже кратко описаны некоторые основные компоненты.
include/ # client applications should target this directory in their build's include paths cutlass/ # CUDA Templates for Linear Algebra Subroutines and Solvers - headers only arch/ # direct exposure of architecture features (including instruction-level GEMMs) conv/ # code specialized for convolution epilogue/ # code specialized for the epilogue of gemm/convolution gemm/ # code specialized for general matrix product computations layout/ # layout definitions for matrices, tensors, and other mathematical objects in memory platform/ # CUDA-capable Standard Library components reduction/ # bandwidth-limited reduction kernels that do not fit the "gemm" model thread/ # simt code that can be performed within a CUDA thread transform/ # code specialized for layout, type, and domain transformations * # core vocabulary types, containers, and basic numeric operations cute/ # CuTe Layout, layout algebra, MMA/Copy atoms, tiled MMA/Copy algorithm/ # Definitions of core operations such as copy, gemm, and operations on cute::tuples arch/ # Bare bones PTX wrapper structs for copy and math instructions atom/ # Meta-information either link to or built from arch/ operators mma_atom.hpp # cute::Mma_Atom and cute::TiledMma copy_atom.hpp # cute::Copy_Atom and cute::TiledCopy *sm*.hpp # Arch specific meta-information for copy and math operations * # Core library types such as Shape, Stride, Layout, Tensor, and associated operations
В примерах CUTLASS SDK шаблоны CUTLASS применяются для реализации базовых вычислений.
tools/ library/ # CUTLASS Instance Library - contains instantiations of all supported CUTLASS templates include/ cutlass/ library/ profiler/ # CUTLASS Profiler - command-line utility for executing operations in the # CUTLASS Library util/ # CUTLASS Utilities - contains numerous helper classes for include/ # manging tensors in device memory, reference cutlass/ # implementations for GEMM, random initialization util/ # of tensors, and I/O.
Каталог test/unit/ состоит из модульных тестов, реализованных с помощью Google Test, которые демонстрируют базовое использование компонентов Core API и полные тесты вычислений CUTLASS GEMM.
Инструкции по созданию и запуску модульных тестов описаны в кратком руководстве.
Каталог tools/profiler/ содержит утилиту командной строки для запуска каждого ядра GEMM. Его можно построить следующим образом:
$ make Cutlass_profiler -j16
По умолчанию для каждого типа данных, математической инструкции и макета создается только один экземпляр размера плитки. Чтобы создать экземпляры всех, установите следующую переменную среды при запуске CMake из пустого каталога build/ . Будьте осторожны, это приводит к десяткам тысяч ядер и длительному времени сборки. Это также приведет к большому двоичному размеру и на некоторых платформах компоновщику не удастся собрать библиотеку. Поэтому настоятельно рекомендуется генерировать только подмножество ядер, как показано в подразделе ниже.
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all ... $ make Cutlass_profiler -j16
Чтобы скомпилировать строго одно ядро или небольшой набор ядер, можно использовать список имен ядер, разделенных запятыми, с подстановочными знаками для сокращения набора ядер. В следующих примерах показано создание одного ядра или его подмножества для архитектуры NVIDIA Ampere и Turing:
Чтобы скомпилировать подмножество ядер Tensor Core GEMM с накоплением FP32 и вводом FP16, ориентированным на архитектуру NVIDIA Ampere и Turing, используйте приведенную ниже командную строку cmake:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8 ... $ make Cutlass_profiler -j16
Пример командной строки для профилирования подмножества ядер Tensor Core GEMM выглядит следующим образом:
./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096
...
============================
Идентификатор проблемы: 1
Поставщик: CUTLASS
Вид операции: gemm
Операция: Cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
Статус: Успех
Проверка: ВКЛ.
Расположение: Пройдено
reference_device: Пройдено
cuBLAS: Пройдено
Аргументы: --gemm_kind=универсальный --m=3456 --n=4096 --k=4096 --A=f16:столбец --B=f16:строка --C=f32:столбец --alpha=1
--beta=0 --split_k_slices=1 --batch_count=1 --op_class=tensorop --accum=f32 --cta_m=256 --cta_n=128
--cta_k=32 --stages=2 --warps_m=4 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=8 --min_cc=75
--max_cc=1024
Байты: 118489088 байт
Флопов: 115992428544 флопа
Время выполнения: 1,55948 мс
Память: 70,7616 ГиБ/с
Математика: 74378,8 ГФЛОПС/с
============================
...Чтобы скомпилировать одно ядро SGEMM, ориентированное на архитектуру NVIDIA Ampere и Turing, используйте следующую командную строку cmake:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1 ... $ make Cutlass_profiler -j16
Пример командной строки для профилирования одного ядра SGEMM CUDA выглядит следующим образом:
$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096
============================
Идентификатор проблемы: 1
Поставщик: CUTLASS
Вид операции: gemm
Операция: Cutlass_simt_sgemm_128x128_8x2_nn_align1
Статус: Успех
Проверка: ВКЛ.
Расположение: Пройдено
cuBLAS: Пройдено
Аргументы: --m=3456 --n=4096 --k=4096 --A=f32:столбец --B=f32:столбец --C=f32:столбец --alpha=1 --beta=0 -- Split_k_slices=1
--batch_count=1 --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4
--warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024
Байты: 180355072 байт
Флопов: 115992428544 флопа
Время выполнения: 6,73655 мс
Память: 24,934 ГиБ/с
Математика: 17218,4 ГФЛОПС/с.
============================Чтобы скомпилировать подмножество ядер свертки тензорного ядра, реализующих прямое распространение (fprop) с накоплением FP32 и вводом FP16, ориентированным на архитектуру NVIDIA Ampere и Turing, используйте приведенную ниже командную строку cmake:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16 ... $ make Cutlass_profiler -j16
Пример командной строки для профилирования подмножества ядер свертки Tensor Core выглядит следующим образом:
$ ./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*fprop_optimized_f16 --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3
...
============================
Идентификатор проблемы: 1
Поставщик: CUTLASS
Вид операции: conv2d
Операция: Cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc
Статус: Успех
Проверка: ВКЛ.
Расположение: Пройдено
reference_device: Пройдено
Аргументы: --conv_kind=fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --p=224 --q =224 --pad_h=1 --pad_w=1
--stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1 --Activation=f16:nhwc --Filter=f16:nhwc --Output=f32:nhwc
--conv_mode=cross --iterator_algorithm=оптимизированный --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1
--eq_gemm_provider=none --op_class=tensorop --accum=f32 --cta_m=128 --cta_n=128 --cta_k=32 --stages=5
--warps_m=2 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=16 --min_cc=80 --max_cc=1024
Байты: 1130659840 байт.
Флопов: 118482796544 флопа
Время выполнения: 0,711496 мс
Память: 1479,99 ГиБ/с
Математика: 166526 ГФЛОПС/с
============================
...Чтобы скомпилировать и запустить одно ядро свертки CUDA Core, реализующее прямое распространение (fprop) с накоплением F32 и вводом FP32, ориентированным на архитектуру NVIDIA Ampere и Turing, используйте приведенную ниже командную строку cmake:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc ... $ make Cutlass_profiler -j16
Пример командной строки для профилирования одного ядра свертки CUDA Core:
$ ./tools/profiler/cutlass_profiler --kernels=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3
============================
Идентификатор проблемы: 1
Поставщик: CUTLASS
Вид операции: conv2d
Операция: Cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
Статус: Успех
Проверка: ВКЛ.
Расположение: Пройдено
reference_device: Пройдено
Аргументы: --conv_kind=fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --p=224 --q =224 --pad_h=1 --pad_w=1
--stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1 --Activation=f32:nhwc --Filter=f32:nhwc --Output=f32:nhwc
--conv_mode=cross --iterator_algorithm=оптимизированный --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1
--eq_gemm_provider=none --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4
--warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024
Байты: 2055798784 байт.
Флопов: 118482796544 флопа
Время выполнения: 7,34266 мс
Память: 260,752 ГиБ/с
Математика: 16136,2 ГФЛОПС/с.
============================Пожалуйста, перейдите по ссылкам для получения дополнительных примеров CMake по выборочной компиляции ядер CUTLASS:
Примеры GEMM CMake
Примеры CMake неявной свертки GEMM
Более подробная информация о профилировщике CUTLASS описана здесь.
CUTLASS выпущен корпорацией NVIDIA как программное обеспечение с открытым исходным кодом под «новой» лицензией BSD из трех пунктов.
Официальный список разработчиков и участников CUTLASS доступен здесь: УЧАСТНИКИ.
Авторские права (c) 2017–2024 КОРПОРАЦИЯ NVIDIA И ДОЧЕРНИЕ ПРЕДСТАВИТЕЛИ. Все права защищены. Идентификатор лицензии SPDX: пункт BSD-3
Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met: 1. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. 2. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution. 3. Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission. THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
