Fonction halves2half2 pour le kit de développement asc-devkit de CANN

__halves2half2

Cette fonction assemble deux valeurs de type half en un seul élément half2, qui est un type vectoriel à deux cmoposantes.

Compatibilité matérielle

Plateforme Support
Ascend 950PR / Ascend 950DT Oui
Atlas série A3 (Entraînement/Inférence) Non
Atlas série A2 (Entraînement/Inférence) Non
Atlas 200I / 500 A2 (Inférence) Non
Atlas série Inférence (AI Core) Non
Atlas série Inférence (Vector Core) Non
Atlas série Entraînement Non

Description

La fonction prend deux scalaires half (demi-précision 16 bits) en entrée. Elle les emballe respectivement dans la composante haute et la composante basse d'un vecteur half2, puis retourne ce vecteur.

Prototype

inline half2 __halves2half2(const half srcLo, const half srcHi)

Paramètres

Nom Direction Description
srcLo Entrée Valeur à placer dans la composante basse du vecteur résultant.
srcHi Entrée Valeur à placer dans la composante haute du vecteur résultant.

Valeur de retour

Un vecteur half2 dont la composente basse vaut srcLo et la composante haute vaut srcHi.

Contraintes

Aucune contrainte particulière.

Fichier d'en-tête requis

Pour utiliser cette fonction intrinsèque, il est nécessaire d'inclure le fichier suivant :

#include "simt_api/asc_fp16.h"

Exemples d'utilisation

Scénario de programmation SIMT :

// L'emploi de vecteurs courts optimise l'efficacité du transfert de données
__aicore__ void assembler_vecteurs(half* tableauA, half* tableauB, half2* resultat, uint32_t taille_totale)
{
    uint32_t position = blockIdx.x * blockDim.x + threadIdx.x;
    if (position >= taille_totale) {
        return;
    }
    resultat[position] = __halves2half2(tableauA[position], tableauB[position]);
}

__global__ __launch_bounds__(1024) void noyau_fusion(half* entree1, half* entree2, half* sortie, uint32_t longueur)
{
    assembler_vecteurs(entree1, entree2, (half2*)sortie, longueur);
}

Scénario de programmation hybride SIMD/SIMT :

// L'emploi de vecteurs courts optimise l'efficacité du transfert de données
__simt_vf__ __launch_bounds__(1024) inline void assembler_vecteurs_mixte(__gm__ half* src1, __gm__ half* src2, __gm__ half2* dest, uint32_t len)
{
    uint32_t pos = blockIdx.x * blockDim.x + threadIdx.x;
    if (pos >= len) {
        return;
    }
    dest[pos] = __halves2half2(src1[pos], src2[pos]);
}

__global__ __vector__ void noyau_fusion_mixte(__gm__ half* in1, __gm__ half* in2, __gm__ half* out, uint32_t total)
{
    asc_vf_call<assembler_vecteurs_mixte>(dim3(1024), in1, in2, (__gm__ half2*)out, total);
}
</assembler_vecteurs_mixte>

Étiquettes: cann Ascend half2 SIMT SIMD

Publié le 27 juin à 18h18