برنامه نویسی

در مورد برنامه ریزی AWS Trainium و Inferentia

ek3nk4r 2024-11-03

0 0 خواندن این مطلب 26 دقیقه زمان میبرد

در مورد برنامه ریزی AWS Trainium و Inferentia

پیشنهاد ویژه

خرید فالوور واقعی خرید لایک اینستاگرام خرید ویو اینستاگرام خرید فالوور اینستاگرام

Summarize this content to 400 words in Persian Lang

تسریع آموزش مدل AI/ML با اپراتورهای سفارشی – قسمت 4

عکس آگاتا برس در Unsplash

در این پست ما به بررسی فرصت‌های بهینه‌سازی زمان اجرا بارهای کاری یادگیری ماشین (ML) از طریق توسعه اپراتور سفارشی ادامه می‌دهیم. این بار، ما بر ابزارهای ارائه شده توسط AWS Neuron SDK برای توسعه و اجرای هسته های جدید در AWS Trainium و AWS Inferentia تمرکز می کنیم. با توسعه سریع اجزای مدل سطح پایین (به عنوان مثال، لایه های توجه) که انقلاب هوش مصنوعی را هدایت می کند، برنامه ریزی شتاب دهنده های مورد استفاده برای آموزش و اجرای مدل های ML بسیار مهم است. تراشه‌های اختصاصی هوش مصنوعی، به‌ویژه، باید جایگزین مناسبی برای چارچوب‌های توسعه پرکاربرد و بسیار تاثیرگذار GPU (GPGPU) مانند CUDA و Triton باشند.

در پست‌های قبلی (به عنوان مثال، اینجا و اینجا) فرصت ساخت و اجرای مدل‌های ML بر روی تراشه‌های AI سفارشی ساخت AWS با استفاده از AWS Neuron SDK را بررسی کردیم. در جدیدترین نسخه SDK (نسخه 2.20.0)، AWS رابط هسته عصبی (NKI) را برای توسعه هسته‌های سفارشی برای NeuronCore-v2، شتاب‌دهنده زیربنایی که هم Trainium و هم Inferentia2 را تامین می‌کند، معرفی کرد. رابط NKI به API دیگری ملحق می شود که برنامه نویسی NeuronCore-v2 را فعال می کند، Neuron Custom C++ Operators. در این پست هر دو فرصت را بررسی کرده و آنها را در عمل نشان خواهیم داد.

سلب مسئولیت

نکته مهم این است که این پست نباید به عنوان جایگزینی برای اسناد رسمی AWS Neuron SDK در نظر گرفته شود. در زمان نوشتن این مقاله، Neuron SDK APIs برای توسعه هسته سفارشی در بتا است و ممکن است با خواندن این مطلب تغییر کند. نمونه هایی که ما به اشتراک می گذاریم، فقط برای اهداف نمایشی در نظر گرفته شده اند. ما هیچ ادعایی در مورد بهینه بودن، استحکام، دوام یا دقت آنها نداریم. لطفاً اشاره ما به پلتفرم‌ها، ابزارها، APIها و غیره را به عنوان تأییدی برای استفاده از آن‌ها تلقی نکنید. بهترین انتخاب برای هر پروژه به ویژگی های مورد استفاده در دست بستگی دارد و تحقیقات و تجزیه و تحلیل مناسب را تضمین می کند.

توسعه هسته های سفارشی برای هسته های عصبی

اگرچه لیست مدل های ML پشتیبانی شده توسط Neuron SDK به طور مداوم در حال افزایش است، برخی از عملیات ها یا پشتیبانی نمی شوند یا به صورت غیربهینه اجرا می شوند. با افشای APIها برای سفارشی‌سازی هسته نورون، SDK به توسعه‌دهندگان اجازه می‌دهد تا عملیات سطح پایین مورد نیاز خود را ایجاد و/یا بهینه‌سازی کنند و فرصت اجرای بارهای کاری ML در Trainium و Inferentia را بسیار افزایش می‌دهد.

همانطور که در پست های قبلی ما در این سری بحث شد، استفاده کامل از قدرت این تراشه های هوش مصنوعی مستلزم درک دقیق معماری سطح پایین آنها است.

