Versnelling van AI/ML-modeltraining met aangepaste operators — Deel 2

Foto door Jas Rolyn op Unsplash

Volgens de Griekse mythologie zou Triton, een god van de zee, het zeewater kalmeren of in beroering brengen door zijn schelp te gebruiken om de getijden en golven te beheersen. In één verhaal in het bijzonder wordt Triton afgebeeld als iemand die zijn krachten gebruikte om de Argonauten door bijzonder gevaarlijk zeewater te leiden. In dit bericht doen we op dezelfde manier een beroep op Triton voor navigatie door complexe reizen, hoewel we dit keer verwijzen naar de Triton-taal en -compiler voor het schrijven van deep learning (DL)-kernels en naar onze reizen door de wereld van AI/ML-ontwikkeling.

Dit is een vervolg op een eerdere post over het onderwerp van het versnellen van AI/ML-toepassingen met aangepaste operators waarin we het potentieel voor prestatieoptimalisatie hebben aangetoond door aangepaste CUDA-kernels te ontwikkelen. Een van onze bedoelingen was om de toegankelijkheid van aangepaste kernelontwikkeling en de kansen die het biedt, zelfs voor niet-deskundige CUDA-ontwikkelaars, te benadrukken. Er zijn echter uitdagingen voor CUDA-ontwikkeling die voor sommigen onoverkomelijk kunnen blijken. Ten eerste zijn veel moderne AI/ML-ontwikkelaars goed thuis in Python, maar voelen ze zich misschien niet op hun gemak bij het ontwikkelen in C++. Bovendien vereist het afstemmen van een CUDA-kernel om volledig te profiteren van de mogelijkheden van de GPU een diepgaand begrip van de onderliggende HW-architectuur en kan het een niet-triviale hoeveelheid werk vergen. Dit geldt met name als u wilt dat uw kernel optimaal draait op verschillende GPU-architecturen. Een groot deel van de complexiteit is het gevolg van het ‘thread-based’ ontwikkelingsmodel van CUDA, waarbij de ontwikkelaar verantwoordelijk is voor het ontwerpen en optimaliseren van alle elementen van de GPU-kerneldraden, inclusief alle details met betrekking tot het gebruik van GPU-geheugen, thread-concurrency, TensorCore-planning en nog veel meer.

De kracht van Triton

De Triton-bibliotheek heeft als doel om GPU-kernelontwikkeling op twee manieren te democratiseren en te vereenvoudigen. Ten eerste biedt het een API voor het bouwen van aangepaste operators in Python (in plaats van C++). Ten tweede maakt het kernelontwikkeling mogelijk op de blok niveau (in plaats van het threadniveau) waardoor alle problemen met betrekking tot het optimaliseren van de prestaties worden geabstraheerd en geautomatiseerd binnenin CUDA-threadblokken. In plaats van de moeizame stappen van het programmeren van de details van de threadaanroep, inclusief de complexiteit van geheugenbeheer, planning van on-chip acceleratie-engines, thread-synchronisatie, etc., kunnen kernelontwikkelaars erop vertrouwen dat Triton het allemaal voor hen doet. Een belangrijk bijproduct van de API-abstractie op hoog niveau van Tritons programmeermodel is dat het de last vermindert van het moeten afstemmen van de kernel voor meerdere verschillende GPU-typen en architecturen.

Natuurlijk heeft het Triton-programmeermodel, zoals meestal het geval is bij het uplevelen van een API, ook zijn nadelen. Sommige kernels kunnen profiteren van de thread-level controle die mogelijk is gemaakt door CUDA (bijvoorbeeld, ze kunnen profiteren van de conditionele uitvoeringsstroom die in onze vorige post is besproken). Andere kernels vereisen mogelijk een zeer gespecialiseerde en delicate behandeling om piekprestaties te bereiken en kunnen lijden onder het geautomatiseerde resultaat van de Triton-compiler. Maar zelfs in gevallen als deze, waarbij de ontwikkeling van een CUDA-kernel uiteindelijk vereist kan zijn, kan de mogelijkheid om snel en eenvoudig een tijdelijke Triton-kernel te maken de ontwikkeling aanzienlijk vergemakkelijken en de productiviteit verhogen.

