Pourquoi connaître le modèle de programmation GPU de plusieurs fournisseurs est un mal nécessaire… ou l'est-il ?

Introduction

Dans mes deux derniers articles sur la programmation parallèle et par accélérateur, j'ai parlé des bases de la programmation par accélérateur et parallèle et de certains des concepts de programmation nécessaires pour garantir l'exactitude de votre code. La question intéressante pour le programmeur d'apprentissage de la programmation des accélérateurs est "Ai-je vraiment besoin d'apprendre CUDA, ROCm et SYCL pour programmer pour les accélérateurs NVIDIA, AMD et Intel?"

Lors de la planification de cette série, je m'attendais à plonger dans les similitudes et les différences entre l'exécution matérielle et les modèles de mémoire de chaque architecture. J'allais parler des distorsions, des fronts d'onde et des groupes de travail, les constructions respectives que CUDA, ROCm et SYCL utilisent pour regrouper l'exécution sur les threads matériels. Cependant, plus j'y pense, plus cela semble un peu trop pour un article de blog. Si ces sujets vous intéressent, consultez le joli SYCL pour les développeurs CUDA de Codeplay.

Au lieu de cela, aujourd'hui, je vais me concentrer sur la façon de vous faire passer de l'absence de code à une exécution sur un GPU, n'importe quel GPU, le plus rapidement possible. Rien n'est aussi amusant que de voir les avantages de performance de votre accélérateur le plus rapidement possible. Pour moi, c'est amusant car cela coïncide avec la sortie de Codeplay oneAPI pour NVIDIA et oneAPI pour AMD (bêta) .

SYCL et plusieurs GPU de fournisseurs

SYCL est un modèle de programmation qui prend en charge le compilateur pour les GPU NVIDIA, AMD et Intel. Vous pouvez écrire votre code en SYCL, puis le créer et l'exécuter sur ces GPU de fournisseurs. Auparavant, cela nécessitait l'utilisation de divers compilateurs communautaires ou la création du compilateur Intel Open Source, en fonction de votre cible GPU. Avec la sortie des dernières chaînes d'outils Codeplay, vous pouvez désormais exécuter rapidement et facilement votre code via une seule chaîne d'outils prédéfinie. Pour plus de détails à ce sujet, vous pouvez lire le blog de Ruyman Reyes, CTO de Codeplay, sur leur nouvelle version.

Une question que vous vous posez peut-être est la suivante : puis-je réellement créer mon code une fois, puis décider au moment de l'exécution sur quel GPU exécuter ? La réponse est OUI ! C'est la bonne chose à propos de ce support, vous construisez votre code une fois et au moment de l'exécution, vous pouvez choisir votre cible manuellement ou laisser la bibliothèque d'exécution choisir automatiquement pour vous.

Puisque vous êtes arrivé jusqu'ici, je suppose que cela vous intéresse et mon objectif est de vous permettre d'être opérationnel sur le GPU de votre choix. Pour mes besoins, je teste les 3 GPU du fournisseur sur Ubuntu 22.04 et je procède à l'installation via APT.

Installation de la chaîne d'outils Codeplay

Tout d'abord, nous avons configuré quelques packages système de base :

sudo apt update
sudo apt -y install cmake pkg-config build-essential

# download the key to system keyring
wget -O- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB \
| gpg --dearmor | sudo tee /usr/share/keyrings/oneapi-archive-keyring.gpg > /dev/null

# add signed entry to apt sources and configure the APT client to use Intel repository:
echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main" | sudo tee /etc/apt/sources.list.d/oneAPI.list

sudo apt-update
sudo apt install intel-basekit

Assurez-vous d'avoir les pilotes/frameworks de base dont vous avez besoin pour chaque GPU que vous avez/voulez cibler, encore une fois, je relie les instructions Ubuntu :

Installation du pilote Intel GPU

Les instructions pour les GPU Intel sur différents systèmes d'exploitation sont disponibles ici . Parce que j'utilise un Intel Arc sur Ubuntu 22.04, j'ai suivi ces instructions .

oneAPI pour l'installation de NVIDIA