معماری هسته نورون

اسناد NKI شامل بخش اختصاصی در مورد طراحی معماری NeuronCore-v2 و پیامدهای آن بر توسعه اپراتور سفارشی است. نکته مهم این است که تفاوت‌های زیادی بین هسته‌های نورون و همتایان شتاب‌دهنده هوش مصنوعی آن‌ها (مانند GPU و TPU) وجود دارد. بهینه‌سازی هسته‌های نورون به مجموعه‌ای از استراتژی‌ها و مهارت‌ها نیاز دارد.

مشابه سایر تراشه های اختصاصی هوش مصنوعی، NeuronCore-v2 شامل چندین موتور شتاب دهنده داخلی است که هر کدام در انجام انواع خاصی از محاسبات تخصص دارند. موتورها می توانند به صورت ناهمزمان و موازی کار کنند. کامپایلر نورون مسئول تبدیل مدل های ML به عملیات سطح پایین و بهینه سازی انتخاب موتور محاسباتی برای هر یک است.

موتور Tensor در ضرب ماتریس تخصص دارد. موتورهای Vector و Scalar هر دو بر روی تانسورها با موتور Vector متخصص در عملیات کاهش و موتور Scalar در توابع غیر خطی کار می کنند. GpSimd یک موتور همه منظوره است که قادر به اجرای برنامه های دلخواه C/C++ است. توجه داشته باشید که در حالی که رابط NKI دسترسی به هر چهار موتور محاسباتی را نشان می دهد، اپراتورهای C++ سفارشی به طور خاص برای GpSimd طراحی شده اند.

جزئیات بیشتر در مورد قابلیت های هر موتور را می توان در مستندات معماری یافت. علاوه بر این، مستندات NKI Instruction Set Architecture (ISA) جزئیاتی را در مورد موتورهایی که عملیات های مختلف سطح پایین بر روی آنها اجرا می شود، ارائه می دهد.

یکی دیگر از جنبه های مهم تراشه نورون، معماری حافظه آن است. دستگاه Neuron شامل سه نوع حافظه HBM، SBUF و PSUM است. درک نزدیک از ظرفیت ها و قابلیت های هر یک برای توسعه بهینه هسته بسیار مهم است.

با توجه به نمای کلی معماری، ممکن است نتیجه بگیرید که توسعه هسته نورون به تخصص بالایی نیاز دارد. اگرچه این ممکن است برای ایجاد هسته‌های کاملاً بهینه‌سازی شده که از تمام قابلیت‌های هسته نورون استفاده می‌کنند صادق باشد، هدف ما نشان دادن قابلیت دسترسی، ارزش و پتانسیل APIهای هسته سفارشی Neuron – حتی برای توسعه‌دهندگان غیرمتخصص است.

هسته های سفارشی NKI

رابط NKI یک API در سطح پایتون است که استفاده از موتورهای محاسباتی هسته نورون و منابع حافظه را در اختیار توسعه دهندگان ML قرار می دهد. راهنمای شروع NKI دستورالعمل‌های راه‌اندازی را به تفصیل شرح می‌دهد و یک فرود نرم را با هسته ساده و “سلام جهان” ارائه می‌دهد. راهنمای مدل برنامه‌نویسی NKI سه مرحله یک هسته معمولی NKI (بارگیری ورودی‌ها، اجرای عملیات روی موتورهای محاسباتی و ذخیره خروجی‌ها) را شرح می‌دهد و عملیات NKI Tile و Tile-based را معرفی می‌کند. آموزش‌های NKI انواع برنامه‌های نمونه هسته NKI را نشان می‌دهند که هر کدام APIها و قابلیت‌های اصلی NKI را معرفی می‌کنند. با توجه به بهینه فرضی هسته‌های نمونه، یک استراتژی ممکن برای توسعه هسته‌های جدید می‌تواند این باشد که 1) نمونه‌ای را شناسایی کنید که مشابه عملیاتی است که می‌خواهید اجرا کنید و سپس 2) از آن به عنوان خط پایه استفاده کنید و به طور مکرر آن را اصلاح و تنظیم کنید. به عملکرد خاصی که نیاز دارید دست پیدا کنید.

