Vue d'ensemble d'Ascend C
Ascend C est un langage de programmation dédié, conçu par Huawei pour exploiter pleinement le potentiel de calcul des processeurs IA de la série Ascend. Basé sur la syntaxe standard du C++, il étend celui-ci avec des bibliothèques d'API spécifiques. En tant que composant clé de l'architecture de calcul hétérogène CANN (Compute Architecture for Neural Networks), il permet le développement de fonctions noyau (kernel functions) s'exécutant sur les AI Cores, offrant des similitudes avec le développement CUDA pour les GPU. L'objectif principal est de fournir un contrôle direct sur l'architecture matérielle tout en conservant la familiarité du C++, pour répondre aux exigences de performance des applications d'IA.
Les principes de conception d'Ascend C reposent sur : le support natif des normes C/C++ pour réduire la courbe d'apprentissage ; une abstraction multi-niveaux qui équilibre efficacité de développement et contrôle matériel ; l'introduction de paradigmes de programmation parallèle ; et l'offre de techniques de débogage jumeau CPU/NPU. Ces caractéristiques en font un outil puissant et accessible pour le développement d'opérateurs IA haute performance.
D'un point de vue architectural, Ascend C s'appuie sur le processeur IA Ascend, basé sur l'architecture Da Vinci, qui intègre plusieurs AI Cores. Chaque core possède des unités de calcul scalaires, vectorielles et matricielles. La pile logicielle CANN assure la liaison entre le matériel et les frameworks d'apprentissage profond, optimisant l'utilisation des ressources de calcul via une co-conception logicielle-matérielle. Les opérateurs développés en Ascend C sont compilés et exécutés sur le processeur IA Ascend.
Comparé au C++ standard, Ascend C se distingue par : un mécanisme de gestion de mémoire spécifique distinguant mémoire globale et mémoire locale ; des qualificateurs de fonction pour préciser l'environnement d'exécution ; des API de transfert de données et de calcul enrichies ; et un paradigme de programmation structuré décomposant l'implémentation d'un opérateur en parties gérables.
Extensions de Syntaxe Propres à Ascend C
Ascend C étend la syntaxe du C++ standard avec des éléments tels que les qualificateurs d'espace d'exécution de fonction, les qualificateurs d'espace d'adresse et l'opérateur d'appel de noyau.
Qualificateurs d'Espace d'Exécution de Fonction
Ces qualificateurs indiquent où une fonction s'exécute (hôte ou périphérique) et par qui elle peut être appelée. Trois principaux qualificateurs existent :
__global__: Désigne une fonction noyau. Elle s'exécute sur le périphérique (AI Core), ne peut être appelée que depuis l'hôte, doit retournervoidet n'être pas membre d'une classe. Son appel sur l'hôte utilise la syntaxe asynchrone<<<>>>. Une synchronisation explicite est nécessaire pour attendre sa complétion.__aicore__: Désigne une fonction s'exécutant sur un AI Core. Elle ne peut être appelée que par une fonction__global__ou une autre fonction__aicore__. Utile pour encapsuler des fonctionnalités communes.__host__: Désigne une fonction s'exécutant côté hôte (souvent implicite). Elle ne peut être appelée que par des fonctions hôte.
Exemple d'utilisation des qualificateurs :
__aicore__ void traitement_accelere() {}
__global__ __aicore__ void fonction_noyau() { traitement_accelere(); }
int programme_hote() {}
Qualificateurs d'Espace d'Adresse
Utilisés dans les déclarations de variables, ils spécifient la région de mémoire où l'objet est alloué. L'AI Core dispose de mémoires hiérarchiques indépendantes.
__gm__: Désigne un objet alloué en mémoire globale (Global Memory) côté périphérique. C'est une mémoire de grande capacité mais relativement lente, servant à stocker les données d'entrée et de sortie. Pour les pointeurs en paramètre de fonction noyau, on utilise la macroGM_ADDRdéfinie comme#define GM_ADDR __gm__ uint8_t*pour uniformiser l'expression.- Espace
private: Espace par défaut pour la plupart des variables locales. Il correspond à la mémoire locale (Local Memory) de l'AI Core, rapide mais de capacité limitée, idéale pour les résultats intermédiaires et les données fréquemment accédées.
Mécanisme d'Appel de Noyau
L'appel d'une fonction noyau depuis l'hôte utilise l'opérateur <<<...>>>. Cet opérateur prend trois paramètres clés :
- blockDim : Le nombre de noyaux (AI Cores) à invoquer pour une exécution parallèle.
- l2ctrl : Paramètre de contrôle du cache L2, généralement
nullptr. - stream : Descripteur de flux pour gérer les opérations asynchrones et la synchronisation.
Exemple d'appel :
// Appel asynchrone de la fonction noyau 'operation_addition' sur 8 coeurs
operation_addition<<<8, nullptr, stream>>>(donnees_in_x, donnees_in_y, donnees_out_z);
L'appel est asynchrone : le contrôle retourne immédiatement à l'hôte. Pour forcer l'attente de la complétion des noyaux, on utilise une fonction de synchronisation comme aclrtSynchronizeStream.
Architecture de l'API Ascend C
Ascend C fournit une vaste bibliothèque d'API, divisée en API de base et API de haut niveau, permettant aux développeurs de choisir le niveau d'abstraction adapté à leurs besoins.
API de Base et Leurs Fonctions
Les API de base abstraient les capacités matérielles fondamentales. Elles incluent :
- API de Calcul : Pour les unités scalaires, vectorielles et matricielles. Elles offrent différents modes : calcul sur un tenseur entier (via surcharge d'opérateurs, ex:
dst = src1 + src2), calcul sur les n premiers éléments (ex:Add(dst, src1, src2, n)), et calcul de découpage haute dimension (plus flexible, avec contrôle des répétitions, des pas d'itération, etc.). - API de Transfert de Données : Comme
DataCopy, pour déplacer les données entre la mémoire globale et la mémoire locale. Un transfert efficace est crucial pour la performance globale. - API de Gestion de Mémoire : Comme
AllocTensoretFreeTensor, pour gérer dynamiquement la mémoire locale limitée de l'AI Core, permettent sa réutilisation. - API de Synchronisation de Tâches : Comme
EnQueetDeQue, pour gérer les dépendances et la communication entre instructions asynchrones exécutées sur différents composants.
API de Haut Niveau : Avantages et Cas d'Usage
Les API de haut niveau encapsulent des algorithmes courants en appelant plusieurs API de base. Elles accélèrent le développement de logique complexe. Par exemple, pour une multiplication matricielle (Matmul), on crée une classe spécifique décrivant les matrices (A, B), le résultat (C), un éventuel biais, ainsi que leurs formats et types de données. L'API de haut niveau gère alors le flux de données et l'opération. Ces API sont particulièrement utiles pour implémenter des opérateurs standards de deep learning (convolution, pooling, normalisation). Pour des algorithmes novateurs ou des schémas de calcul particuliers, les API de base restent nécessaires.
Stockage des Données et Objets Tensor
Ascend C abstrait la mémoire de l'AI Core. La mémoire interne est la Local Memory, la mémoire externe est la Global Memory. Les données y sont manipulées via des objets :
GlobalTensor: Représente des données dans la Global Memory. On initialise son buffer avecSetGlobalBuffer.LocalTensor: Représente des données dans la Local Memory, utilisé pour les calculs.
Cette abstraction unifiée simplifie la conception des API. Le modèle de programmation typique consiste à charger des blocs de données de la Global Memory vers la Local Memory pour les traiter, puis à écrire les résultats en retour.
Développement Détaillé des Fonctions Noyau
Une fonction noyau est l'entrée côté périphérique d'un opérateur Ascend C. Lors de son appel, plusieurs noyaux exécutent le même code de manière parallèle sur des données différentes (paradigme SPMD).
Définition et Règles d'Appel
La définition doit utiliser les qualificateurs __global__ et __aicore__. Les paramètres de type pointeur doivent être qualifiés avec __gm__ (ou utiliser GM_ADDR). La fonction doit retourner void et n'accepter que des types de données primitifs ou des pointeurs.
Exemple simplifié d'une fonction noyau pour une addition :
extern "C" __global__ __aicore__ void operateur_addition(GM_ADDR entree_a, GM_ADDR entree_b, GM_ADDR sortie_c)
{
OperateurAdditionnel operateur;
operateur.Initialiser(entree_a, entree_b, sortie_c);
operateur.Executer();
}
L'appel depuis l'hôte :
void lancer_addition(uint32_t nb_noeuds, void* ctrl_l2, void* flux, uint8_t* a, uint8_t* b, uint8_t* c)
{
operateur_addition<<<nb_noeuds, ctrl_l2, flux>>>(a, b, c);
}
Paradigme de Programmation et Pipeline à Trois Étages
Ascend C promeut un paradigme structuré, notamment le modèle "Entrée, Calcul, Sortie" (CopyIn, Compute, CopyOut). Ce modèle clarifie le flux de données :
- CopyIn (Entrée) : Transfert des données de la Global Memory vers la Local Memory via des API de transfert.
- Compute (Calcul) : Exécution des opérations (vectorielles, matricielles, etc.) sur les données en Local Memory.
- CopyOut (Sortie) : Transfert des résultats de la Local Memory vers la Global Memory.
Ce pipeline permet une paralélisation efficace : pendant qu'un lot de données est calculé, un autre peut être chargé et les résultats d'un lot précédent peuvent être déchargés, masquant ainsi les latences de transfert.
Parallélisme Multi-Cœur et Découpage des Données
L'exploitation des multiples AI Cores repose sur le découpage des données. Si 8 noyaux sont utilisés pour traiter un vecteur de 16 384 éléments, chaque noyau en traite 2 048. L'adresse de départ pour chaque noyau est décalée de GetBlockIdx() * 2048. Ce décalage est appliqué lors de l'initialisation des GlobalTensor.
À l'intérieur d'un noyau, les données peuvent être à nouveau divisées en blocs plus petits (Tiling) pour activer le double buffering, exécutant ainsi les étapes de chargement et de calcul en parallèle sur des blocs de données successifs.
Implémentation Pratique d'un Opérateur d'Addition
Analyse et Spécifications
Considérons un opérateur d'addition élémentaire : z = x + y. L'analyse définit :
- Entrées :
x,yde typehalf, de forme(8, 2048), formatND. - Sortie :
zde même type, forme et format. - Fonction Noyau :
add_customavec trois paramètres GM_ADDR. - API Utilisées :
DataCopypour les transferts,Addpour le calcul,AllocTensor/FreeTensorpour la mémoire locale, et des files d'attente (Queues) pour la synchronisation.
Implémentation de la Fonction Noyau et de la Classe
La fonction noyau crée et utilise une classe d'opérateur :
extern "C" __global__ __aicore__ void add_custom(GM_ADDR src_a, GM_ADDR src_b, GM_ADDR dst_c)
{
CalculAddition operateur;
operateur.PrendreEnCharge(src_a, src_b, dst_c);
operateur.Traiter();
}
La classe d'opérateur encapsule la logique et la mémoire :
class CalculAddition {
public:
__aicore__ inline CalculAddition() {}
__aicore__ inline void PrendreEnCharge(GM_ADDR a, GM_ADDR b, GM_ADDR c)
{
// Découpage des données pour le noyau courant
uint32_t decalage = TAILLE_BLOC * GetBlockIdx();
tamponA.SetGlobalBuffer((__gm__ half*)a + decalage, TAILLE_BLOC);
tamponB.SetGlobalBuffer((__gm__ half*)b + decalage, TAILLE_BLOC);
tamponC.SetGlobalBuffer((__gm__ half*)c + decalage, TAILLE_BLOC);
// Initialisation des files d'attente pour le double buffering
tuyau.InitialiserFile(fileEntreeA, 2, TAILLE_TUILE * sizeof(half));
tuyau.InitialiserFile(fileEntreeB, 2, TAILLE_TUILE * sizeof(half));
tuyau.InitialiserFile(fileSortieC, 2, TAILLE_TUILE * sizeof(half));
}
__aicore__ inline void Traiter()
{
for (int32_t i = 0; i < NOMBRE_TUILES * 2; i++) { // 2 pour le double buffering
ChargerEntree(i);
EffectuerCalcul(i);
DechargerSortie(i);
}
}
private:
__aicore__ inline void ChargerEntree(int32_t etape)
{
LocalTensor<half> tensA = fileEntreeA.AllouerTensor();
DataCopy(tensA, tamponA[etape * TAILLE_TUILE], TAILLE_TUILE);
fileEntreeA.MettreEnFile(tensA);
LocalTensor<half> tensB = fileEntreeB.AllouerTensor();
DataCopy(tensB, tamponB[etape * TAILLE_TUILE], TAILLE_TUILE);
fileEntreeB.MettreEnFile(tensB);
}
__aicore__ inline void EffectuerCalcul(int32_t etape)
{
LocalTensor<half> entreeA = fileEntreeA.RetirerDeFile<half>();
LocalTensor<half> entreeB = fileEntreeB.RetirerDeFile<half>();
LocalTensor<half> resultat = fileSortieC.AllouerTensor<half>();
Add(resultat, entreeA, entreeB, TAILLE_TUILE); // Opération vectorielle
fileEntreeA.LibererTensor(entreeA);
fileEntreeB.LibererTensor(entreeB);
fileSortieC.MettreEnFile<half>(resultat);
}
__aicore__ inline void DechargerSortie(int32_t etape)
{
LocalTensor<half> resultat = fileSortieC.RetirerDeFile<half>();
DataCopy(tamponC[etape * TAILLE_TUILE], resultat, TAILLE_TUILE);
fileSortieC.LibererTensor(resultat);
}
TPipe tuyau;
TQue<:vecin> fileEntreeA, fileEntreeB;
TQue<:vecout> fileSortieC;
GlobalTensor<half> tamponA, tamponB, tamponC;
};
</half></:vecout></:vecin></half></half></half></half></half></half></half></half></half></half></half>
Compilation et Validation
La validation peut se faire en mode CPU (pour le débogage logique) ou en mode NPU (pour une exécution réelle et une évaluation des performances). Le flux de développement typique comprend : la préparation de l'environnement, l'analyse de l'opérateur, l'implémentation du noyau, et les étapes de validation sur CPU puis NPU. Ce processus systématique permet de développer et déployer des opérateurs Ascend C efficacement.