Sobre a programabilidade do AWS Trainium e Inferen…

Grupo Shoslh
novembro 1, 2024
4:19 pm
Sem comentários

Acelerando o treinamento de modelos de IA/ML com operadores personalizados – Secção 4

12 minutos de leitura

13 horas detrás

Foto de Ágata Bres sobre Remover respingo

Nesta postagem, continuamos nossa exploração das oportunidades de otimização do tempo de realização de cargas de trabalho de aprendizagem de máquina (ML) por meio do desenvolvimento de operadores personalizados. Desta vez, nos concentramos nas ferramentas fornecidas pelo SDK do AWS Neuron para desenvolver e executar novos kernels em Treinamento AWS e Inferência AWS. Com o rápido desenvolvimento dos componentes do protótipo de insignificante nível (por exemplo, camadas de atenção) impulsionando a revolução da IA, a programabilidade dos aceleradores usados para treinar e executar modelos de ML é crucial. Os chips de IA dedicados, em privado, devem oferecer uma escolha válida às estruturas de desenvolvimento de GPU de uso universal (GPGPU), amplamente utilizadas e altamente impactantes, uma vez que CUDA e Tritão.

Em postagens anteriores (por exemplo, cá e cá), exploramos a oportunidade de erigir e executar modelos de ML em chips de IA personalizados da AWS usando o devotado SDK do AWS Neuron. Em sua versão mais recente do SDK (versão 2.20.0), a AWS apresentou o Interface do Kernel do Neurônio (NKI) para desenvolver kernels personalizados para NeuronCore-v2o acelerador subjacente alimentando ambos Treinamento e Inferência2. A interface NKI se junta a outra API que permite NeuronCore-v2 programabilidade, Operadores C++ personalizados do Neuron. Neste post iremos explorar ambas as oportunidades e demonstrá-las em ação.

Isenções de responsabilidade

É importante ressaltar que esta postagem não deve ser vista uma vez que um substituto da postagem solene Documentação do SDK do AWS Neuron. No momento em que oriente item foi escrito, as APIs do Neuron SDK para desenvolvimento de kernel personalizado estavam na versão beta e podem mudar no momento em que você ler isto. Os exemplos que compartilhamos são unicamente para fins demonstrativos. Não fazemos nenhuma reivindicação quanto à sua otimização, robustez, espaço ou precisão. Por obséquio, não veja nossa menção a quaisquer plataformas, ferramentas, APIs, etc., uma vez que um endosso para seu uso. As melhores escolhas para qualquer projeto dependem das especificidades do caso de uso em questão e garantem investigação e estudo adequadas.

Embora a lista de modelos de ML suportados pelo Neuron SDK esteja crescendo continuamente, algumas operações permanecem sem suporte ou implementadas de forma aquém do ideal. Ao expor APIs para personalização do kernel do Neuron, o SDK permite que os desenvolvedores criem e/ou otimizem as operações de insignificante nível necessárias, aumentando significativamente a oportunidade de realização de cargas de trabalho de ML no Trainium e no Inferentia.

Conforme discutido em nossas postagens anteriores desta série, aproveitar totalmente o poder desses chips de IA requer uma compreensão detalhada de sua arquitetura de insignificante nível.

A arquitetura do núcleo do neurônio

A documentação do NKI inclui um seção dedicada sobre o design da arquitetura do NeuronCore-v2 e suas implicações no desenvolvimento de operadores customizados. É importante ressaltar que existem muitas diferenças entre os núcleos do Neuron e seus equivalentes aceleradores de IA (por exemplo, GPUs e TPUs). A otimização para núcleos do Neuron requer um conjunto individual de estratégias e habilidades.

Semelhante a outros chips de IA dedicados, o NeuronCore-v2 inclui vários motores de aceleraçãocada um deles especializado em realizar certos tipos de cálculos. Os motores podem funcionar de forma assíncrona e em paralelo. O Compilador de neurônios é responsável por transformar modelos de ML em operações de insignificante nível e otimizar a escolha do mecanismo de computação para cada um.

