De bijdrage van Codeplay aan DPC++ brengt SYCL-ondersteuning voor NVIDIA GPU's

Codeplay maakt vanaf het begin deel uit van de SYCL™-gemeenschap en ons team heeft de afgelopen 5 jaar samengewerkt met collega's van enkele van de grootste halfgeleiderleveranciers, waaronder Intel en Xilinx, om de SYCL-standaard te definiëren. In die tijd is de interface die SYCL biedt aan ontwikkelaars geëvolueerd en we verwachten dit jaar te kunnen praten over de nieuwste versie van de standaard. ComputeCpp, onze implementatie van SYCL, was de eerste implementatie die SYCL v1.2.1-conform was, iets waar we erg trots op zijn en we blijven dit product ontwikkelen om de prestaties te verbeteren en het aantal ondersteunde apparaten te vergroten.

Vorig jaar waren we erg enthousiast toen we hoorden dat Intel enorm veel moeite doet voor SYCL en een API heeft ontwikkeld met DPC++ (een implementatie van SYCL met extensies) voor zijn CPU's, GPU's en FPGA's. Met de extra ondersteuning voor SYCL van Xilinx, Renesas en Imagination Technologies, kunnen softwareontwikkelaars zich nu richten op een breed scala aan apparaten met SYCL.