NKI API Reference Manual API Python برای توسعه هسته را شرح می دهد. با نحو و معنایی که مشابه Triton و NumPy است، تعریف زبان NKI به حداکثر رساندن دسترسی و سهولت استفاده است. با این حال، توجه به این نکته مهم است که توسعه هسته NKI به عملیات تعریف شده در کتابخانه NKI محدود می شود، که (تا زمان نگارش این مقاله) نسبت به کتابخانه هایی مانند Triton و NumPy کمتر و محدودتر هستند.

نمونه اسباب بازی – هسته GIOU

همانطور که در پست‌های قبلی‌مان، استفاده از NKI را با ایجاد یک پیاده‌سازی سفارشی از عملیات تقاطع عمومی بر روی اتحادیه (GIOU) روی یک جفت دسته از جعبه‌های ورودی ارزیابی می‌کنیم. از آنجایی که GIOU شامل عملیات مبتنی بر پیکسل است، ما از هسته *exp* از راهنمای برنامه نویسی NKI به عنوان نقطه مرجع استفاده کردیم و استفاده از نمایه سازی تانسور پیشرفته NKI را در پیاده سازی خود به کار بردیم. برای تسهیل اشکال‌زدایی در یک محیط CPU، ما همچنین گزینه‌هایی برای اجرای کد با استفاده از APIهای nki.simulate_kernel و nki.language.device_print.html اضافه کردیم.

import torch
import neuronxcc.nki as nki
import neuronxcc.nki.language as nl
import numpy as np

simulate = 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, p_2 = 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 = targets[i_p, 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)

برای اجرای هسته GIOU خود، دو دسته از جعبه های تصادفی تولید می کنیم و آنها را به تابع خود تغذیه می کنیم:

# 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)

برای ارزیابی عملکرد هسته NKI خود، آن را با اجرای ساده GIOU زیر در 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

ما از ابزار سنجش زیر برای مقایسه عملکرد زمان اجرا دو عملکرد خود استفاده می کنیم:

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_time

return 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}’)

محیط زمان اجرا

ما اسکریپت خود را روی یک نمونه Amazon EC2 inf2.xlarge (شامل دو هسته Neuron و چهار vCPU) اجرا کردیم. ما از جدیدترین نسخه Deep Learning AMI برای Neuron که در زمان نگارش این مقاله موجود بود، “Deep Learning AMI Neuron (Ubuntu 22.04) 20241027” با AWS Neuron 2.20.1 و PyTorch 2.1 استفاده کردیم.

نتایج

هسته GIOU سفارشی ما میانگین زمان اجرا 0.211 میلی ثانیه را در مقایسه با 0.293 نشان می دهد که به میزان 39٪ افزایش عملکرد را نشان می دهد. به خاطر داشته باشید که این نتایج برای نمونه اسباب بازی ما منحصر به فرد است. سایر عملگرها، به ویژه آنهایی که شامل ضرب ماتریس هستند (و از موتور Tensor استفاده می کنند) احتمالاً نتایج مقایسه ای متفاوتی را نشان می دهند.

بهینه سازی عملکرد هسته NKI

گام بعدی در توسعه هسته ما – فراتر از محدوده این پست – تجزیه و تحلیل عملکرد هسته GIOU با استفاده از Neuron Profiler اختصاصی به منظور شناسایی تنگناها و بهینه سازی پیاده سازی ما خواهد بود. لطفاً راهنمای عملکرد NKI را برای جزئیات بیشتر ببینید.

Neuron Custom C++ Operators

روش دوم برای ایجاد یک هسته Neuron سفارشی، ساخت یک اپراتور C++ برای موتور GpSimd است. این روش در Neuron Custom C++ Operators Developer Guide و در Neuron Custom C++ Operators در MLP و Neuron Custom C++ Operators Performance Optimization نمایش داده شده است.

Neuron Custom C++ Operators با تسهیل ترکیب چندین عملیات سطح پایین در یک هسته واحد، فرصتی را برای “تلفیق هسته” در موتور GpSimd ارائه می دهد. این رویکرد می‌تواند هزینه‌های سربار مرتبط با موارد زیر را به میزان قابل توجهی کاهش دهد: 1) بارگیری چندین هسته جداگانه، و 2) انتقال داده بین مناطق مختلف حافظه.