O Motor tensor é perito em multiplicação de matrizes. O Vetor e Escalar ambos os motores operam em tensores com o motor Vector especializado em operações de redução e o motor Escalar em funções não lineares. GpSimd é um mecanismo de uso universal capaz de executar programas C/C++ arbitrários. Observe que enquanto o NKI interface expõe aproximação a todos os quatro mecanismos de computação, operadores C++ personalizados são projetados especificamente para GpSimd.

Mais detalhes sobre as capacidades de cada mecanismo podem ser encontrados na documentação da arquitetura. Outrossim, o Arquitetura do conjunto de instruções NKI (ISA) a documentação fornece detalhes sobre os mecanismos nos quais diferentes operações de insignificante nível são executadas.

Outro paisagem importante do chip Neuron é sua arquitetura de memória. Um dispositivo Neuron inclui três tipos de memória, HBM, SBUF e PSUM. Uma compreensão íntima das capacidades e capacidades de cada um é crucial para o desenvolvimento ideal do kernel.

Dada a visão universal da arquitetura, você pode concluir que o desenvolvimento do kernel do Neuron requer basta conhecimento. Embora isso possa ser verdade para a geração de kernels totalmente otimizados que aproveitam todos os recursos do núcleo do Neuron, nosso objetivo é provar a acessibilidade, o valor e o potencial das APIs de kernel personalizadas do Neuron — mesmo para desenvolvedores não especialistas.

O NKI interface é uma API em nível Python que expõe o uso dos principais mecanismos de computação e recursos de memória do Neuron para desenvolvedores de ML. O Introdução ao NKI O guia detalha as instruções de forma e fornece um pouso suave com um kernel simples, “olá mundo”. O Padrão de Programação NKI O guia detalha os três estágios de um kernel NKI típico (carregamento de entradas, realização de operações nos mecanismos de computação e armazenamento de saídas) e apresenta o NKI Tile e Operações baseadas em blocos. O Tutoriais NKI provar uma variedade de aplicativos de modelo do kernel NKI, com cada um apresentando novas APIs e recursos principais do NKI. Dada a suposta otimização dos kernels de modelo, uma estratégia provável para o desenvolvimento de novos kernels poderia ser 1) identificar uma modelo que seja semelhante à operação que você deseja implementar e logo 2) usá-la uma vez que risca de base e refiná-la iterativamente e ajustá-la para inferir a funcionalidade específica que você precisa.

O Manual de referência da API NKI detalha a API Python para desenvolvimento do kernel. Com uma sintaxe e semiologia semelhantes a Tritão e NumPyo QUAL linguagem definição visa maximizar a acessibilidade e facilidade de uso. No entanto, é importante notar que o desenvolvimento do kernel NKI está restringido às operações definidas no NKI livraria, que (no momento em que oriente livro foi escrito) são menos e mais restritas do que em bibliotecas uma vez que Tritão e NumPy.

Exemplo de brinquedo – um kernel GIOU

Porquê em nossos posts anteriores, avaliamos o uso do NKI construindo uma implementação customizada do Intersecção Generalizada sobre União (GIOU) operação em um par de lotes de caixas de ingressão. Porquê GIOU envolve operações em pixels, usamos o experiência núcleo do Programação NKI guia uma vez que ponto de referência e incorporou o uso do NKI indexação avançada de tensores em nossa implementação. Para facilitar a depuração em um envolvente de CPU, também adicionamos opções para executar o código usando o what.simulate_kernel e nki.linguagem.device_print.html APIs.