Voor meer informatie over de motivaties achter Triton en de details van het programmeermodel, zie de aankondiging van Triton, de officiële documentatie van Triton en het originele whitepaper van Triton.

Vrijwaring

Net als in ons vorige bericht is het onze bedoeling om een ​​eenvoudige demonstratie te geven van de mogelijkheden die Triton biedt. Zie dit bericht niet als een vervanging voor de officiële Triton-documentatie of de bijbehorende tutorials. We gebruiken hetzelfde gezichtsdetectiemodel als in ons vorige bericht als basis voor onze demonstratie en voeren onze experimenten uit in dezelfde Google Cloud-omgeving: een g2-standard-16 VM (met een enkele L4 GPU) met een speciale deep learning VM-image en PyTorch 2.4.0. Zoals eerder doen we geen moeite om onze voorbeelden te optimaliseren en/of hun robuustheid, duurzaamheid of nauwkeurigheid te verifiëren. Er moet worden opgemerkt dat hoewel we onze experimenten uitvoeren op een PyTorch-model en op een NVIDIA GPU, de ontwikkeling van de Triton-kernel wordt ondersteund door aanvullende frameworks en onderliggende HW’s.

Triton als onderdeel van Torch-compilatie

In eerdere berichten (bijvoorbeeld hier) hebben we het gebruik van PyTorch-compilatie en de mogelijke impact op runtime-prestaties gedemonstreerd. De standaardcompiler die door de torch.compiler wordt gebruikt, is TorchInductor, die sterk afhankelijk is van Triton-kernels voor zijn GPU-versnelling. Daarom lijkt het alleen maar passend dat we onze Triton-verkenning beginnen met het beoordelen van de automatische Triton-ondersteunde optimalisatie die torch.compile biedt. Het onderstaande codeblok bevat dezelfde forward pass van het gezichtsdetectiemodel dat we in ons vorige bericht hebben geïntroduceerd, samen met de gecompileerde GIOU-verliesfunctie. Om het kort te houden, hebben we een deel van de ondersteunende code weggelaten. Raadpleeg ons vorige bericht voor de volledige implementatie.

def loss_with_padding(pred, targets):
    mask = (targets(...,3) > 0).to(pred.dtype)
    total_boxes = mask.sum()
    loss = generalized_box_iou(targets, pred)
    masked_loss = loss*mask
    loss_sum = masked_loss.sum()
    return loss_sum/torch.clamp(total_boxes, 1)


device = torch.device("cuda:0")
model = torch.compile(Net()).to(device).train()
loss_fn = torch.compile(loss_with_padding)

# forward portion of training loop wrapped with profiler object
with torch.profiler.profile(
   schedule=torch.profiler.schedule(wait=5, warmup=5, active=10, repeat=1)
) as prof:
    for step, data in enumerate(train_loader):

        with torch.profiler.record_function('copy data'):
            images, boxes = data_to_device(data, device)
            torch.cuda.synchronize(device)

        with torch.profiler.record_function('forward'):
            with torch.autocast(device_type="cuda", dtype=torch.bfloat16):
                outputs = model(images)
            torch.cuda.synchronize(device)

        with torch.profiler.record_function('calc loss'):
            loss = loss_fn(outputs, boxes)
            torch.cuda.synchronize(device)
        prof.step()
        if step > 30:
            break

    # filter and print profiler results
    event_list = prof.key_averages()
    for i in range(len(event_list) - 1, -1, -1):
        if event_list(i).key not in ('forward', 'calc loss', 'copy data'):
            del event_list(i)
    print(event_list.table())

De prestatieresultaten (gemiddeld over meerdere runs) worden hieronder weergegeven:

-------------  ------------  ------------
         Name     CPU total  CPU time avg