مثال اسباب بازی – هسته GIOU C++

در بلوک کد زیر یک عملگر C++ GIOU را برای Neuron پیاده سازی می کنیم و آن را در فایلی به نام ذخیره می کنیم giou.cpp. هسته ما از دسترسی TCM برای بهینه سازی عملکرد خواندن و نوشتن حافظه استفاده می کند و تنظیم *چند هسته ای * را برای استفاده از هر هشت پردازنده داخلی GpSimd اعمال می کند.

#include <stdint.h>
#include <stdlib.h>
#include <torch/torch.h>
#include <neuron/neuron-utils.hpp>
#include <algorithm>

// input boxes of shape 1024x256x4
// output scores of shape 1024×256
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 index 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 local memory
t_pred_tcm_acc.tensor_to_tcm<float>(tcm_pred,
partition*cpu_id + i*tcm_in_size,
tcm_in_size);
t_target_tcm_acc.tensor_to_tcm<float>(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<float>(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;
}

ما نیاز به یک جداگانه داریم shape.cpp فایلی که شکل خروجی تابع GIOU ما را تعریف می کند و عملگر سفارشی ما را با کتابخانه Neuron ثبت می کند:

#include <stdint.h>
#include <stdlib.h>
#include <torch/torch.h>
#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”);
}

را build.py اسکریپت عملگر C++ را کامپایل می کند و آن را به عنوان یک API پایتون نمایش می دهد:

import os
import torch_neuronx
from torch_neuronx.xla_impl import custom_op

custom_op.load(
name=”giou”,
compute_srcs=[‘giou.cpp’],
shape_srcs=[‘shape.cpp’],
build_directory=os.getcwd(),
multicore=True,
verbose=True
)

اسکریپت کامپایل یک کتابخانه *libgiou.so * ایجاد می کند که شامل اجرای عملگر C++ GIOU ما است. در بلوک کد زیر، کتابخانه را بارگذاری می کنیم و عملکرد هسته سفارشی خود را با استفاده از ابزار سنجش تعریف شده در بالا اندازه گیری می کنیم:

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}’)

محیط زمان اجرا

ما از همان محیط Neuron از آزمایش‌های NKI خود برای کامپایل و آزمایش هسته C ++ خود استفاده کردیم. لطفاً به مراحل نصبی که برای توسعه اپراتور C++ سفارشی لازم است توجه کنید.

نتایج

هسته C++ GIOU ما میانگین زمان اجرا 0.061 میلی ثانیه را نشان داد – تقریباً پنج برابر سریعتر از اجرای خط پایه ما. این احتمالاً نتیجه “همجوشی هسته” است، همانطور که در بالا مورد بحث قرار گرفت.

نتیجه گیری

جدول زیر نتایج زمان اجرا آزمایش های ما را خلاصه می کند.
میانگین زمان پیاده سازی های مختلف GIOU (کمتر بهتر است) – توسط نویسنده

لطفاً به خاطر داشته باشید که این نتایج مختص نمونه اسباب بازی و محیط زمان اجرا مورد استفاده در این مطالعه است. نتایج مقایسه سایر هسته ها ممکن است بسیار متفاوت باشد – بسته به درجه ای که آنها می توانند از موتورهای محاسباتی داخلی هسته نورون استفاده کنند.

جدول زیر برخی از تفاوت‌هایی را که بین دو روش سفارشی‌سازی هسته نورون AWS مشاهده کردیم، خلاصه می‌کند.

مقایسه بین ابزار سفارشی سازی هسته (توسط نویسنده)

از طریق رابط سطح بالای پایتون، API های NKI، قدرت موتورهای شتاب دهنده نورون را به شکلی در دسترس و کاربرپسند در اختیار توسعه دهندگان ML قرار می دهند. کتابخانه سطح پایین C++ Custom Operators برنامه ریزی حتی بیشتر را امکان پذیر می کند، اما به موتور GpSimd محدود می شود. با ترکیب موثر هر دو ابزار، توسعه دهندگان می توانند به طور کامل از قابلیت های معماری AWS Neuron بهره ببرند.