Vous pouvez aller ici pour obtenir les dernières instructions officielles sur la façon d'installer CUDA, j'intègre ce que j'utilise pour le faire fonctionner:

1. Installez le trousseau de clés CUDA, CUDA et redémarrez

wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.0-1_all.deb
sudo dpkg -i cuda-keyring_1.0-1_all.deb

sudo apt-get update
sudo apt-get install cuda-11-7
sudo apt-get install nvidia-gds
sudo reboot

export PATH=/usr/local/cuda-11.7/bin${PATH:+:${PATH}}
export LD_LIBRARY_PATH=/usr/local/cuda-11.7/lib64\
                         ${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}

sh oneapi-for-nvidia-gpus-2023.0.0-linux.sh

. /opt/intel/oneapi/setvars.sh --include-intel-llvm
export PATH=/usr/local/cuda-11.7/bin${PATH:+:${PATH}}
export LD_LIBRARY_PATH=/usr/local/cuda-11.7/lib64\
                         ${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}

Vous pouvez aller ici pour obtenir les dernières instructions officielles sur la façon d'installer ROCm, encore une fois en ligne ici pour vous faciliter un peu les choses :

1. Installez le programme AMD amdgpu-install via APT

sudo apt-get update
wget https://repo.radeon.com/amdgpu-install/5.4.1/ubuntu/jammy/amdgpu-install_5.4.50401-1_all.deb
sudo apt-get install ./amdgpu-install_5.4.50401-1_all.deb

amdgpu-install --usecase="dkms,graphics,opencl,hip,hiplibsdk"

sh oneapi-for-amd-gpus-2023.0.0-linux.sh

. /opt/intel/oneapi/setvars.sh --include-intel-llvm

export PATH=/PATH_TO_ROCM_ROOT/bin:$PATH
export LD_LIBRARY_PATH=/PATH_TO_ROCM_ROOT/lib:$LD_LIBRARY_PATH

Il existe de nombreux exemples de codes différents que vous pouvez essayer dans le référentiel GitHub des exemples open source oneAPI :

À partir de l'exemple Codeplay, vous pouvez voir qu'ils ont créé ce simple-sycl-app.cpp :

#include <sycl/sycl.hpp>

int main() {
  // Creating buffer of 4 ints to be used inside the kernel code
  sycl::buffer<sycl::cl_int, 1> Buffer(4);

  // Creating SYCL queue
  sycl::queue Queue;

  // Size of index space for kernel
  sycl::range<1> NumOfWorkItems{Buffer.size()};

  // Submitting command group(work) to queue
  Queue.submit([&](sycl::handler &cgh) {
    // Getting write only access to the buffer on a device
    auto Accessor = Buffer.get_access<sycl::access::mode::write>(cgh);
    // Executing kernel
    cgh.parallel_for<class FillBuffer>(
        NumOfWorkItems, [=](sycl::id<1> WIid) {
          // Fill buffer with indexes
          Accessor[WIid] = (sycl::cl_int)WIid.get(0);
        });
  });

  // Getting read only access to the buffer on the host.
  // Implicit barrier waiting for queue to complete the work.
  const auto HostAccessor = Buffer.get_access<sycl::access::mode::read>();

  // Check the results
  bool MismatchFound = false;
  for (size_t I = 0; I < Buffer.size(); ++I) {
    if (HostAccessor[I] != I) {
      std::cout << "The result is incorrect for element: " << I
                << " , expected: " << I << " , got: " << HostAccessor[I]
                << std::endl;
      MismatchFound = true;
    }
  }

  if (!MismatchFound) {
    std::cout << "The results are correct!" << std::endl;
  }

  return MismatchFound;
}

Généralement, pour construire pour nos différentes cibles, nous pourrions utiliser des lignes de commande qui ressemblent à ceci :

Nvidia :

clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda simple-sycl-app.cpp -o simple-sycl-app

clang++ -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=<ARCH> simple-sycl-app.cpp -o simple-sycl-app

clang++ -fsycl -fsycl-targets=spir64 simple-sycl-app.cpp -o simple-sycl-app