import torch
import neuronxcc.nki as nki
import neuronxcc.nki.language as nl
import numpy as npsimulate = False
try:
# if torch libraries are installed assume that we are running on Neuron
import torch_xla.core.xla_model as xm
import torch_neuronx
from torch_neuronx import nki_jit
device = xm.xla_device()
# empty implementation 
def debug_print(*args, **kwargs):
pass
except:
# if torch libraries are not installed assume that we are running on CPU
# and program script to use nki simulation
simulate = True
nki_jit = nki.trace
debug_print = nl.device_print
device = 'cpu'
@nki_jit
def giou_kernel(preds_ptr,
targets_ptr,
output_ptr):
epsilon = 1e-5
TILE_M = nl.tile_size.pmax  # 128
TILE_N = nl.tile_size.psum_fmax  # 512
TILE_N_OUT = TILE_N // 4
p_1, Sua visita nos ajuda a continuar oferecendo o melhor para você! = preds_ptr.shape
t_1, t_2 = targets_ptr.shape
o_1, o_2 = output_ptr.shape
#  verify input
# batch size must be multiple of 128
assert p_1 % TILE_M == 0
assert p_1 == t_1
assert p_1 == o_1
# num boxes box *4 must be multiple of 512
assert p_2 % TILE_N == 0
assert p_2 == t_2
assert p_2 // 4 == o_2
num_tiles_m = p_1 // TILE_M
num_tiles_n = p_2 // TILE_N
# Generate tensors for advanced indexing
i_p = nl.arange(TILE_M)[:, None]
i_f = nl.arange(TILE_N // 4)[None, :]
i_f_0 = (4 * i_f)
i_f_1 = (4 * i_f + 1)
i_f_2 = (4 * i_f + 2)
i_f_3 = (4 * i_f + 3)
# Use affine_range to loop over tiles
for m in nl.affine_range(num_tiles_m):
for n in nl.affine_range(num_tiles_n):
# Load input data from HBM
preds = nl.load(preds_ptr[m * TILE_M:(m + 1) * TILE_M,
n * TILE_N:(n + 1) * TILE_N])
targets = nl.load(targets_ptr[m * TILE_M:(m + 1) * TILE_M,
n * TILE_N:(n + 1) * TILE_N])
debug_print('preds', preds)
preds_left = preds[i_p, i_f_0]
preds_top = preds[i_p, i_f_1]
preds_right = preds[i_p, i_f_2]
preds_bottom = preds[i_p, i_f_3]
gt_left = targets[i_p, i_f_0]
gt_top = targets[i_p, i_f_1]
gt_right = targets[i_p, i_f_2]
gt_bottom = Sua visita nos ajuda a continuar oferecendo o melhor para você! i_f_3]
# Compute the area of each box
area1 = (preds_right - preds_left) * (preds_bottom - preds_top)
area2 = (gt_right - gt_left) * (gt_bottom - gt_top)
# Compute the intersection
left = nl.maximum(preds_left, gt_left)
top = nl.maximum(preds_top, gt_top)
right = nl.minimum(preds_right, gt_right)
bottom = nl.minimum(preds_bottom, gt_bottom)
inter_w = nl.maximum(right - left, 0)
inter_h = nl.maximum(bottom - top, 0)
inter_area = inter_w * inter_h
union_area = area1 + area2 - inter_area
iou_val = inter_area / nl.maximum(union_area, epsilon)
# Compute the smallest enclosing box
enclose_left = nl.minimum(preds_left, gt_left)
enclose_top = nl.minimum(preds_top, gt_top)
enclose_right = nl.maximum(preds_right, gt_right)
enclose_bottom = nl.maximum(preds_bottom, gt_bottom)
enclose_w = nl.maximum(enclose_right - enclose_left, 0)
enclose_h = nl.maximum(enclose_bottom - enclose_top, 0)
enclose_area = enclose_w * enclose_h
# Compute GIOU
delta_area = (enclose_area - union_area)
enclose_area = nl.maximum(enclose_area, epsilon)
giou = iou_val - delta_area / enclose_area
# Store results
nl.store(output_ptr[m * TILE_M:(m + 1) * TILE_M,
n * TILE_N_OUT:(n + 1) * TILE_N_OUT],
giou)

Para executar nosso kernel GIOU, geramos dois lotes de caixas aleatórias e os alimentamos em nossa função:

# generate random data in np
np.random.seed(0)
batch_size = 1024
n_boxes = 256
img_size = 256
boxes = []for i in range(2):
# Randomly generate box sizes and positions
box_sizes = np.random.randint(1, img_size, size=(batch_size,n_boxes,2))
top_left = np.random.randint(0, img_size-1, size=(batch_size,n_boxes,2))
bottom_right = np.clip(top_left + box_sizes, 0, img_size - 1)
# Concatenate top-left and bottom-right coordinates
rand_boxes = np.concatenate((top_left, bottom_right), axis=2)
boxes.append(rand_boxes.astype(np.float32))
out = np.empty((batch_size, n_boxes), np.float32)
# convert tensors to PyTorch
t_boxes_0 = torch.tensor(boxes[0]).to(device)
t_boxes_1 = torch.tensor(boxes[1]).to(device)
t_out = torch.tensor(out).to(device)
if simulate:
# the simulation API requires numpy input
nki.simulate_kernel(giou_kernel, 
boxes[0].reshape((batch_size, -1)),
boxes[1].reshape((batch_size, -1)),
out)
else:
giou_kernel(t_boxes_0.view((batch_size, -1)),
t_boxes_1.view((batch_size, -1)),
t_out)

