WCudaMSE/Student_Cuda/src/cpp/core/01b_AddVector/AddVector.cu

   1 #include <iostream>
   2
   3 #include "Indice2D.h"
   4 #include "cudaTools.h"
   5 #include "Device.h"
   6
   7 using std::cout;
   8 using std::endl;
   9
  10 static __global__ void add(float* ptrDevV1, float* ptrDevV2, int n, float* ptrDevResult);
  11 static __device__ float work(float v1, float v2);
  12
  13 __global__ void add(float* ptrDevV1, float* ptrDevV2, int n, float* ptrDevResult)
  14     {
  15     const int NB_THREAD = Indice2D::nbThread();
  16     const int TID = Indice2D::tid();
  17
  18     int s = TID;
  19
  20     while (s < n)
  21         {
  22         ptrDevResult[s] = work(ptrDevV1[s], ptrDevV2[s]);
  23         s += NB_THREAD;
  24         }
  25     }
  26
  27 __device__ float work(float v1, float v2)
  28     {
  29     return v1 + v2;
  30     }
  31
  32 bool addVectors()
  33     {
  34     cout << "addVectors() ..." << endl;
  35
  36     // Inputs (passé en paramètre de la fonction dans un cas général).
  37     float v1[] = { 1, 2, 3 };
  38     float v2[] = { 10, 20, 30 };
  39
  40     // Outputs (renvoyer de la fonction dans un cas général).
  41     float vRes[3];
  42
  43     // Allocation coté GPU.
  44     float* ptrDevV1, *ptrDevV2, *ptrDevVResult = 0;
  45     const size_t vecSize = 3 * sizeof(float);
  46     HANDLE_ERROR(cudaMalloc(&ptrDevV1, vecSize));
  47     HANDLE_ERROR(cudaMalloc(&ptrDevV2, vecSize));
  48     HANDLE_ERROR(cudaMalloc(&ptrDevVResult, vecSize));
  49
  50     HANDLE_ERROR(cudaMemset(ptrDevV1, 0, vecSize));
  51     HANDLE_ERROR(cudaMemset(ptrDevV2, 0, vecSize));
  52     HANDLE_ERROR(cudaMemset(ptrDevVResult, 0, vecSize));
  53
  54     HANDLE_ERROR(cudaMemcpy(ptrDevV1, v1, vecSize, cudaMemcpyHostToDevice));
  55     HANDLE_ERROR(cudaMemcpy(ptrDevV2, v2, vecSize, cudaMemcpyHostToDevice));
  56
  57     const dim3 dg(2, 2, 1);
  58     const dim3 db(2, 2, 1);
  59     Device::assertDim(dg, db);
  60
  61     add<<<dg, db>>>(ptrDevV1, ptrDevV2, 3, ptrDevVResult);
  62
  63     // Barrière implicite de synchronisation ('cudaMemCpy').
  64     HANDLE_ERROR(cudaMemcpy(vRes, ptrDevVResult, vecSize, cudaMemcpyDeviceToHost));
  65
  66     return vRes[0] == 11 && vRes[1] == 22 && vRes[2] == 33;
  67     }