Création pour plusieurs GPU de fournisseurs

Disons que je voulais courir pour les GPU NVIDIA et Intel en même temps, je peux en fait créer un gros binaire qui fonctionnera sur un GPU NVIDIA ou Intel. L'astuce ici consiste à définir correctement vos paramètres de ligne de commande de construction pour avoir plusieurs cibles SYCL.

clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda,spir64 simple-sycl-app.cpp -o simple-sycl-app

tonym@LianLi-Linux:~/dev/2023.0Test$ clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda,spir64 simple-sycl-app.cpp -o simple-sycl-app
 
tonym@LianLi-Linux:~/dev/2023.0Test$ SYCL_DEVICE_FILTER=cuda SYCL_PI_TRACE=1 ./simple-sycl-app
SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_cuda.so [ PluginVersion: 11.15.1 ]
SYCL_PI_TRACE[all]: Selected device: -> final score = 1500
SYCL_PI_TRACE[all]:   platform: NVIDIA CUDA BACKEND
SYCL_PI_TRACE[all]:   device: NVIDIA GeForce RTX 3080 Ti
The results are correct!
tonym@LianLi-Linux:~/dev/2023.0Test$
 
tonym@LianLi-Linux:~/dev/2023.0Test$ SYCL_DEVICE_FILTER=ext_oneapi_level_zero:gpu:0 SYCL_PI_TRACE=1 ./simple-sycl-app
SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_level_zero.so [ PluginVersion: 11.15.1 ]
SYCL_PI_TRACE[all]: Selected device: -> final score = 1550
SYCL_PI_TRACE[all]:   platform: Intel(R) Level-Zero
SYCL_PI_TRACE[all]:   device: Intel(R) Graphics [0x56a0]
The results are correct!

Travail futur

Certains lecteurs avisés remarqueront peut-être que je n'ai pas inclus d'exemple AMD dans mon exemple de GPU multifournisseur. C'est parce que j'ai récemment remplacé le GPU AMD 6800XT par un tout nouveau GPU AMD RX 7900XT. Malheureusement, ROCm ne s'installe pas correctement sur mon système Linux, quel que soit le noyau que j'utilise avec la carte RDNA3. Une fois que cela sera mis à jour, je reviendrai et fournirai des instructions avec un exemple fonctionnant sur les 3 cartes.

Notez qu'auparavant une version de ce flux fonctionnait pour le 6800XT, mais c'était avant que le logiciel Codeplay ne soit publié à l'aide du flux open source DPC++ documenté ici :

Conclusion

La oneAPI pour les GPU NVIDIA de Codeplay m'a permis de créer facilement des binaires pour les GPU NVIDIA ou Intel. Le temps de configuration de l'oneAPI supplémentaire pour les GPU NVIDIA était d'environ 10 minutes en plus du temps d'installation du kit d'outils Intel oneAPI Base. Cela comprenait l'installation de CUDA et du logiciel Codeplay.

Le nouveau logiciel Codeplay ouvre une toute nouvelle option pour les particuliers et les entreprises qui envisagent des options GPU multifournisseurs. Auparavant, les options étaient des compilateurs open source pris en charge par la communauté. C'est une option fantastique et ce que j'utilise généralement pour le travail quotidien.

Cependant, pour les personnes qui envisagent une assistance à plus long terme, c'est-à-dire les entreprises qui dépendent des chaînes d'outils pour leur subsistance, le fait d'avoir une option officiellement prise en charge par Codeplay et Intel peut aider à assurer la tranquillité d'esprit.

Si vous voulez voir quelles nouvelles technologiques aléatoires je lis, vous pouvez me suivre sur Twitter. Consultez également Code Together , un podcast Intel pour les développeurs que j'héberge et où nous parlons de technologie.

Tony est architecte logiciel et évangéliste technique chez Intel. Il a été l'architecte d'Intel VTune Profiler et d'autres outils de performance et a récemment dirigé l'équipe d'ingénierie logicielle qui a construit la plate-forme de centre de données qui a permis la solution évolutive MLPerf de Habana.