Para determinar o desempenho de nosso kernel NKI, iremos compará-lo com a seguinte implementação ingênua de GIOU em PyTorch:

def torch_giou(boxes1, boxes2):
# loosely based on torchvision generalized_box_iou_loss code
epsilon = 1e-5# Compute areas of both sets of boxes
area1 = (boxes1[...,2]-boxes1[...,0])*(boxes1[...,3]-boxes1[...,1])
area2 = (boxes2[...,2]-boxes2[...,0])*(boxes2[...,3]-boxes2[...,1])
# Corners of intersection
lt = torch.max(boxes1[..., :2], boxes2[..., :2])
rb = torch.min(boxes1[..., 2:], boxes2[..., 2:])
# Width and height of intersection
wh = (rb - lt).clamp(min=0)
# Area of the intersection
inter = wh[..., 0] * wh[..., 1]
# Union of the two boxes
union = area1 + area2 - inter
iou = inter / union.clamp(epsilon)
# Corners of enclosing box
lti = torch.min(boxes1[..., :2], boxes2[..., :2])
rbi = torch.max(boxes1[..., 2:], boxes2[..., 2:])
# Width and height of the enclosing box
whi = (rbi - lti).clamp(min=0)
# Area of the enclosing box
areai = (whi[..., 0] * whi[..., 1]).clamp(epsilon)
return iou - (areai - union) / areai

Usamos o seguinte utilitário de benchmarking para confrontar o desempenho de tempo de realização de nossas duas funções:

import time
def benchmark(f, warmup_iters=20, ntrials: int = 100):
def run(*args, **kwargs):
# warmup
for _ in range(warmup_iters):
f(*args, **kwargs)
start_time = time.time()
for _ in range(ntrials):
f(*args, **kwargs)
end_time = time.time()
# Calculate average time per iteration
avg_time = (end_time - start_time) / ntrials
return avg_timereturn run
avg_time = benchmark(torch_giou)(t_boxes_0, t_boxes_1)
print(f'torch_giou: {avg_time}')
avg_time = benchmark(giou_kernel)(t_boxes_0.view((batch_size, -1)),
t_boxes_1.view((batch_size, -1)),
t_out)
print(f'giou_kernel: {avg_time}')

Envolvente de realização

Executamos nosso script em um Amazon EC2 inf2.xlarge instância (contendo dois Núcleos de neurônios e quatro vCPUs). Usamos a versão mais recente do AMI de aprendizagem profundo para Neuron disponível no momento da redação deste item, “Deep Learning AMI Neuron (Ubuntu 22.04) 20241027”, com AWS Neuron 2.20.1 e PyTorch 2.1.

Resultados

Nosso kernel GIOU personalizado demonstrou um tempo de realização médio de 0,211 milissegundos em conferência com 0,293, totalizando um aumento de desempenho de 39%. Tenha em mente que estes resultados são exclusivos do nosso exemplo de brinquedo. Outros operadores, principalmente aqueles que incluem multiplicações de matrizes (e utilizam o mecanismo Tensor), provavelmente exibirão resultados comparativos diferentes.

Otimizando o Desempenho do Kernel NKI

O próximo passo no desenvolvimento do nosso kernel — além do escopo deste post — seria explorar o desempenho do kernel GIOU usando o devotado Perfilador de neurônios para identificar gargalos e otimizar nossa implementação. Por obséquio veja o Guia de desempenho NKI para mais detalhes.

O segundo método para gerar um kernel Neuron personalizado é erigir um operador C++ para o Motor GpSimd. Levante método é descrito no Guia do desenvolvedor de operadores C++ personalizados do Neuron e demonstrado no Operadores C++ personalizados do Neuron em MLP e Otimização de desempenho de operadores C++ personalizados do Neuron tutoriais.