Terwijl ComputeCpp experimentele ondersteuning biedt voor NVIDIA® GPU's met behulp van OpenCL™ en NVIDIA's PTX, bood DPC++ (Intel®'s SYCL-implementatie) de mogelijkheid om volledige ondersteuning toe te voegen voor NVIDIA GPU's geïntegreerd in de LLVM-compiler zonder door te gaan OpenCL, en vandaag zijn we verheugd om de eerste, experimentele fase van onze implementatie open source te maken waarmee SYCL-ontwikkelaars zich kunnen richten op NVIDIA GPU's.

De codebase voor deze implementatie bevindt zich in een aparte vork van zowel het hoofd-LLVM-compilerproject als de DPC++-tak, wat betekent dat je dit project voorlopig moet gebruiken om de ondersteuning voor NVIDIA GPU's met SYCL uit te proberen. We streven ernaar om met Intel samen te werken om de NVIDIA GPU-ondersteuning toe te voegen aan de upstream Intel/llvm-compiler.

Dit is een vroege, onvolledige release en er wordt verder gewerkt aan de integratie van meer functies en verbeterde prestaties.

Hoe is de NVIDIA GPU-ondersteuning geïmplementeerd?

De oorspronkelijke benadering die we in ComputeCpp gebruikten om NVIDIA-platforms te ondersteunen, was om te vertrouwen op de NVIDIA OpenCL 1.2-implementatie en onze apparaatcompiler te gebruiken om PTX (de NVIDIA Intermediate Representation) uit te zenden in plaats van SPIR-V. Deze aanpak was voldoende om te demonstreren dat SYCL op verschillende platforms draait met behulp van open standaarden, maar we realiseerden ons al snel dat de mogelijkheden van de OpenCL 1.2-ondersteuning van NVIDIA beperkt waren.

De ondersteuning voor NVIDIA-platforms die we toevoegen aan de DPC++-compiler is rechtstreeks gebaseerd op NVIDIA's CUDA™, in plaats van OpenCL. DPC++ gebruikt een Plugin Interface (PI) om verschillende backends te targeten. Intel biedt een OpenCL 2.2-plug-in waarmee DPC++ kan worden uitgevoerd op OpenCL-platforms, en we hebben een plug-in geïmplementeerd die tijdens runtime kan worden geselecteerd (door de omgevingsvariabele SYCL_BE=PI_CUDA in te stellen).
Door native CUDA te gebruiken, doet DPC++ dat niet moeten vertrouwen op de OpenCL-ondersteuning van NVIDIA, waardoor meer functies mogelijk zijn en mogelijk betere algehele prestaties.

Verder hebben we verschillende wijzigingen doorgevoerd in de libclc-bibliotheek om ondersteuning toe te voegen voor PTX-ingebouwde onderdelen die zijn gemaakt met SPIR-V en dit maakt het gebruik van SYCL-ingebouwde onderdelen op NVIDIA GPU's mogelijk.

Aangezien een SYCL voor CUDA-toepassing een native CUDA-toepassing is , zullen alle bestaande tooling en bibliotheken van het CUDA-ecosysteem werken met een SYCL-applicatie die is gebouwd met behulp van deze back-end.

Waar kan ik dit project voor gebruiken?

Met dit project kunt u NVIDIA GPU's targeten met SYCL-code, zonder dat u door de OpenCL-laag in het systeem hoeft te gaan. Als u een NVIDIA GPU hebt, kunt u nu DPC++ op uw systeem uitvoeren om SYCL-toepassingen te compileren. Als u een bestaande CUDA-toepassing hebt, kunt u deze stapsgewijs overzetten naar SYCL met behulp van de CUDA-ondersteuning en deze vervolgens uitvoeren op een platform dat geen CUDA heeft. Hierdoor kunnen ontwikkelaars hun applicatie snel overzetten naar andere platforms.

Hoe gebruik ik het project?

Er zijn instructies in het README-bestand van het project waarin wordt uitgelegd hoe de NVIDIA-back-end voor DPC++ moet worden gebruikt. Deze zijn gemakkelijk te volgen en vereisen het gebruik van enkele vlaggen bij het compileren, en wat code om uw apparaatkiezer in te stellen om ervoor te zorgen dat de runtime weet op welk apparaat het moet worden gericht. In het bijzonder de sectie “Build SYCL toolchain met ondersteuning voor NVIDIA CUDA” en de Clang-compileropties hebben specifieke instructies.

In SYCL wordt een apparaatkiezer gebruikt om je doel te vertellen op welk apparaat je je code wilt uitvoeren. Voor NVIDIA-apparaten moet uw SYCL-apparaatselector er ongeveer zo uitzien:

class CUDASelector: public cl::sycl::device_selector {
public:
int operator()(const cl::sycl::device &Device) const overschrijven {
met gebruik van naamruimte cl::sycl::info;
const std::string DeviceName = Device.get_info() ;
const std::string DeviceVendor = Device.get_info();
if (Device.is_gpu() && (DeviceName.find(“NVIDIA”) != std::string:: npos)) {
return 1;
};
return -1;
}
};

Wat werkt en wat niet ?

Dit project is getest met Ubuntu 18.04 met behulp van CUDA 10.1 op een Titan RTX GPU (rekenmogelijkheden 7.5). Het zou echter ook moeten werken op andere Linux-versies met elke NVIDIA GPU die compatibel is met SM 5.0 of hoger.

Op dit moment kan de gecompileerde SYCL-toepassing alleen CUDA of OpenCL targeten, niet beide op dezelfde tijd. Om een ​​SYCL-toepassing voor de CUDA-back-end te bouwen, is de vlag nvptx64-nvidia-cuda-sycldevice vereist, zoals hieronder weergegeven:

$ clang++ -fsycl -fsycl-targets=nvptx64-nvidia- cuda-sycldevice sycl-app.cpp -o sycl-app.exe

Vervolgens moet bij het uitvoeren van de toepassing de CUDA SYCL-back-end tijdens runtime worden geselecteerd:

$ SYCL_BE=PI_CUDA sycl-app.exe

Dit wordt later vervangen zodra de De PI-interface kan meerdere back-ends tegelijkertijd ondersteunen.

We zijn ons ook bewust van enkele situaties waarin de NVIDIA OpenCL-headers een conflict veroorzaken met de OpenCL-headers die nodig zijn voor het gebruik van DPC++. De NVIDIA OpenCL-headers zijn voor een oudere versie van OpenCL die niet wordt ondersteund door DPC++. Hoewel de CMake-configuratie voor DPC++ zijn eigen headers installeert, hebben we in sommige configuraties geconstateerd dat de NVIDIA OpenCL-headers nog steeds zijn geselecteerd.

Over het algemeen komt de geboden functionaliteit overeen met wat u vandaag kunt doen met de NVIDIA OpenCL-implementatie , maar we werken aan uitbreidingen van SYCL die meer mogelijkheden van de onderliggende hardware beschikbaar stellen.

Welke prestaties kan ik verwachten?

Deze initiële release van het project is niet geoptimaliseerd, dus de prestaties zijn mogelijk niet onder alle omstandigheden optimaal.

Momenteel past de codegeneratie geen specifieke optimalisatiepassen toe, en de CUDA-implementatie van de SYCL specificatie in de plug-ininterface gebruikt geen geavanceerde CUDA-functies.

Hoe kan ik ondersteuning krijgen voor dit project?

Als u een probleem tegenkomt dat specifiek verband houdt met NVIDIA GPU-ondersteuning , geef dat dan aan in uw uitgave in de Intel/llvm-repository.

Wat zijn de toekomstplannen voor dit project?

De huidige focus is om zoveel mogelijk functionaliteit te implementeren, zodat ontwikkelaars een reeks SYCL-applicaties op NVIDIA GPU-platforms kunnen uitvoeren. Zodra we er zeker van zijn dat we een goede dekking hebben, zullen we samenwerken met de gemeenschap om prestatieknelpunten te identificeren die moeten worden aangepakt.

Ons doel is om de ervaring voor SYCL-ontwikkelaars te verbeteren en de ondersteuning van SYCL voor apparaten die op het CUDA-platform zijn gebouwd, te vergroten. De componenten die open source zijn door Codeplay zijn herbruikbaar door andere implementaties, waaronder ComputeCpp, dus verwacht dat de ecosysteemondersteuning voor NVIDIA-hardware in 2020 zal toenemen.

Hoe kan ik bijdragen?

We zijn geïnteresseerd om te zien hoe uw code presteert op verschillende NVIDIA GPU's om te identificeren waar prestatieknelpunten zich bevinden en welke variaties optreden bij verschillende hardwaremodellen. Als je je eigen resultaten wilt geven, dien deze dan in via de Issue tracker in de Intel/llvm-repository.

Ruyman Reyes
Principal Software Engineer, Programming Models
Codeplay Software

Ruyman Reyes
Principal Software Engineer, Programming Models
Codeplay Software

p>

SYCL en SPIR zijn handelsmerken van de Khronos Group Inc. Nvidia en CUDA zijn geregistreerde handelsmerken van NVIDIA Corporation. OpenCL en het OpenCL-logo zijn handelsmerken van Apple Inc. die met toestemming van Khronos worden gebruikt. Imagination Imagination Technologies is een geregistreerd handelsmerk van Imagination Technologies Limited. Intel is een handelsmerk van Intel Corporation in de VS en/of andere landen. Linux is het gedeponeerde handelsmerk van Linus Torvalds in de VS en andere landen.