cuda (2)

CUDA 1.

W oparciu o środowisko CUDA napisać równoległy program szyfrujący plik tekstowy szyfrem Cezara. Należy zaproponować kod kernela oraz jego wywołanie z kodem odpowiedzialnym za przydział pamięci z podziałem na bloki w dwóch wymiarach. Proszę ominąć obsługę błędów. #include "cuda_runtime.h" #include "device_launch_parameters.h" #include #include #include #define KROK 4 #define SIZE 1000 __global__ void SzyfrCezara(char * tekst, char * szyfrowany) { int i = threadIdx.x; //szyfr cezara to alfabet przesunięty o jakiś krok //jeżeli wybieramy krok = 4 to A będzie E, B będzie F itp. szyfrowany[i] = tekst[i] + KROK; } int main() { //obliczamy wczesniej rozmiar, którego bedziemy wszedzie uzywac size_t size = SIZE * sizeof(char); //tekst oryginalny, alokacja w pamieci procesora char * tekst = (char*) malloc(size);

cudaMalloc(&cTekst, size); cudaMalloc(&cSzyfrowany, size); //kopiowanie tablic wejsciowych na urzadzenie cudaMemcpy(cTekst, tekst, size, cudaMemcpyHostToDevice); //tego nie trzeba kopiowac bo to jest puste, jedynie zaalokowane //i tylko bedziemy do tego pisać //cudaMemcpy(cSzyfrowany, szyfrowany, size, cudaMemcpyHostToDevice); int threadsPerBlock = 256; int blocksPerGrid=(SIZE + threadsPerBlock - 1 ) / threadsPerBlock; SzyfrCezara<<>>(tekst, szyfrowany); //kopiowanie wyników do pamieci procesora cudaMemcpy(szyfrowany, cSzyfrowany, size, cudaMemcpyDeviceToHost); //tutaj tak samo w druga strone oryginalnego tekstu nie trzeba kopiwoac //oryginalnego tekstu z urzadzenia do procesora bo sie nie zmienil

//tekst zaszyfrowany, alokacja w pamieci procesora char * szyfrowany = (char*) malloc(size);

//zapis wyniku do pliku FILE * fp2; char * filename = "zaszyfrowany.txt"; fp = fopen(filename, "w");

//skoro ma byc to plik tekstowy to trzeba go wczytac FILE * fp; char * filename = "dozaszyfrowania.txt"; fp = fopen(filename, "r");

//czytujemy zawartosc pliku do tablicy fprintf(fp, "%s", szyfrowany); //zwalnianie pamieci procesora free(tekst); free(zaszyfrowany);

//czytujemy zawartosc pliku do tablicy fscanf(fp, "%s", tekst); //alokacja w pamieci urzadzenia char * cTekst, cSzyfrowany; //cudaMalloc(cTekst, size); //cudaMalloc(cSzyfrowany, szie); //w instrukcji jest z & mimo iż to jest wskaźnik błąd ?

//zwalnianie pamieci urzadzenia cudaFree(cTekst, size); cudaFree(cSzyfrowany, size); //zamykamy plik fclose(fp); fclose(fp2); }

2.

W oparciu o środowisko CUDA napisać równoległy program zapewniajacy transpozycje macierzy kwadratowej. Należy zaproponować kod kernela oraz jego wywołanie z kodem odpowiedzialnym za przydział pamięci z podziałem na bloki w dwóch wymiarach. #include "cuda_runtime.h" #include "device_launch_parameters.h" #include #include #include #define N 10 //w instrukcji był przykład gdzie parametry zostaly podane w ten sposób __global__ void Transponuj(float in[N] [N], float out[N][N]) { int i = threadIdx.x; int j = threadIdx.y; //transponowanie out[i][j] = in[j][i]; } cudaError_t kernel(float ** macierz, float ** macierzTransponowana) { //rozmiar do zaalokowania int size = N * sizeof(float); //macierze na urządzeniu float ** d_in; float ** d_out; //status błędów cudy cudaError_t cudaStatus; cudaStatus = cudaSetDevice(0); if(cudaStatus != cudaSuccess) goto Error; //alokacja pamieci na urzadzeni //alokacja pierwszego wymiaru cudaStatus = cudaMalloc(d_in, N * sizeof(float*)); if(cudaStatus != cudaSuccess) goto Error; //kopiowanie pierwszego wymiaru cudaStatus = cudaMemcpy(d_in, macierz, N * sizeof(float*), cudaMemcpyHostToDevice); if(cudaStatus != cudaSuccess) goto Error; //druga macierz

