API tools faq
paste
Advertisement
Guest User

Untitled

a guest
Apr 30th, 2017
310
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
C++ 5.57 KB | None | 0 0
raw download clone embed print report
  1. void __global__ transpose_before_z_solve(double (*lhs_prev)[P_SIZE][P_SIZE][5],
  2.                                          double (*lhsp_prev)[P_SIZE][P_SIZE][5],
  3.                                          double (*lhsm_prev)[P_SIZE][P_SIZE][5],
  4.                                          double (*rhs_prev)[P_SIZE][P_SIZE][5],
  5.                                          double (*ws_prev)[P_SIZE][P_SIZE],
  6.                                          double (*rho_i_prev)[P_SIZE][P_SIZE],
  7.                                          double (*speed_prev)[P_SIZE][P_SIZE],
  8.                                          double *lhs_,
  9.                                          double *lhsp_,
  10.                                          double *lhsm_,
  11.                                          double *rhs,
  12.                                          double *ws,
  13.                                          double *rho_i,
  14.                                          double *speed,
  15.                                          int nx, int ny, int nz)
  16. {
  17.     register const int j = threadIdx.x + blockIdx.x * blockDim.x;
  18.     register const int i = threadIdx.y + blockIdx.y * blockDim.y;
  19.     register const int k = threadIdx.z + blockIdx.z * blockDim.z;
  20.    
  21.     if((k < nz) && (j < ny) && (i < nx))
  22.     {
  23.         #pragma unroll 5
  24.         for(int m = 0; m < 5; m++)
  25.         {
  26.             rhs(m, k, i, j) = rhs_prev[k][j][i][m];
  27.         }
  28.        
  29.         ws(k, i, j) = ws_prev[k][j][i];
  30.         rho_i(k, i, j) = rho_i_prev[k][j][i];
  31.         speed(k, i, j) = speed_prev[k][j][i];
  32.     }
  33. }
  34.  
  35. ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
  36.  
  37. void __global__ transpose_after_z_solve(double (*lhs_prev)[P_SIZE][P_SIZE][5],
  38.                                         double (*lhsp_prev)[P_SIZE][P_SIZE][5],
  39.                                         double (*lhsm_prev)[P_SIZE][P_SIZE][5],
  40.                                         double (*rhs_prev)[P_SIZE][P_SIZE][5],
  41.                                         double (*ws_prev)[P_SIZE][P_SIZE],
  42.                                         double (*rho_i_prev)[P_SIZE][P_SIZE],
  43.                                         double (*speed_prev)[P_SIZE][P_SIZE],
  44.                                         double *lhs_,
  45.                                         double *lhsp_,
  46.                                         double *lhsm_,
  47.                                         double *rhs,
  48.                                         double *ws,
  49.                                         double *rho_i,
  50.                                         double *speed,
  51.                                         int nx, int ny, int nz)
  52. {
  53.     register const int j = threadIdx.x + blockIdx.x * blockDim.x;
  54.     register const int i = threadIdx.y + blockIdx.y * blockDim.y;
  55.     register const int k = threadIdx.z + blockIdx.z * blockDim.z;
  56.    
  57.     if((k < nz) && (j < ny) && (i < nx))
  58.     {
  59.         #pragma unroll 5
  60.         for(int m = 0; m < 5; m++)
  61.         {
  62.             rhs_prev[k][j][i][m] = rhs(m, k, i, j);
  63.         }
  64.        
  65.         ws_prev[k][j][i] = ws(k, i, j);
  66.         rho_i_prev[k][j][i] = rho_i(k, i, j);
  67.         speed_prev[k][j][i] = speed(k, i, j);
  68.     }
  69. }
  70.  
  71. void __global__ transpose_before_x_solve(double (*rhs_prev)[P_SIZE][P_SIZE][5],
  72.                                          double (*us_prev)[P_SIZE][P_SIZE],
  73.                                          double (*rho_i_prev)[P_SIZE][P_SIZE],
  74.                                          double (*speed_prev)[P_SIZE][P_SIZE],
  75.                                          double *rhs,
  76.                                          double *us,
  77.                                          double *rho_i,
  78.                                          double *speed,
  79.                                          int nx, int ny, int nz)
  80. {
  81.     register const int k = threadIdx.x + blockIdx.x * blockDim.x;
  82.     register const int j = threadIdx.y + blockIdx.y * blockDim.y;
  83.     register const int i = threadIdx.z + blockIdx.z * blockDim.z;
  84.    
  85.     if((k < nz) && (j < ny) && (i < nx))
  86.     {
  87.         #pragma unroll 5
  88.         for(int m = 0; m < 5; m++)
  89.         {
  90.             rhs(m, i, j, k) = rhs_prev[k][j][i][m];
  91.         }
  92.        
  93.         us(i, j, k) = us_prev[k][j][i];
  94.         rho_i(i, j, k) = rho_i_prev[k][j][i];
  95.         speed(i, j, k) = speed_prev[k][j][i];
  96.     }
  97. }
  98.  
  99. ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
  100.  
  101. void __global__ transpose_after_x_solve(double (*rhs_prev)[P_SIZE][P_SIZE][5],
  102.                                         double (*us_prev)[P_SIZE][P_SIZE],
  103.                                         double (*rho_i_prev)[P_SIZE][P_SIZE],
  104.                                         double (*speed_prev)[P_SIZE][P_SIZE],
  105.                                         double *rhs,
  106.                                         double *us,
  107.                                         double *rho_i,
  108.                                         double *speed,
  109.                                         int nx, int ny, int nz)
  110. {
  111.     register const int k = threadIdx.x + blockIdx.x * blockDim.x;
  112.     register const int j = threadIdx.y + blockIdx.y * blockDim.y;
  113.     register const int i = threadIdx.z + blockIdx.z * blockDim.z;
  114.    
  115.     if((k < nz) && (j < ny) && (i < nx))
  116.     {
  117.         #pragma unroll 5
  118.         for(int m = 0; m < 5; m++)
  119.         {
  120.             rhs_prev[k][j][i][m] = rhs(m, i, j, k);
  121.         }
  122.        
  123.         us_prev[k][j][i] = us(i, j, k);
  124.         rho_i_prev[k][j][i] = rho_i(i, j, k);
  125.         speed_prev[k][j][i] = speed(i, j, k);
  126.     }
  127. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement
Public Pastes
Advertisement
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy.  OK, I Understand
Not a member of Pastebin yet?
Sign Up, it unlocks many cool features!