Laget en fizzbuzz-sak med C og NVIDIA CUDA, hvor utregningen skjer på GPU. Fizzbuzz er her parallellisert over 10 blokker og 10 tråder, og med en GPU som har Compute Level på 2.0 eller høyere går det faktisk an å bruke printf direkte (tidligere ville man lagret resultatet i en array på GPU og kopiert den til CPU for utskrift). Begrensningen med å printe direkte er at trådene fullfører i litt vilkårlig rekkefølge, derfor har jeg slengt på tallet foran i output for hver linje. fizz og buzz instansieres først på CPU og kopieres over til GPU. Kunne helt fint bare skrevet "fizz" og "buzz" inn i kernelen, men ville vise hvordan man går frem for å få data over fra CPU til GPU med cudaMemcpy og cudaMalloc (om du selv skal drive med CUDA anbefaler jeg Thrust-biblioteket, som gjør slik minnehåndtering mye simplere).
Kode
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
__global__ void fizzbuzzKernel(char*, char*);
void main()
{
// Allokerer fizz og buzz på CPU
char *fizz = "fizz\0";
char *buzz = "buzz\0";
// Device-pointere for fizz og buzz
char *dev_fizz = 0;
char *dev_buzz = 0;
// Allokerer plass til fizz og buzz på GPU
cudaMalloc((void**)&dev_fizz, 5 * sizeof(char));
cudaMalloc((void**)&dev_buzz, 5 * sizeof(char));
// Kopierer fizz og buzz fra CPU til GPU
cudaMemcpy(dev_fizz, fizz, 5 * sizeof(char), cudaMemcpyHostToDevice);
cudaMemcpy(dev_buzz, buzz, 5 * sizeof(char), cudaMemcpyHostToDevice);
// Kjører kernel med 10 blokker, som hver har 10 threads
fizzbuzzKernel<<<10, 10>>>(dev_fizz, dev_buzz);
// Synkronisering
cudaDeviceSynchronize();
// ???
// Profit!
}
__global__ void fizzbuzzKernel(char *fizz, char *buzz)
{
int number = threadIdx.x * blockDim.x + blockIdx.x + 1;
if(number % 3 == 0 && number % 5 == 0)
{
printf("%i:\t%s%s\n", number, fizz, buzz);
}
else if(number % 3 == 0)
{
printf("%i:\t%s\n", number, fizz);
}
else if(number % 5 == 0)
{
printf("%i:\t%s\n", number, buzz);
}
else
{
printf("%i:\t%i\n", number, number);
}
}
http://pastebin.com/SNADa3mM