cudaStatus = cudaMalloc(d_out, N * sizeof(float*)); if(cudaStatus != cudaSuccess) goto Error; //alokacja drugiego wymiaru int i = 0; for(i = 0; i < N; i++) { //pierwsza macierz cudaStatus = cudaMalloc(d_in[i], size); if(cudaStatus != cudaSuccess) goto Error; //przenoszenie kopiowanie pamieci z PC do urzeadzenia cudaStatus = cudaMemcpy(d_in[i], macierz[i], size, cudaMemcpyHostToDevice); if(cudaStatus != cudaSuccess) goto Error; //druga macierz cudaStatus = cudaMalloc(d_out[i], size); if(cudaStatus != cudaSuccess) goto Error; } //zlecenie wykonania zadania int threadsPerblock = 256; int blocksPerGrid = (SIZE + threadsPerBlock - 1 ) / threadsPerBlock; Transponuj<<>>(macierz, macierzTransponowana); //kopiowanie wyników pierwszego wymiaru cudaStatus = cudaMemcpy(macierzTransponowana, d_out, N * sizeof(float*), cudaMemcpyDeviceToHost); if(cudaStatus != cudaSuccess) goto Error; for(i = 0; i < N; i++) { cudaStatus = cudaMemcpy(macierzTransponowana[i], d_out[i], size, cudaMemcpyDeviceToHost);

if(cudaStatus != cudaSuccess) goto Error;

cudaFree(d_in[i]); cudaFree(d_out[i]);

} Error: //albo też end //zwalnianie pamieci drugiego wymiaru for(i = 0; i < N; i++) {

} //zwalnianie pamieci pierwszego wymiaru cudaFree(d_in); cudaFree(d_out); }

Albo ## transpozycja macierzy kwadratowej #include "cuda_runtime.h" #include "device_launch_parameters.h"

cudaStatus = cudaMalloc((void**)&dev_in, bytes); cudaStatus = cudaMalloc((void**)&dev_out, bytes);

#include // main() jest na końcu, nie trzeba niczego wstępnie deklarować; // po każdym wywołaniu funkcji cuda... powinno się sprawdzić błąd, np.: // cudaError_t cudaStatus = cudaSetDevice(0); // if (cudaStatus != cudaSuccess) goto Error; // tutaj sprawdzany tylko raz, poglądowo. __global__ void transposeKernel(int *a_out, const int *a_in, size_t size) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int i = y * size + x; int j = x * size + y; if (i < (size * size) && j < (size * size)) { a_out[i] = a_in[j]; // zamiana miejscami i <-> j powinna dać to samo } } cudaError_t transposeWithCuda(int *a_out, const int *a_in, size_t size) { int *dev_in = 0; int *dev_out = 0; int N = size * size; int bytes = N * sizeof(int); cudaError_t cudaStatus = cudaSetDevice(0); if (cudaStatus != cudaSuccess) goto Error; // przykładowe sprawdzenie błędu

cudaStatus = cudaMemcpy(dev_in, a_in, bytes, cudaMemcpyHostToDevice); int blocksize = 16; // ustalona wartość int nblocks = N / blocksize + (N % blocksize == 0 ? 0 : 1); // ceil() dim3 dimGrid(nblocks, nblocks); dim3 dimBlock(blocksize, blocksize); transposeKernel<<>>(dev_out, dev_in, size); cudaStatus = cudaThreadSynchronize(); cudaStatus = cudaMemcpy(a_out, dev_out, bytes, cudaMemcpyDeviceToHost); Error: cudaFree(dev_in); cudaFree(dev_out); return cudaStatus; } int main() { const size_t size = 7; // ustalona wartość int a_in[size][size]; int a_out[size][size] = { 0 }; // wypełnij_danymi(); cudaError_t cudaStatus = transposeWithCuda((int*)a_out, (int*)a_in, size); if (cudaStatus != cudaSuccess) return 1; // wypisz_wynik(); cudaStatus = cudaThreadExit();

if (cudaStatus != cudaSuccess) return 1;

return 0; }

3.

Dodawanie macierzy

#include "cuda_runtime.h" #include "device_launch_parameters.h" #include cudaError_t laddWithCuda(int *c, const int *a, const int *b, size_t size); __global__ void laddKernel(int *c, const int *a, const int *b, const int size) { /* obliczenie wspolrzednych blockIdx.x*blockDim.x*blockDim.y - wybór bloku threadIdx.x*blockDim.y - wspólrzedna X watku threadIdx.y - wspólrzedna Y watku */ int i = blockIdx.x*blockDim.x*blockDim.y+threadIdx.x* blockDim.y+threadIdx.y; // jeżeli wykraczamy poza zakres tablicy to rezygnujemy z obliczeń if(i>size*size) return; c[i] = a[i] | b[i]; } /* * Glówna funkcja programu */ int main() { // Deklaracja macierzy const int arraySize = 7; int a[arraySize][arraySize]; int b[arraySize][arraySize]; int c[arraySize][arraySize] = { 0 }; // Wypelnienie macierzy A i B danymi for(int i=0; i
for(int j=0; j< arraySize; ++j) printf("%5X ",a[i][j]); printf("\t"); for(int j=0; j< arraySize; ++j) printf("%5X ",b[i][j]); printf("\t"); for(int j=0; j< arraySize; ++j) printf("%5X ",c[i][j]); printf("\n"); } // cudaThreadExit must be called before exiting in order for profiling and // tracing tools such as Parallel Nsight and Visual Profiler to show complete traces. cudaStatus = cudaThreadExit(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaThreadExit failed!"); return 1; } return 0; } //Funkcja wspomagajaca wywolanie sumy logicznej na CUDA cudaError_t laddWithCuda(int *c, const int *a, const int *b, size_t size) { int *dev_a = 0; int *dev_b = 0; int *dev_c = 0; cudaError_t cudaStatus // Choose which GPU to run on, change this on a multi-GPU system. cudaStatus = cudaSetDevice(0); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); goto Error; } // Alokacja pamiêci na potrzeby macierzy cudaStatus = cudaMalloc((void**)&dev_c, size * size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } cudaStatus = cudaMalloc((void**)&dev_a, size * size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error;

} cudaStatus = cudaMalloc((void**)&dev_b, size * size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } // Kopiowanie danych macierzy z pamiêci procesora do pamiêci GPU cudaStatus = cudaMemcpy(dev_a, a, size * size * sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } cudaStatus = cudaMemcpy(dev_b, b, size * size * sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } // Launch a kernel on the GPU with one thread for each element. dim3 tpb(4,4); //Struktura dim3 reprezentuj¹ca wymiary bloku (iloœæ w¹tków w dwóch wymiarach)

laddKernel<<<4, tpb>>>(dev_c, dev_a, dev_b,size); //wywo³anie z wykorzystaniem czterech bloków // cudaThreadSynchronize waits for the kernel to finish, and returns // any errors encountered during the launch. cudaStatus = cudaThreadSynchronize(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus); goto Error; } // Copy output vector from GPU buffer to host memory. cudaStatus = cudaMemcpy(c, dev_c, size * size * sizeof(int), cudaMemcpyDeviceToHost); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } // zwolnienie pamiêci Error: cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b); return cudaStatus; }