cutlass
CUTLASS
CUTLASS 3.6.0 – outubro de 2024
CUTLASS é uma coleção de abstrações de modelos CUDA C++ para implementar multiplicação de matrizes de alto desempenho (GEMM) e cálculos relacionados em todos os níveis e escalas dentro de CUDA. Ele incorpora estratégias para decomposição hierárquica e movimentação de dados semelhantes àquelas usadas para implementar cuBLAS e cuDNN. CUTLASS decompõe essas "partes móveis" em componentes de software modulares e reutilizáveis, abstraídos por classes de modelo C++. Primitivos para diferentes níveis de uma hierarquia de paralelização conceitual podem ser especializados e ajustados por meio de tamanhos de blocos personalizados, tipos de dados e outras políticas algorítmicas. A flexibilidade resultante simplifica seu uso como blocos de construção em kernels e aplicativos personalizados.
Para suportar uma ampla variedade de aplicações, o CUTLASS fornece amplo suporte para cálculos de precisão mista, fornecendo movimentação de dados especializada e abstrações de acumulação múltipla para ponto flutuante de meia precisão (FP16), BFloat16 (BF16), Tensor Float 32 (TF32), ponto flutuante de precisão simples (FP32), emulação FP32 via instrução de núcleo tensor, tipos de ponto flutuante de precisão dupla (FP64), tipos de dados inteiros (4b e 8b) e tipos de dados binários (1b). CUTLASS demonstra operações de multiplicação de matriz síncrona de warp visando os Tensor Cores programáveis e de alto rendimento implementados pelas arquiteturas Volta, Turing, Ampere e Hopper da NVIDIA.
Consulte o Guia de início rápido para começar rapidamente.
Consulte a listagem de funcionalidades para obter a lista de operações suportadas em cada nível da hierarquia do modelo de execução.
CUTLASS 3.0 introduziu uma nova biblioteca central, CuTe, para descrever e manipular tensores de threads e dados. CuTe é uma coleção de abstrações de modelos C++ CUDA para definir e operar em layouts hierarquicamente multidimensionais de threads e dados. CuTe fornece objetos Layout e Tensor que empacotam de forma compacta o tipo, forma, espaço de memória e layout dos dados, enquanto executa a indexação complicada para o usuário. Isso permite que os programadores se concentrem nas descrições lógicas de seus algoritmos enquanto CuTe faz a contabilidade mecânica para eles. Com essas ferramentas, podemos projetar, implementar e modificar rapidamente todas as operações de álgebra linear densa.
As abstrações centrais do CuTe são layouts hierarquicamente multidimensionais que podem ser compostos com matrizes de dados para representar tensores. A representação de layouts é poderosa o suficiente para representar quase tudo que precisamos para implementar álgebra linear densa e eficiente. Os layouts também podem ser combinados e manipulados por meio da composição funcional, na qual construímos um grande conjunto de operações comuns, como agrupamento e particionamento.
CUTLASS 3.0 e posteriores adotam CuTe em toda a hierarquia GEMM em seus modelos. Isso simplifica muito o design e melhora a composição e a legibilidade do código. Mais documentação específica do CuTe pode ser encontrada em seu diretório de documentação dedicado.
Além dos GEMMs, o CUTLASS implementa convolução de alto desempenho por meio do algoritmo GEMM implícito. GEMM implícito é a formulação de uma operação de convolução como um GEMM, aproveitando assim as vantagens do pipeline GEMM modular do CUTLASS. Isso permite que o CUTLASS construa convoluções reutilizando componentes GEMM altamente otimizados.
CUTLASS 3.6.0 é uma atualização do CUTLASS adicionando:
Hopper estruturado GEMM esparso.
FP16
8º PQ
INT8
TF32
Uma refatoração da API de convolução kernel::ConvUniversal do CUTLASS 3.x para alinhá-la com gemm::GemmUniversal . Agora a API de convolução 3.x não é mais considerada uma API beta.
Um GEMM de entrada mista aprimorado e uma implementação de tabela de pesquisa para o modo somente escala INT4 x FP8 .
Nós EVT para seleção Top-K e exemplo de softmax e GEMM usando-os.
Programmatic Dependent Launch (PDL) que aproveita um novo recurso Hopper para acelerar dois kernels consecutivos e suas documentações correspondentes.
Uma nova ferramenta de depuração, synclog, para despejar todos os eventos de sincronização de um kernel para um arquivo. Consulte a documentação do synclog para obter detalhes.
Um novo epílogo habilitado para TMA para GEMM agrupado que traz melhoria significativa de desempenho, bem como suporte EVT.
Um epílogo de array de ponteiros habilitado para SIMT.
Um novo cronograma de kernel Ping-Pong para GEMM agrupado e algumas outras otimizações.
Uma nova estratégia de instanciação para kernels do profiler CUTLASS junto com documentação aprimorada para o nível de instanciação no profiler CUTLASS.
Um novo suporte de hardware para comparações e cálculos de cutlass::bfloat16_t
Corrigido o uso de isnan no Windows para half_t .
Requisitos mínimos:
Arquitetura: Volta
Compilador: deve suportar pelo menos C++17
Versão do kit de ferramentas CUDA: 11.4
A partir do CUTLASS 3.0, o CUTLASS removeu o suporte para o seguinte:
Arquiteturas de GPU Maxwell e Pascal
Ubuntu 16.04
CUDA10.2
Versões da linguagem C++ inferiores a 17.
Consulte o CHANGELOG para uma lista detalhada de lançamentos e atualizações.
As primitivas CUTLASS são muito eficientes. Quando usados para construir kernels GEMM em todo o dispositivo, eles exibem desempenho máximo comparável ao cuBLAS para cálculos GEMM escalares. A figura acima mostra as melhorias contínuas de desempenho do CUTLASS em um NVIDIA H100 (arquitetura NVIDIA Hopper) desde o CUTLASS 3.1. CUTLASS 3.5.1 foi compilado com o CUDA 12.5u1 Toolkit. As operações do Tensor Core são implementadas usando as instruções mma e wgmma do CUDA.
Ao usar blocos de construção CUTLASS para construir kernels gemm implícitos em todo o dispositivo (Fprop, Dgrad e Wgrad), o desempenho do CUTLASS também é comparável ao cuDNN ao executar camadas Resnet-50 em um NVIDIA A100, conforme mostrado na figura acima. As operações do Tensor Core são implementadas usando a instrução mma do CUDA.
CUTLASS requer um compilador host C++17 e tem melhor desempenho quando construído com o CUDA 12.4 Toolkit . Também é compatível com 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 e CUDA 12.3.2.
Testamos os seguintes ambientes.
| Sistema operacional | Compilador |
|---|---|
| Ubuntu 18.04 | CGC 7.5.0 |
| Ubuntu 20.04 | CGC 10.3.0 |
| Ubuntu 22.04 | CGC 11.2.0 |
| Ubuntu 22.04 | Clang 10.0.0 |
| Ubuntu 22.04 | Clang 14.0.6 |
| Ubuntu 22.04 | Clang 17.0.6 |
| Janelas 10.0 | Visual Studio 2019 v16.11.27 |
Nota: O GCC 8.5.0 possui regressões conhecidas relacionadas a expressões dobradas e operadores sobrecarregados. Recomenda-se o uso do GCC 7.5.0 ou (preferencial) GCC >= 9.
CUTLASS é executado com sucesso nas seguintes GPUs NVIDIA e espera-se que seja eficiente nas GPUs NVIDIA baseadas nas arquiteturas Volta, Turing, Ampere, Ada e Hopper.
| GPU | Capacidade de computação CUDA | Kit de ferramentas CUDA mínimo exigido pelo CUTLASS-3 |
|---|---|---|
| GPU Tensor Core NVIDIA V100 | 7,0 | 11.4 |
| NVIDIA Titan V | 7,0 | 11.4 |
| NVIDIA GeForce RTX 2080TI, 2080, 2070 | 7,5 | 11.4 |
| NVIDIA T4 | 7,5 | 11.4 |
| GPU Tensor Core NVIDIA A100 | 8,0 | 11.4 |
| NVIDIA A10 | 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 |
| GPU Tensor Core NVIDIA H100 | 9,0 | 11.8 |
Em geral, o código PTX gerado para uma arquitetura alvo pode ser executado em arquiteturas futuras (ou seja, é compatível com versões futuras). No entanto, CUDA 12.0 introduziu o conceito de "recursos acelerados por arquitetura" cujo PTX não possui garantias de compatibilidade futura. Várias instruções Hopper PTX se enquadram nesta categoria de recursos acelerados por arquitetura e, portanto, requerem uma arquitetura de destino sm_90a (observe o "a" anexado). Para obter mais detalhes sobre esta e outras instruções aceleradas por arquitetura, consulte a documentação CUDA.
As informações da arquitetura alvo são passadas para o CUTLASS através do sinalizador cmake CUTLASS_NVCC_ARCHS . Para maximizar o desempenho do Hopper GH100, os usuários são obrigados a construir o CUTLASS com 90a como arquitetura alvo. Se um usuário construir acidentalmente um kernel que usa recursos SM90a (por exemplo, instruções Hopper Tensor Core), usando o alvo SM90 (observe a falta de "a"), com CUDA Toolkit 12 ou 11.8, espera-se que o kernel falhe com um tempo de execução erro.
cmake .. -DCUTLASS_NVCC_ARCHS="90a"
Consulte a documentação da funcionalidade para obter detalhes sobre quais kernels requerem quais arquiteturas de destino.
CUTLASS é descrito nos documentos a seguir e na documentação que acompanha o Doxygen.
Guia de início rápido - crie e execute o CUTLASS
Funcionalidade - resume a funcionalidade disponível no CUTLASS
GEMM eficiente em CUDA - descreve como os kernels GEMM podem ser implementados de forma eficiente em CUDA
Design CUTLASS 3.x - descreve o design CUTLASS 3.x, seus benefícios e como o CuTe nos permite escrever muito mais componentes combináveis
GEMM API 3.x - descreve o modelo GEMM do CUTLASS 3.x e os conceitos do modelo C++
GEMM API 2.x - descreve o modelo GEMM do CUTLASS 2.x e os conceitos do modelo C++
Convolução GEMM implícita - descreve a convolução 2-D e 3-D no CUTLASS
Organização do Código - descreve a organização e conteúdo do projeto CUTLASS
Terminologia - descreve os termos usados no código
Diretrizes de programação - diretrizes para escrever CUDA C++ moderno e eficiente
Tipos fundamentais - descreve classes C++ básicas usadas em CUTLASS para representar quantidades numéricas e matrizes
Layouts - descreve layouts de matrizes e tensores na memória
Tile Iterators - descreve conceitos C++ para iterar blocos de matrizes na memória
CUTLASS Profiler - aplicativo de criação de perfil orientado por linha de comando
Utilitários CUTLASS - modelos adicionais usados para facilitar o desenvolvimento rápido
Lançamento de kernel dependente - descreve um novo recurso no Hopper que permite a sobreposição de kernels dependentes no mesmo fluxo e como ele é usado no CUTLASS.
Também descrevemos a estrutura de um GEMM eficiente em nossa palestra na GPU Technology Conference 2018.
CUTLASS: Primitivos de software para álgebra linear densa em todos os níveis e escalas dentro de CUDA
Desenvolvendo Kernels CUDA para Levar Tensor Cores ao Limite Absoluto na NVIDIA A100
Acelerando a convolução com núcleos tensores em CUTLASS
Acelerando o gradiente de dados retroativos aumentando a utilização do Tensor Core no CUTLASS
CUTLASS: API Python, melhorias e NVIDIA Hopper
CUTLASS é uma biblioteca de modelos somente de cabeçalho e não precisa ser construída para ser usada por outros projetos. As aplicações clientes devem ter como alvo o diretório include/ do CUTLASS em seus caminhos de inclusão.
Testes de unidade, exemplos e utilitários CUTLASS podem ser criados com CMake. A versão mínima do CMake é fornecida no guia de início rápido. Certifique-se de que a variável de ambiente CUDACXX aponte para NVCC no CUDA Toolkit instalado em seu sistema.
$ exportar CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc Crie um diretório de construção dentro do projeto CUTLASS e execute o CMake. Por padrão, o CUTLASS construirá kernels para as versões 5.0, 6.0, 6.1, 7.0, 7.5, 8.0, 8.6, 8.9 e 9.0 da arquitetura CUDA. Para reduzir o tempo de compilação, você pode especificar as arquiteturas para construir o CUTLASS alterando a configuração do CMake CUTLASS_NVCC_ARCHS .
$ compilação mkdir && compilação cd $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # compila para a arquitetura Ampere da NVIDIA
No diretório build/ , compile e execute os testes de unidade CUTLASS construindo o destino test_unit com make.
Os testes unitários são organizados como vários binários que espelham os namespaces de nível superior do CUTLASS e podem ser executados em paralelo através do argumento de linha de comando -j do make.
$ faça test_unit -j ... ... ... [----------] Desmontagem do ambiente de teste global [==========] 946 testes de 57 casos de teste foram executados. (10.812 ms no total) [APROVADO] 946 testes.
Todos os testes devem ser aprovados em plataformas suportadas, embora o número exato de testes possa variar com o tempo.
CUTLASS é organizado como uma biblioteca somente de cabeçalho junto com utilitários, ferramentas, exemplos e testes de unidade. A documentação do Doxygen fornece uma lista completa de arquivos, classes e conceitos de modelo definidos no projeto CUTLASS.
Uma explicação detalhada da organização do código-fonte pode ser encontrada na documentação do CUTLASS, mas vários componentes principais estão resumidos abaixo.
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
Os exemplos do CUTLASS SDK aplicam modelos CUTLASS para implementar cálculos básicos.
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.
O diretório test/unit/ consiste em testes de unidade implementados com o Google Test que demonstram o uso básico dos componentes principais da API e testes completos dos cálculos CUTLASS GEMM.
As instruções para construir e executar os testes de unidade estão descritas no guia de início rápido.
O diretório tools/profiler/ contém um utilitário de linha de comando para iniciar cada um dos kernels GEMM. Pode ser construído da seguinte forma:
$faça cutlass_profiler -j16
Por padrão, apenas um tamanho de bloco é instanciado para cada tipo de dados, instrução matemática e layout. Para instanciar tudo, defina a seguinte variável de ambiente ao executar o CMake a partir de um diretório build/ vazio. Cuidado, isso resulta em dezenas de milhares de kernels e longos tempos de construção. Isso também resultaria em um tamanho binário grande e, em algumas plataformas, o vinculador falharia na construção da biblioteca. Portanto, é altamente recomendado gerar apenas um subconjunto de kernels conforme demonstrado na subseção abaixo.
$ cmake.. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=todos ... $faça cutlass_profiler -j16
Para compilar estritamente um kernel ou um pequeno conjunto de kernels, uma lista delimitada por vírgulas de nomes de kernel com caracteres curinga pode ser usada para reduzir o conjunto de kernels. Os exemplos a seguir mostram a construção de exatamente um ou um subconjunto de kernels para as arquiteturas NVIDIA Ampere e Turing:
Para compilar um subconjunto de kernels Tensor Core GEMM com acumulação FP32 e entrada FP16 visando a arquitetura NVIDIA Ampere e Turing, use a linha de comando cmake abaixo:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8 ... $faça cutlass_profiler -j16
Um exemplo de linha de comando para criar o perfil de um subconjunto de kernels GEMM do Tensor Core é o seguinte:
./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096
...
=============================
ID do problema: 1
Provedor: CUTLASS
Tipo de operação: gemm
Operação: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
Status: Sucesso
Verificação: LIGADO
Disposição: Aprovado
reference_device: Aprovado
cuBLAS: Aprovado
Argumentos: --gemm_kind=universal --m=3456 --n=4096 --k=4096 --A=f16:coluna --B=f16:linha --C=f32:coluna --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
Bytes: 118489088 bytes
FLOPs: 115992428544 flops
Tempo de execução: 1,55948ms
Memória: 70,7616 GiB/s
Matemática: 74378,8 GFLOP/s
=============================
...Para compilar um kernel SGEMM direcionado às arquiteturas NVIDIA Ampere e Turing, use a linha de comando cmake abaixo:
$ cmake. ... $faça cutlass_profiler -j16
Um exemplo de linha de comando para criação de perfil de kernel SGEMM CUDA único é o seguinte:
$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096
=============================
ID do problema: 1
Provedor: CUTLASS
Tipo de operação: gemm
Operação: cutlass_simt_sgemm_128x128_8x2_nn_align1
Status: Sucesso
Verificação: LIGADO
Disposição: Aprovado
cuBLAS: Aprovado
Argumentos: --m=3456 --n=4096 --k=4096 --A=f32:coluna --B=f32:coluna --C=f32:coluna --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
Bytes: 180355072 bytes
FLOPs: 115992428544 flops
Tempo de execução: 6,73655 ms
Memória: 24.934 GiB/s
Matemática: 17218,4 GFLOP/s
=============================Para compilar um subconjunto de kernels de convolução do núcleo Tensor implementando propagação direta (fprop) com acumulação FP32 e entrada FP16 visando a arquitetura NVIDIA Ampere e Turing, use a linha de comando cmake abaixo:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16 ... $faça cutlass_profiler -j16
Um exemplo de linha de comando para criar o perfil de um subconjunto de kernels de convolução do Tensor Core é o seguinte:
$ ./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
...
=============================
ID do problema: 1
Provedor: CUTLASS
Tipo de operação: conv2d
Operação: cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc
Status: Sucesso
Verificação: LIGADO
Disposição: Aprovado
reference_device: Aprovado
Argumentos: --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 --Ativação=f16:nhwc --Filter=f16:nhwc --Saída=f32:nhwc
--conv_mode=cruzado --iterator_algorithm=otimizado --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
Bytes: 1130659840 bytes
FLOPs: 118482796544 flops
Tempo de execução: 0,711496ms
Memória: 1479,99 GiB/s
Matemática: 166526 GFLOP/s
=============================
...Para compilar e executar um kernel de convolução CUDA Core implementando propagação direta (fprop) com acumulação F32 e entrada FP32 visando a arquitetura NVIDIA Ampere e Turing, use a linha de comando cmake abaixo:
$ cmake. ... $faça cutlass_profiler -j16
Exemplo de linha de comando para criação de perfil de um kernel de convolução 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
=============================
ID do problema: 1
Provedor: CUTLASS
Tipo de operação: conv2d
Operação: cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
Status: Sucesso
Verificação: LIGADO
Disposição: Aprovado
reference_device: Aprovado
Argumentos: --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 --Ativação=f32:nhwc --Filtro=f32:nhwc --Saída=f32:nhwc
--conv_mode=cruzado --iterator_algorithm=otimizado --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
Bytes: 2055798784 bytes
FLOPs: 118482796544 flops
Tempo de execução: 7,34266 ms
Memória: 260,752 GiB/s
Matemática: 16136,2 GFLOP/s
=============================Por favor, siga os links para mais exemplos do CMake sobre a compilação seletiva de kernels CUTLASS:
Exemplos GEMM CMake
Exemplos implícitos de CMake de convolução GEMM
Mais detalhes sobre o CUTLASS Profiler são descritos aqui.
CUTLASS é lançado pela NVIDIA Corporation como software de código aberto sob a "Nova" licença BSD de 3 cláusulas.
A lista oficial de desenvolvedores e contribuidores do CUTLASS está disponível aqui: CONTRIBUIDORES.
Copyright (c) 2017 - 2024 NVIDIA CORPORATION E AFILIADAS. Todos os direitos reservados. Identificador de licença SPDX: Cláusula 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.
