BookGPU/Chapters/chapter5/code/laplacian.cu

   1 template <typename T>\r
   2 __global__ void set_dirichlet_bc(T* u, int Nx)\r
   3 {\r
   4         int i = blockDim.x*blockIdx.x+threadIdx.x;\r
   5         if(i<Nx)\r
   6         {\r
   7                 u[i]           = 0.0;\r
   8                 u[(Nx-1)*Nx+i] = 0.0;\r
   9                 u[i*Nx]        = 0.0;\r
  10                 u[i*Nx+Nx-1]   = 0.0;\r
  11         }\r
  12 };\r
  13 \r
  14 template <typename T>\r
  15 struct laplacian\r
  16 {\r
  17         gpulab::FD::stencil_2d<T>       m_stencil;\r
  18 \r
  19         laplacian(int alpha) : m_stencil(2,alpha) {}\r
  20 \r
  21         template <typename V>\r
  22         void operator()(T t, V const& u, V & rhs) const\r
  23         {\r
  24                 m_stencil.mult(u,rhs); // rhs = du/dxx + du/dyy\r
  25 \r
  26                 // Make sure bc is correct\r
  27                 dim3 block = BLOCK1D(rhs.Nx());\r
  28                 dim3 grid  = GRID1D(rhs.Nx());\r
  29                 set_dirichlet_bc<<<grid,block>>>(RAW_PTR(rhs),rhs.Nx());\r
  30         }\r
  31 };