Advertisement
Guest User

Untitled

a guest
Apr 8th, 2020
336
0
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
C++ 2.22 KB | None | 0 0
  1. template <unsigned int blockSize>
  2. __global__ void scalar_7(float *Mat, int ind_x, int ind_y, int Len, int Size, float *resMat ){
  3.     size_t s = threadIdx.x + blockIdx.x * blockDim.x;
  4.     int sID = threadIdx.x;
  5.     int i;
  6.     int stride;
  7.     float sum=0;
  8.     int indx;
  9.     int ad_x=ind_x*Len;
  10.     int ad_y=ind_y*Len;
  11.     int loc_x;
  12.     int loc_y;
  13.    
  14.     stride = Len / blockDim.x +1;
  15.    
  16.    
  17.     // We use a single block to process single pair of vectors
  18.    
  19.     volatile __shared__ float pom[BlockSize];
  20.    
  21.     pom[sID] = 0;
  22.     // Najpierw każdy wątek czyta 'stride' wartości z wektorów x i y
  23.     // i wykonuje fragment obliczeń do iloczynu skalarnego
  24.     // Potem wszystkie wątki robią redukcję.
  25.     // Pamiętamy, że każdy wektor jest w rzeczywistości fragmentem pamięci w jednowymiarowej
  26.     // tablicy Mat. Dlatego musimy policzyć odpowiedni adres każdego elementu z tej tablicy
  27.     for (i=0;i<stride;i++) {
  28.         indx = s*stride+i; // Wyliczamy dla każdego wątku numer elementu w wektorach x i y
  29.         loc_x = ad_x + indx; // Wyliczamy numery elementów w tablicy Mat odpowiadające
  30.         loc_y = ad_y + indx; // elementom o numerze indx w wektorach x i y
  31.         if (indx<Len) {
  32.             sum+=Mat[loc_x]*Mat[loc_y];
  33.         }
  34.     }
  35.     pom[sID] = sum;
  36.     __syncthreads(); // synchronizujemy wątki - wszystkie wątki muszą skończyć swoje wczytywanie
  37.     // zanim przejdziemy dalej
  38.  
  39.     if (blockSize >= 1024) {
  40.         if (sID < 512) { pom[sID] += pom[sID + 512];} __syncthreads();
  41.     }
  42.  
  43.     if (blockSize >= 512) {
  44.         if (sID < 256) { pom[sID] += pom[sID + 256]; } __syncthreads();
  45.     }
  46.     if (blockSize >= 256) {
  47.         if (sID < 128) { pom[sID] += pom[sID + 128]; } __syncthreads();
  48.     }
  49.     if (blockSize >= 128) {
  50.         if (sID < 64) { pom[sID] += pom[sID + 64]; } __syncthreads();
  51.     }
  52.  
  53.     if(sID < 32) {
  54.         pom[sID] += pom[sID + 32];
  55.         pom[sID] += pom[sID + 16];
  56.         pom[sID] += pom[sID + 8];
  57.         pom[sID] += pom[sID + 4];
  58.         pom[sID] += pom[sID + 2];
  59.         pom[sID] += pom[sID + 1];
  60.     }
  61.  
  62.     if (sID==0) {
  63.         // wpisujemy odpowiedni wynik do tablicy wynikowej
  64.         resMat[ind_x*Size+ind_y]=pom[0];
  65.     }
  66.  
  67. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement