t0re199

Transfer

Oct 16th, 2021 (edited)
208
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
  1. #include "time.h"
  2. #include "stdio.h"
  3. #include "stdlib.h"
  4. #include "assert.h"
  5. #include "cuda.h"
  6.  
  7.  
  8. #define tx threadIdx.x
  9. #define ty threadIdx.y
  10.  
  11.  
  12. #define KER_MONO 0x0
  13. #define KER_GRID 0x1
  14. #define MEM_STD 0x0
  15. #define MEM_UNI 0x1
  16.  
  17. #define KER_SIZE 0x20
  18.  
  19.  
  20. const int OP1_ROWS = 2000,
  21.           OP1_COLS = 500,
  22.           OP2_COLS = 2000,
  23.           SIZE = OP1_ROWS * OP1_COLS,
  24.           ACC_SIZE = OP1_ROWS * OP2_COLS;
  25.  
  26.  
  27.  
  28. ////////////////////////////DEVICE CODE HERE////////////////////////////////////
  29. //////////////////////////////////////////////////////////////////////////////
  30. __global__
  31. void cudaFill(float *buffer, float value, const int len) //aka matrixInit
  32. {
  33.     int i = blockIdx.x * blockDim.x + threadIdx.x;
  34.  
  35.     if(i < len)
  36.     {
  37.         buffer[i] = value;
  38.     }
  39. }
  40.  
  41.  
  42. __global__
  43. void cudaMatrixMultiplication(float *op1, float *op2, float *acc, const int dim, const int op1Rows, const int op2Cols)
  44. {
  45.     int i = blockIdx.x * blockDim.x + threadIdx.x,
  46.         j = 0x0,
  47.         k = 0x0,
  48.  
  49.         rowStride = blockDim.x * gridDim.x,
  50.         colStride = blockDim.y * gridDim.y;
  51.  
  52.  
  53.     for(; i < op1Rows; i += rowStride)
  54.     {
  55.         for(j = blockIdx.y * blockDim.y + threadIdx.y; j < op2Cols; j += colStride)
  56.         {
  57.             float scalarProduct = 0.0f;
  58.  
  59.             for(k = 0x0; k < dim; k++)
  60.             {
  61.                 scalarProduct += op1[i * dim + k] * op2[k * dim + j];
  62.             }
  63.             acc[i * op1Rows + j] = scalarProduct;
  64.         }
  65.     }
  66. }
  67.  
  68.  
  69. __global__
  70. void cudaTiledMatrixMultiplication(float *op1, float *op2, float *acc, const int dim, const int op1Rows, const int op2Cols)
  71. {
  72.     __shared__ int s_op1[KER_SIZE][KER_SIZE];
  73.     __shared__ int s_op2[KER_SIZE][KER_SIZE];
  74.  
  75.     int row = blockIdx.y * blockDim.y + ty,
  76.         col = 0x0,
  77.         p = 0x0,
  78.         k = 0x0,
  79.    
  80.         rowStride = blockDim.y * gridDim.y,
  81.         colStride = blockDim.x * gridDim.x,
  82.  
  83.         rowLimit = op1Rows + (op1Rows % rowStride),
  84.         colLimit = op2Cols + (op2Cols % colStride);
  85.  
  86.     for(; row < rowLimit; row += rowStride)
  87.     {
  88.         for(col = blockIdx.x * blockDim.x + tx; col < colLimit; col += colStride)
  89.         {    
  90.             float scalarProduct = 0.0f;
  91.            
  92.             for(p = 0x0; p < (KER_SIZE + dim - 0x1) / KER_SIZE; p++) //TODO redefine condition!
  93.             {
  94.                 if(p * KER_SIZE + tx < dim && row < op1Rows)
  95.                 {
  96.                     s_op1[ty][tx] = op1[row * dim + k * KER_SIZE + tx];
  97.                 }
  98.                 else
  99.                 {
  100.                     s_op1[ty][tx] = 0.0f;
  101.                 }
  102.  
  103.                 if(p * KER_SIZE + ty < dim && col < op2Cols)
  104.                 {
  105.                     s_op2[ty][tx] = op2[(p * KER_SIZE + ty) * op2Cols + col];
  106.                 }
  107.                 else
  108.                 {
  109.                     s_op2[ty][tx] = 0.0f;
  110.                 }
  111.  
  112.                 __syncthreads();
  113.  
  114.                 for(k = 0x0; k < KER_SIZE; k++)
  115.                 {
  116.                     scalarProduct += s_op1[ty][k] * s_op2[k][tx];
  117.                 }
  118.  
  119.                 __syncthreads();
  120.             }
  121.  
  122.             if(row < op1Rows && col < op2Cols)
  123.             {
  124.                 acc[row * op2Cols + col] = scalarProduct;
  125.             }
  126.         }
  127.     }
  128. }
  129. //////////////////////////////////////////////////////////////////////////////
  130.  
  131.  
  132.  
  133. ////////////////////////////HOST CODE HERE////////////////////////////////////
  134. //////////////////////////////////////////////////////////////////////////////
  135. void matrixMultiplication(float *op1, float *op2, float *acc, const int dim, const int op1Rows, const int op2Cols);
  136. void fill(float *buffer, const float value, const int len);
  137. void assertfy(float *buffer, const float value, const int len);
  138. void printUpTo(float *buffer, const int len);
  139.  
  140.  
  141. void matrixMultiplication(float *op1, float *op2, float *acc, const int dim, const int op1Rows, const int op2Cols)
  142. {
  143.     int i = 0x0,
  144.         j = 0x0,
  145.         k = 0x0;
  146.  
  147.     for(; i < op1Rows; i++)
  148.     {
  149.         for(j = 0x0; j < op2Cols; j++)
  150.         {
  151.             float scalarProduct = 0.0f;
  152.  
  153.             for(k = 0x0; k < dim; k++)
  154.             {
  155.                 scalarProduct += op1[i * dim + k] * op2[k * dim + j];
  156.             }
  157.             acc[i * op1Rows + j] = scalarProduct;
  158.         }
  159.     }
  160. }
  161.  
  162.  
  163. void fill(float *buffer, const float value, const int len)
  164. {
  165.     int i = 0x0;
  166.     for(; i < len; i++)
  167.     {
  168.         buffer[i] = value;
  169.     }
  170. }
  171.  
  172.  
  173. void assertfy(float *buffer, const float value, const int len)
  174. {
  175.     int i = 0x0;
  176.     for(; i < len; i++)
  177.     {
  178.         if(buffer[i] != value)
  179.         {
  180.             printf("[x] Illegal Found: %1.2f\n", buffer[i]);
  181.          return;
  182.         }
  183.         assert(buffer[i] == value);
  184.     }
  185. }
  186.  
  187.  
  188. void printUpTo(float *buffer, const int len)
  189. {
  190.     int i = 0x0;
  191.     for(; i < len; i++)
  192.     {
  193.         printf("%1.2f\t", buffer[i]);
  194.     }
  195.     printf("\n");
  196. }
  197. //////////////////////////////////////////////////////////////////////////////
  198.  
  199.  
  200.  
  201. //////////////////////EXECTUTION TYPE HERE////////////////////////////////////
  202. //////////////////////////////////////////////////////////////////////////////
  203. void exec_host_only();
  204. void exec_cuda_std(const int ky, const int kx);
  205. void exec_cuda_tiled(const int ky, const int kx);
  206.  
  207.  
  208. void exec_host_only()
  209. {
  210.         float *op1,
  211.           *op2,
  212.           *acc;
  213.  
  214.     op1 = (float*) malloc(SIZE * sizeof(float));
  215.     op2 = (float*) malloc(SIZE * sizeof(float));
  216.     acc = (float*) malloc(ACC_SIZE * sizeof(float));
  217.  
  218.     clock_t begin = clock();
  219.  
  220.     fill(op1, 1.0f, SIZE);
  221.     fill(op2, 2.0f, SIZE);
  222.  
  223.     matrixMultiplication(op1, op2, acc, OP1_COLS, OP1_ROWS, OP2_COLS);
  224.  
  225.     clock_t end = clock();
  226.     float elapsedTime = (float)(end - begin) / CLOCKS_PER_SEC;
  227.  
  228.     printf("Elapsed Time: %1.2f\n", elapsedTime);
  229.  
  230.     assertfy(acc, 1000.0f, ACC_SIZE);
  231.     printUpTo(acc, 0xa);
  232.  
  233.     free(op1);
  234.     free(op2);
  235.     free(acc);
  236. }
  237.  
  238.  
  239. void exec_cuda_std(const int ky, const int kx)
  240. {
  241.     float *op1,
  242.           *op2,
  243.           *acc;
  244.  
  245.     cudaMallocManaged((void**) &op1, SIZE * sizeof(float), cudaMemAttachGlobal);
  246.     cudaMallocManaged((void**) &op2, SIZE * sizeof(float), cudaMemAttachGlobal);
  247.     cudaMallocManaged((void**) &acc, ACC_SIZE * sizeof(float), cudaMemAttachGlobal);
  248.  
  249.     dim3 blockDim(0x400, 0x1, 0x1);
  250.     dim3 gridDim(ceil(SIZE / (float) 0x400), 0x1, 0x1);
  251.  
  252.     cudaFill<<<gridDim, blockDim>>>(op1, 1.0f, SIZE);
  253.     cudaFill<<<gridDim, blockDim>>>(op2, 2.0f, SIZE);
  254.  
  255.     dim3 blockDim_(ky, kx, 0x1);
  256.     dim3 gridDim_(0x20, 0x20, 0x1);
  257.  
  258.     printf("[*] gridDim(%d, %d)\n", gridDim_.x, gridDim_.y);
  259.     printf("[*] blockDim(%d, %d)\n", blockDim_.x, blockDim_.y);
  260.  
  261.     cudaMatrixMultiplication<<<gridDim_, blockDim_>>>(op1, op2, acc, OP1_COLS, OP1_ROWS, OP2_COLS);
  262.  
  263.     cudaDeviceSynchronize();
  264.  
  265.     assertfy(acc, 1000.0f, ACC_SIZE);
  266.     printUpTo(acc, 0xa);
  267.  
  268.     cudaFree(op1);
  269.     cudaFree(op2);
  270.     cudaFree(acc);
  271. }
  272.  
  273.  
  274. void exec_cuda_tiled(const int ky, const int kx)
  275. {
  276.     float *op1,
  277.           *op2,
  278.           *acc;
  279.  
  280.     cudaMallocManaged((void**) &op1, SIZE * sizeof(float), cudaMemAttachGlobal);
  281.     cudaMallocManaged((void**) &op2, SIZE * sizeof(float), cudaMemAttachGlobal);
  282.     cudaMallocManaged((void**) &acc, ACC_SIZE * sizeof(float), cudaMemAttachGlobal);
  283.  
  284.     dim3 blockDim(0x400, 0x1, 0x1);
  285.     dim3 gridDim(ceil(SIZE / (float) 0x400), 0x1, 0x1);
  286.  
  287.     cudaFill<<<gridDim, blockDim>>>(op1, 1.0f, SIZE);
  288.     cudaFill<<<gridDim, blockDim>>>(op2, 2.0f, SIZE);
  289.  
  290.     dim3 blockDim_(ky, kx, 0x1);
  291.     dim3 gridDim_(0x200, 0x200, 0x1);
  292.  
  293.     printf("[*] gridDim(%d, %d)\n", gridDim_.x, gridDim_.y);
  294.     printf("[*] blockDim(%d, %d)\n", blockDim_.x, blockDim_.y);
  295.  
  296.     cudaTiledMatrixMultiplication<<<gridDim_, blockDim_>>>(op1, op2, acc, OP1_COLS, OP1_ROWS, OP2_COLS);
  297.  
  298.     cudaDeviceSynchronize();
  299.  
  300.     assertfy(acc, 100.0f, ACC_SIZE);
  301.     printUpTo(acc, 0xa);
  302.  
  303.     cudaFree(op1);
  304.     cudaFree(op2);
  305.     cudaFree(acc);
  306. }
  307. //////////////////////////////////////////////////////////////////////////////
  308.  
  309.  
  310.  
  311. int main(int argc, char const *argv[])
  312. {
  313.     exec_cuda_tiled(KER_SIZE, KER_SIZE);
  314.     return 0x0;
  315. }
RAW Paste Data