-------------  ------------  ------------
    copy data      56.868ms       5.687ms
      forward        1.329s     132.878ms
    calc loss       8.282ms     828.159us
-------------  ------------  ------------

Bedenk dat de gemiddelde tijd van de originele verliesfunctie (op opgevulde invoer) 1,844 ms was. De prestatieverbetering die voortvloeit uit torchcompilatie is dus groter dan 2X(!!).

De Triton-kernels die automatisch door torch.compile worden gegenereerd, kunnen daadwerkelijk worden bekeken door de TORCH_LOGS-omgevingsvariabele in te stellen, zoals uitgelegd in deze PyTorch-zelfstudie. Sommigen hebben zelfs voorgesteld om deze kernels te gebruiken als startpunt voor Triton-ontwikkeling (zie bijvoorbeeld hier). In onze ervaring kunnen deze kernels echter lastig te ontcijferen zijn.

In het volgende gedeelte proberen we de resultaten van PyTorch-compilatie verder te verbeteren door een GIOU Triton-kernel te implementeren.

Een aangepaste Triton-kernel maken

Een geweldige plek om uw Triton-ontwikkelingsreis te beginnen is met de officiële Triton-zelfstudies. De zelfstudies worden geïntroduceerd in toenemende volgorde van complexiteit, waarbij elke zelfstudie een of meer van Tritons unieke functies uitbreidt. Onze GIOU Triton-kernel lijkt het meest op het meest basale voorbeeld van vectoroptelling. Net als in onze CUDA-implementatie wijzen we een blok toe aan elk sample in de invoerbatch en programmeren we het om te werken op alle bounding boxes in het sample. Let op het gebruik van tl.load en tl.store voor het lezen en schrijven van gegevens van en naar het geheugen, evenals het gebruik van vectoriële rekenkunde door blokprogramma’s.

import triton
import triton.language as tl