Os operadores Neuron Custom C++ apresentam uma oportunidade para “fusão de kernel” no mecanismo GpSimd, facilitando a combinação de várias operações de insignificante nível em uma única realização de kernel. Essa abordagem pode reduzir significativamente a sobrecarga associada a: 1) carregamento de vários kernels individuais e 2) transferência de dados entre diferentes regiões de memória.

Exemplo de brinquedo – um kernel GIOU C++

No conjunto de código aquém, implementamos um operador C++ GIOU para Neuron e o salvamos em um registro chamado giou.cpp. Nosso kernel usa o Acessador TCM para otimizar o desempenho de leitura e gravação de memória e aplica o multicore forma para use todos os oito processadores internos do GpSimd.

#include 
#include 
#include 
#include 
#include // input boxes of shape 1024x256x4
// output scores of shape 1024x256
torch::Tensor giou(const torch::Tensor& t_pred, 
const torch::Tensor& t_target) {
size_t num_samples = t_pred.sizes()[0];
size_t num_boxes = t_pred.sizes()[1];
torch::Tensor t_out = get_dst_tensor();
// get the number of GpSimd processors (8 in NeuronCoreV2) 
uint32_t cpu_count = get_cpu_count();
// get índice of current processor
uint32_t cpu_id = get_cpu_id();
// divide the batch size into 8 partitions 
uint32_t partition = num_samples / cpu_count;
// use tcm buffers to load and write data
size_t tcm_in_size = num_boxes*4;
size_t tcm_out_size = num_boxes;
float *tcm_pred = (float*)torch::neuron::tcm_malloc(
sizeof(float)*tcm_in_size);
float *tcm_target = (float*)torch::neuron::tcm_malloc(
sizeof(float)*tcm_in_size);
float *tcm_output = (float*)torch::neuron::tcm_malloc(
sizeof(float)*tcm_in_size);
auto t_pred_tcm_acc = t_pred.tcm_accessor();
auto t_target_tcm_acc = t_target.tcm_accessor();
auto t_out_tcm_acc = t_out.tcm_accessor();
// iterate over each of the entries in the partition
for (size_t i = 0; i < partition; i++) {
// load the pred and target boxes into lugar memory
t_pred_tcm_acc.tensor_to_tcm(tcm_pred,
partition*cpu_id + i*tcm_in_size,
tcm_in_size);
t_target_tcm_acc.tensor_to_tcm(tcm_target,
partition*cpu_id + i*tcm_in_size,
tcm_in_size);
// iterate over each of the boxes in the entry
for (size_t j = 0; j < num_boxes; j++) {
const float epsilon = 1e-5;
const float* box1 = &tcm_pred[j * 4];
const float* box2 = &tcm_target[j * 4];
// Compute area of each box
float area1 = (box1[2] - box1[0]) * (box1[3] - box1[1]);
float area2 = (box2[2] - box2[0]) * (box2[3] - box2[1]);
// Compute the intersection
float left = std::max(box1[0], box2[0]);
float top = std::max(box1[1], box2[1]);
float right = std::min(box1[2], box2[2]);
float bottom = std::min(box1[3], box2[3]);
float inter_w = std::max(right - left, 0.f);
float inter_h = std::max(bottom - top, 0.f);
float inter_area = inter_w * inter_h;
// Compute the union area
float union_area = area1 + area2 - inter_area;
// IoU
float iou_val = inter_area / std::max(union_area, epsilon);
// Compute the smallest enclosing box
float enclose_left = std::min(box1[0], box2[0]);
float enclose_top = std::min(box1[1], box2[1]);
float enclose_right = std::max(box1[2], box2[2]);
float enclose_bottom = std::max(box1[3], box2[3]);
float enclose_w = std::max(enclose_right - enclose_left, 0.f);
float enclose_h = std::max(enclose_bottom - enclose_top, 0.f);
float enclose_area = std::max(enclose_w * enclose_h, epsilon);
float result = iou_val - (enclose_area-union_area)/enclose_area;
tcm_output[j] = result;
}
// write the giou scores of all boxes in the current entry
t_out_tcm_acc.tcm_to_tensor(tcm_output,
partition*cpu_id + i*tcm_out_size,
tcm_out_size);
}
torch::neuron::tcm_free(tcm_pred);
torch::neuron::tcm_free(tcm_target);
return t_out;
}

