Matrix Multiplikation auf der Grafikkarte mit .NET und Cudafy
Die Parallelisierung von Algorithmen und Programmteilen hat mir schon immer Spaß gemacht und ich habe mir sogar zu Hause einen Rechencluster gebaut, mit dem ich experimentieren konnte. In der aktuellen dotnetpro Ausgabe (07/2013) beschreibt Bernd Marquardt, bei dem ich auf der Parallel 2012 Konferenz einen .NET TPL Workshop mitmachen durfte, in einem Artikel die Parallelisierung von Algorithmen mit AMP unter C++, die dann auf der Grafikkarte ausgeführt werden. Leider muss man hier immer noch den Umweg über C++ gehen, aber glücklicherweise gibt es für .NET mit Cudafy ein Framework, mit dem man diesen Umweg nicht gehen muss. Cudafy unterstützt, neben dem namensgebenden CUDA von Nvidia, auch OpenCL, sodass man damit plattformübergreifende Parallelisierungen vornehmen kann. Ich habe in den letzten Tage ein wenig mit Cudafy herumgespielt und möchte hier einmal ein recht simples Beispiel der Matrixmultiplikation beschreiben.
Um Cudafy zu verwenden, benötigen wir zuerst ein GPGPU Objekt, das unsere zu verwendende Hardware repräsentiert. Cudafy unterstützt dabei sowohl CUDA von Nvidia als auch OpenCL der Khronos Group. Ich habe mich in meinem Beispiel für OpenCL entschieden, denn so konnte ich sowohl die CPU als auch die beiden Grafikkarten in meinem System (Intel und Nvidia) zur Berechnung auswählen. Folgender Codeschnippsel zeigt die Instanziierung des GPGPU Objektes, ein ganz rudimentäres Exception Handling (jaja, ich weiss ;-)) und die Instanziierung meiner Matrixmultiplikationsklasse.
CudafyModes.Target = eGPUType.OpenCL;
CudafyModes.DeviceId = 2;
CudafyTranslator.Language = eLanguage.OpenCL;
try
{
var gpu = CudafyHost.GetDevice(eGPUType.OpenCL, CudafyModes.DeviceId);
Console.WriteLine("Running examples using {0}", gpu.GetDeviceProperties().Name);
Console.WriteLine("Available devices {0}", CudafyHost.GetDeviceCount(eGPUType.OpenCL));
var matrixMult = new MatrixMultiplication(gpu, 512);
matrixMult.Execute();
matrixMult.CpuCalculation();
}
catch (Exception ex)
{
Console.WriteLine(ex);
}
Console.ReadKey();
Wichtigster Grundstein meiner Entwicklungsumgebung ist eine abstrakte Basisklasse, von denen meine implementierten Algorithmenklassen erben. Dieser Schritt ist nicht unbedingt notwendig, macht aber in meinem Fall Sinn, da ich beim “Herumspielen” nicht nur die Matrixmultiplikation implementiert habe, sondern auch andere Algorithmen.
using System;
using System.Collections.Generic;
using System.Linq;
using System.Text;
using System.Threading.Tasks;
using Cudafy;
using Cudafy.Host;
using Cudafy.Translator;
using System.Diagnostics;
namespace CudafyTest.Common
{
internal abstract class GpuCalculation
{
protected GPGPU Gpu { get; set; }
public GpuCalculation(GPGPU gpu)
{
Gpu = gpu;
}
public void Execute()
{
CudafyModule km = CudafyTranslator.Cudafy(this.GetType());
Gpu.LoadModule(km);
var stopWatch = new Stopwatch();
stopWatch.Start();
OnExecute();
stopWatch.Stop();
Console.WriteLine("Execution time = {0} ms", stopWatch.ElapsedMilliseconds);
}
protected abstract void OnExecute();
}
}
Im Konstruktor muss ein GPGPU Objekt übergeben werden, das die ausgewählte Client Hardware repräsentiert (die Instanziierung wurde bereits weiter oben beschrieben). Die Execute Methode enthält nun den Code, der benötigt wird, m einen Kernel mit Cudafy für die zu verwendende Client Hardware vorzubereiten (Aufruf von Cudafy(…) und LoadModule(..)). Außerdem wird noch eine Zeitmessung mit der allseits bekannten Stopwatch() durchgeführt. In diesem Block wird die innerhalb der Algorithmenklasse zu implementierenden Methode OnExecute() aufgerufen, sodass nach Ausführung die Dauer der Berechnung (inklusive der Übertragungszeit der Daten zur berechnenden Hardware) in der Konsole ausgegeben wird.
Die Klasse zur Matrixmultiplikation sieht dann so aus.
using Cudafy;
using Cudafy.Host;
using CudafyTest.Common;
using System;
using System.Collections.Generic;
using System.Diagnostics;
using System.Linq;
using System.Text;
using System.Threading.Tasks;
namespace CudafyTest.Calculations
{
class MatrixMultiplication : GpuCalculation
{
private float\[,\] m\_A;
private float\[,\] m\_B;
private float\[,\] m\_Result;
private float\[,\] m\_ResultCpu;
private int m\_Dimension;
public MatrixMultiplication(GPGPU gpgpu, int number) : base(gpgpu)
{
m\_Dimension = number;
InitExampleArrays(number);
}
private void InitExampleArrays(int number)
{
m\_A = new float\[number, number\];
m\_B = new float\[number, number\];
m\_Result = new float\[number, number\];
//fill the array
for (int i = 0; i < number; i++)
{
for (int j = 0; j < number; j++)
{
m\_A\[i, j\] = 1.0f + (float)i \* (float)(number - j) / (float)number;
m\_B\[i, j\] = 1.0f + (float)(number - i) \* (float)j / (float)number;
}
}
//PrintResult(m\_A);
//PrintResult(m\_B);
}
protected override void OnExecute()
{
float\[,\] devA = Gpu.CopyToDevice(m\_A);
float\[,\] devB = Gpu.CopyToDevice(m\_B);
float\[,\] devResult = Gpu.Allocate(m\_Result);
Gpu.Launch(new dim3(m\_Dimension, m\_Dimension), 1).Multiply(m\_Dimension, devA, devB, devResult);
Gpu.CopyFromDevice(devResult, m\_Result);
Gpu.FreeAll();
//PrintResult(m\_Result);
}
\[Cudafy\]
public static void Multiply(GThread gthread, int dimension, float\[,\] a, float\[,\] b, float\[,\] result)
{
int x = gthread.blockIdx.x \* gthread.blockDim.x + gthread.threadIdx.x;
int y = gthread.blockIdx.y \* gthread.blockDim.y + gthread.threadIdx.y;
if (x >= dimension || y >= dimension)
{
return;
}
float sum = 0.0f;
for (int k = 0; k < dimension; k++)
{
sum += a\[x, k\] \* b\[k, y\];
}
result\[x, y\] = sum;
}
public void CpuCalculation()
{
Stopwatch watch = new Stopwatch();
watch.Start();
m\_ResultCpu = new float\[m\_Dimension, m\_Dimension\];
for (int x = 0; x < m\_Dimension; x++)
{
for (int y = 0; y < m\_Dimension; y++)
{
float sum = 0.0f;
for (int k = 0; k < m\_Dimension; k++)
{
sum += m\_A\[x, k\] \* m\_B\[k, y\];
}
m\_ResultCpu\[x, y\] = sum;
}
}
Console.WriteLine("CPU Calculation:");
//PrintResult(m\_Result);
watch.Stop();
Console.WriteLine("Time elapsed {0} ms", watch.ElapsedMilliseconds);
}
private void PrintResult(float\[,\] m\_Result)
{
for (int y = 0; y < m\_Dimension; y++)
{
for (int x = 0; x < m\_Dimension; x++)
{
Console.Write("{0:F} ", m\_Result\[x, y\]);
}
Console.WriteLine();
}
Console.WriteLine();
}
}
}
Wie bereits erwähnt, ist diese Klasse während meiner Tests entstanden, sodass ich in diesem Beispiel Testdaten verwenden, die ich bei Aufruf des Konstruktors generiere. Dem Konstruktor wird in diesem Beispiel neben dem GPGPU Objekt auch noch die Dimension der Matrix mit übergeben. Die wichtigsten Teile des Codes stecken aber in der Methode OnExecute() und Multiply(…). Die Methode OnExecute() ist dafür verantwortlich, den Speicher auf der Client Hardware anzulegen und die Arrays vom Host auf den Client zu übertragen. Dies wird hier durch die Methode CopyToDevice(…) vorgenommen. Das Ergebnisarray wird nicht auf die Client Hardware kopiert, sondern nur der Speicher reserviert, da es zu Beginn sowieso leer ist und deshalb keine Daten benötigt werden. Die Launch(…) Methode startet dann die Berechnung auf der Client Hardware. Den Aufruf werde ich hier nicht weiter erläutern, mehr zur Launch(…) Methode findet man aber in der Cudafy Dokumentation. Nach dem Aufruf wird das Ergebnis von der Hardware wieder auf den Host übertragen und der belegte Speicher auf dem Client wieder freigegeben. Das war’s auch schon.
Die Client Implementierung des Algorithmus steckt in der mit dem “Cudafy” Attribut markierten Methode Multiply(…). Die Methodenparameter sind GThread (wird standardmäßig von Cudafy hinzugefügt und enthält Statusinformationen), die Dimension der Matrix und die Matrizen für die Matrixmultiplikation. In dieser Methode wird zuerst die x und y Position innerhalb der Matrix errechnet. Die Berechnung orientiert sich hier an den Blöcken, die über die Launch Methode mit übergeben werden. Genauere Informationen dazu findet ihr in den Grundlagen zum Thema OpenCL und CUDA. Im nächsten Schritt wird vorsichtshalber noch überprüft, dass x und y auch wirklich innerhalb der Matrix liegen und anschließend die Matrixmultiplikation durchgeführt, die dann im Ergebnisarray gespeichert wird.
Die Methode CpuCalculation() habe ich implementiert, um die Geschwindigkeit einer einzelnen CPU mit der Ausführung auf einem Client zu vergleichen. Auf meinem Rechner war eine Nvidia GT650M bei einer Dimension von 256 bereits 3x so schnell wie ein Kern meiner i7 CPU. Bei einer Dimension von 512 war sie schon um den Faktor 7,5 schneller.
Abschließend kann man sagen, dass sich die Nutzung einer Grafikkarte als Client bei einem hinreichend großen Problem lohnt und man damit eine große Beschleunigung erreichen kann. Dank Cudafy funktioniert das alles auch in .NET ohne den Umweg über C++. In meinem Beispiel bleibt außerdem noch viel Platz für Optimierungen, da ich das Beispiel möglichst simpel halten wollte - und das gilt sowohl für die Berechnung der Matrixmultiplikation auf der CPU (hier wäre bspw. eine Implementierung mit der TPL sinnvoll, die dann alle Kerne nutzt) als auch auf der GPU.