@triton.jit
def giou_kernel(preds_ptr,
                    targets_ptr,
                    output_ptr,
                    valid_ptr,
                    BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(axis=0)
    box_id = tl.arange(0, BLOCK_SIZE)
    
    box_offsets = pid * BLOCK_SIZE + box_id
    
    preds_left = tl.load(preds_ptr + 0 + 4 * box_offsets)
    preds_top = tl.load(preds_ptr + 1 + 4 * box_offsets)
    preds_right = tl.load(preds_ptr + 2 + 4 * box_offsets)
    preds_bottom = tl.load(preds_ptr + 3 + 4 * box_offsets)
    
    gt_left = tl.load(targets_ptr + 0 + 4 * box_offsets)
    gt_top = tl.load(targets_ptr + 1 + 4 * box_offsets)
    gt_right = tl.load(targets_ptr + 2 + 4 * box_offsets)
    gt_bottom = tl.load(targets_ptr + 3 + 4 * box_offsets)
    
    epsilon = 1e-5
    
    # 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 = tl.maximum(preds_left, gt_left)
    top = tl.maximum(preds_top, gt_top)
    right = tl.minimum(preds_right, gt_right)
    bottom = tl.minimum(preds_bottom, gt_bottom)
    
    inter_w = right - left
    inter_h = bottom - top
    inter_area = inter_w * inter_h
    
    union_area = area1 + area2 - inter_area
    
    iou_val = inter_area / tl.maximum(union_area, epsilon)
    
    # Compute the smallest enclosing box
    enclose_left = tl.minimum(preds_left, gt_left)
    enclose_top = tl.minimum(preds_top, gt_top)
    enclose_right = tl.maximum(preds_right, gt_right)
    enclose_bottom = tl.maximum(preds_bottom, gt_bottom)
    
    enclose_w = enclose_right - enclose_left
    enclose_h = enclose_bottom - enclose_top
    enclose_area = enclose_w * enclose_h
    
    # Compute GIOU
    delta_area = (enclose_area - union_area)
    enclose_area = tl.maximum(enclose_area, epsilon)
    giou = iou_val - delta_area / enclose_area
    
    # Store results
    tl.store(output_ptr + (box_offsets),
             tl.where(gt_bottom > 0, giou, 0))
    tl.store(valid_ptr + (box_offsets), gt_bottom > 0)


def loss_with_triton(pred, targets):
    batch_size = pred.shape(0)
    n_boxes = pred.shape(1)
    
    # convert to float32 (remove to keep original dtypes)
    pred = pred.to(torch.float32)
    targets = targets.to(torch.float32)

    # allocate output tensors
    output = torch.empty_strided(pred.shape(0:2), 
                                 stride=(n_boxes,1),
                                 dtype = pred.dtype,
                                 device = pred.device)
    valid = torch.empty_strided(pred.shape(0:2),
                                stride=(n_boxes,1),
                                dtype = torch.bool,
                                device = pred.device)
 
    # call Triton kernel
    giou_kernel((batch_size,))(pred, targets, output, valid,
                               BLOCK_SIZE=n_boxes)

    total_valid = valid.sum()
    loss_sum = output.sum()
    return loss_sum/total_valid.clamp(1)

De resultaten van het draaien met onze Triton-kernel worden hieronder vastgelegd. Hoewel het iets slechter is dan in ons vorige experiment, kan dit het resultaat zijn van extra optimalisaties die door torch.compile zijn uitgevoerd.

-------------  ------------  ------------
         Name     CPU total  CPU time avg
-------------  ------------  ------------
    copy data      57.089ms       5.709ms
      forward        1.338s     133.771ms
    calc loss       8.908ms     890.772us
-------------  ------------  ------------

Naar aanleiding van de aanbeveling van PyTorch’s documentatie over het gebruik van Triton kernels, beoordelen we de prestaties van onze kernel verder, dit keer in combinatie met PyTorch-compilatie. De resultaten (gemiddeld over meerdere runs) zijn iets beter dan het automatisch gecompileerde verlies van ons eerste experiment.

-------------  ------------  ------------
         Name     CPU total  CPU time avg
-------------  ------------  ------------
    copy data      57.008ms       5.701ms
      forward        1.330s     132.951ms
    calc loss       7.189ms     718.869us
-------------  ------------  ------------

Bij het ontwikkelen van onze aangepaste GIOU CUDA-kernel merkten we de overhead op van het converteren van de invoertensoren naar float32, en de noodzaak om onze kernel te verbeteren om verschillende invoertypen te ondersteunen om deze conversie te vermijden. In het geval van onze Triton-kernel kan dit vrij eenvoudig worden bereikt door de conversiebewerkingen te verwijderen. De aangepaste kernel wordt automatisch gegenereerd (JIT-gecompileerd) met de originele typen.

-------------  ------------  ------------
         Name     CPU total  CPU time avg
-------------  ------------  ------------
    copy data      57.034ms       5.703ms
      forward        1.325s     132.456ms
    calc loss       6.219ms     621.950us
-------------  ------------  ------------

Onze uiteindelijke resultaten komen overeen met de CUDA-kernelresultaten die we in ons vorige bericht zagen.

Resultaten

De volgende tabel vat de resultaten van onze experimenten samen. De resultaten werden gemiddeld over meerdere runs vanwege enige variantie die we observeerden. We hebben de resultaten van onze aangepaste CUDA-kernel uit onze vorige post opgenomen ter referentie. Houd er rekening mee dat de vergelijkende resultaten waarschijnlijk sterk kunnen variëren op basis van de details van de kernel en de runtime-omgeving.

Samenvatting van gemiddelde verlieslooptijden (per auteur)

Hoewel ons eerste Triton-kernelexperiment resulteerde in een lagere prestatie vergeleken met onze aangepaste CUDA-operator, konden we de snelheid ervan evenaren door compilatie toe te passen en de gegevenstypeconversies te verwijderen.

Deze bevindingen komen overeen met wat je van Triton zou verwachten: enerzijds impliceert de API-abstractie op hoog niveau een zeker verlies van controle over de low-level flow, wat kan resulteren in verminderde runtime-prestaties. Anderzijds stellen de (relatieve) eenvoud en kracht van de API’s gebruikers in staat om de prestatiekloof te dichten door functies veel gemakkelijker te implementeren dan in CUDA.

Je zou een sterk argument kunnen aanvoeren dat de Triton-kernel die we hebben gekozen om te evalueren is wat de documentatie zou aanduiden als “beschamend parallel”, d.w.z. samengesteld uit element-gewijze bewerkingen, en dat is als zodanig een vreselijke kernel om de waarde van Triton op te demonstreren. Sterker nog, een complexer programma, dat geavanceerder geheugenbeheer, planning, synchronisatie, etc. vereist, kan nodig zijn om de volledige kracht van Triton te laten zien.

Volgende stappen

Er zijn meerdere extra stappen nodig om onze taak te voltooien. Deze omvatten het afstemmen van onze aangepaste kernel en het implementeren van de backward-functie.

1. Kerneloptimalisatie

Hoewel Triton veel van de low-level kernel-optimalisatie abstraheert, blijven er veel controles over die een grote impact kunnen hebben op de runtime-prestaties. Deze omvatten de grootte van elk blok, het aantal thread warps dat moet worden gebruikt (zoals gedemonstreerd in de softmax-tutorial) en hoe L2-geheugen wordt benaderd (zie de matrixvermenigvuldigingstutorial voor een voorbeeld van sissend). Triton bevat een autotuning-functie voor het optimaliseren van de keuze van hyperparameters (zoals gedemonstreerd in de matrixvermenigvuldigingstutorial en in het PyTorch Triton-voorbeeld). Hoewel we autotuning uit ons voorbeeld hebben weggelaten, is het een essentiële stap in de ontwikkeling van de Triton-kernel.

2. Implementatie van de backward pass

We hebben ons voorbeeld beperkt tot alleen de forward pass van de GIOU-verliesfunctie. Een volledige oplossing zou vereisen dat er ook een kernel wordt gemaakt voor de backward pass (zoals gedemonstreerd in de tutorial over laagnormalisatie). Dit is meestal iets ingewikkelder dan de forward pass. Je kunt je afvragen waarom de high-level kernel development API die door Triton wordt blootgelegd deze uitdaging niet aanpakt door automatische differentiatie te ondersteunen. Zoals blijkt, om redenen die buiten het bereik van dit bericht vallen (zie bijvoorbeeld hier), is automatische differentiatie van aangepaste kernels extreem moeilijk te implementeren. Niettemin zou dit een absolute killer van een functie voor Triton zijn en we kunnen alleen maar hopen dat dit op enig moment in de toekomst wordt ondersteund.

Samenvatting

Triton is met gemak een van de belangrijkste en meest impactvolle AI/ML-bibliotheken van de afgelopen jaren. Hoewel het moeilijk is om de hoeveelheid innovatie en vooruitgang die het mogelijk heeft gemaakt op het gebied van AI te beoordelen, zijn de voetafdrukken overal te vinden — van de kernimplementatie van PyTorch 2 en de afhankelijkheden ervan, tot de gespecialiseerde aandachtslagen binnen de geavanceerde LLM-modellen die langzaam ons dagelijks leven perforeren.

Triton’s populariteit is te danken aan zijn innovatieve programmeermodel voor kernelontwikkeling. Triton was ooit beperkt tot het domein van CUDA-experts, maar maakt het maken van aangepaste DL-primitieven toegankelijk voor elke Python-ontwikkelaar.

In dit bericht hebben we alleen de oppervlakte van Triton en zijn mogelijkheden aangeraakt. Bekijk zeker de online documentatie van Triton en andere bronnen om meer te weten te komen.

De kracht van Triton ontketenen: GPU-kerneloptimalisatie in Python onder de knie krijgen werd oorspronkelijk gepubliceerd in Towards Data Science op Medium. Mensen zetten het gesprek voort door dit verhaal te markeren en erop te reageren.