Exigimos um separado forma.cpp registro que define o formato de saída de nossa função GIOU e registra nosso operador personalizado na livraria Neuron:

#include 
#include 
#include 
#include "torchneuron/register.h"torch::Tensor giou_shape(torch::Tensor boxes1, torch::Tensor boxes2) {
torch::Tensor t_out = torch::zeros({boxes1.sizes()[0],
boxes1.sizes()[1]},
torch::kFloat);
return t_out;
}
NEURON_LIBRARY(my_ops, m) {
m.def("giou", &giou_shape, "giou");
}

O erigir.py script compila o operador C++ e o expõe uma vez que uma API Python:

import os
import torch_neuronx
from torch_neuronx.xla_impl import custom_opcustom_op.load(
name='giou',
compute_srcs=['giou.cpp'],
shape_srcs=['shape.cpp'],
build_directory=os.getcwd(),
multicore=True,
verbose=True
)

O script de compilação gera um libgiou.so livraria contendo a implementação do nosso operador C++ GIOU. No conjunto de código aquém carregamos a livraria e medimos o desempenho do nosso kernel personalizado usando o utilitário de benchmarking definido supra:

from torch_neuronx.xla_impl import custom_op
custom_op.load_library('libgiou.so')avg_time = benchmark(torch.ops.my_ops.giou)(t_boxes_0, t_boxes_1)
print(f'C++ giou: {avg_time}')

Envolvente de realização

Usamos o mesmo envolvente Neuron de nossos experimentos NKI para compilar e testar nosso kernel C++. Por obséquio, observe o etapas de instalação que são necessários para o desenvolvimento de operadores C++ customizados.

Resultados

Nosso kernel C++ GIOU demonstrou um tempo de realização médio de 0,061 milissegundos — quase cinco vezes mais rápido que nossa implementação básica. Presumivelmente, isso é resultado da “fusão do kernel”, conforme discutido supra.

A tábua aquém resume os resultados do tempo de realização de nossos experimentos.

Tenha em mente que esses resultados são específicos do exemplo do brinquedo e do envolvente de realização usado neste estudo. Os resultados comparativos de outros kernels podem ser muito diferentes — dependendo do intensidade em que eles podem aproveitar os mecanismos de computação internos do núcleo do Neuron.

A tábua aquém resume algumas das diferenças que observamos entre os dois métodos de personalização do kernel AWS Neuron.

Confrontação entre ferramentas de customização do kernel (por responsável)

Por meio de sua interface Python de basta nível, as APIs NKI expõem o poder dos mecanismos de aceleração Neuron aos desenvolvedores de ML de maneira alcançável e fácil de usar. A livraria de operadores personalizados C++ de insignificante nível permite uma programação ainda maior, mas é limitada ao Motor GpSimd. Ao combinar efetivamente as duas ferramentas, os desenvolvedores podem aproveitar totalmente os recursos da arquitetura AWS Neuron.

Com a revolução da IA em pleno curso, muitas empresas estão desenvolvendo novos chips avançados de IA para atender à crescente demanda por computação. Embora os anúncios públicos muitas vezes destaquem o desempenho de tempo de realização, a economia de custos e a eficiência energética desses chips, vários recursos básicos são essenciais para tornar esses chips e suas pilhas de software verdadeiramente viáveis para o desenvolvimento de ML. Esses recursos incluem ferramentas robustas de depuração, estudo de desempenho e utilitários de otimização, programabilidadee muito mais.

Nesta postagem nos concentramos nos utilitários disponíveis para programar os aceleradores de IA desenvolvidos internamente pela AWS Treinamento e Inferênciae demonstraram seu uso na geração de operações de ML personalizadas. Essas ferramentas capacitam os desenvolvedores a otimizar o desempenho de seus modelos de ML nos chips de IA da AWS e abrem novas oportunidades para inovação e originalidade.

Tags:

Crédito: Manancial Original