خلاصه

با رشد کامل انقلاب هوش مصنوعی، بسیاری از شرکت‌ها در حال توسعه تراشه‌های پیشرفته هوش مصنوعی برای پاسخگویی به تقاضای رو به رشد برای محاسبات هستند. در حالی که اطلاعیه های عمومی اغلب عملکرد این تراشه ها، صرفه جویی در هزینه و بهره وری انرژی را برجسته می کند، چندین قابلیت اصلی ضروری است تا این تراشه ها و پشته های نرم افزاری آنها واقعاً برای توسعه ML قابل دوام باشند. این قابلیت ها شامل ابزارهای اشکال زدایی قوی، ابزارهای تجزیه و تحلیل عملکرد و بهینه سازی، قابلیت برنامه ریزی و غیره است.

در این پست، ما بر ابزارهای موجود برای برنامه‌نویسی شتاب‌دهنده‌های هوش مصنوعی بومی AWS، Trainium و Inferentia تمرکز کردیم و استفاده از آنها را در ساخت عملیات ML سفارشی نشان دادیم. این ابزارها توسعه دهندگان را قادر می سازد تا عملکرد مدل های ML خود را بر روی تراشه های هوش مصنوعی AWS بهینه کنند و فرصت های جدیدی را برای نوآوری و خلاقیت باز کنند.

فهرست مطالب

تسریع آموزش مدل AI/ML با اپراتورهای سفارشی – قسمت 4

عکس آگاتا برس در Unsplash

سلب مسئولیت

توسعه هسته های سفارشی برای هسته های عصبی

معماری هسته نورون

هسته های سفارشی NKI

نمونه اسباب بازی – هسته GIOU

import torch
import neuronxcc.nki as nki
import neuronxcc.nki.language as nl
import numpy as np

simulate = 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, p_2 = 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 = targets[i_p, 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)

برای اجرای هسته GIOU خود، دو دسته از جعبه های تصادفی تولید می کنیم و آنها را به تابع خود تغذیه می کنیم:

# 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)

برای ارزیابی عملکرد هسته NKI خود، آن را با اجرای ساده GIOU زیر در 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

ما از ابزار سنجش زیر برای مقایسه عملکرد زمان اجرا دو عملکرد خود استفاده می کنیم:

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_time

    return 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}')

محیط زمان اجرا

نتایج

بهینه سازی عملکرد هسته NKI

Neuron Custom C++ Operators

مثال اسباب بازی – هسته GIOU C++

در بلوک کد زیر یک عملگر C++ GIOU را برای Neuron پیاده سازی می کنیم و آن را در فایلی به نام ذخیره می کنیم giou.cpp. هسته ما از دسترسی TCM برای بهینه سازی عملکرد خواندن و نوشتن حافظه استفاده می کند و تنظیم *چند هسته ای * را برای استفاده از هر هشت پردازنده داخلی GpSimd اعمال می کند.

#include <stdint.h>
#include <stdlib.h>
#include <torch/torch.h>
#include <neuron/neuron-utils.hpp>
#include <algorithm>

// 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 index 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 local memory
    t_pred_tcm_acc.tensor_to_tcm<float>(tcm_pred,
                                        partition*cpu_id + i*tcm_in_size,
                                        tcm_in_size);
    t_target_tcm_acc.tensor_to_tcm<float>(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<float>(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;
}

ما نیاز به یک جداگانه داریم shape.cpp فایلی که شکل خروجی تابع GIOU ما را تعریف می کند و عملگر سفارشی ما را با کتابخانه Neuron ثبت می کند:

#include <stdint.h>
#include <stdlib.h>
#include <torch/torch.h>
#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");
}

را build.py اسکریپت عملگر C++ را کامپایل می کند و آن را به عنوان یک API پایتون نمایش می دهد:

import os
import torch_neuronx
from torch_neuronx.xla_impl import custom_op

custom_op.load(
    name="giou",
    compute_srcs=['giou.cpp'],
    shape_srcs=['shape.cpp'],
    build_directory=os.getcwd(),
    multicore=True,
    verbose=True
)

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}')