Advertisement
phystota

unrolling_kernel_coarsening

Dec 1st, 2024
26
0
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
C++ 10.04 KB | None | 0 0
  1. #include <cmath>
  2. #include <iostream>
  3. #include "gpu-new-forward.h"
  4.  
  5. #define TILE_WIDTH 16
  6. #define BLOCK_SIZE 512
  7. #define COARSE_FACTOR 3
  8.  
  9. __global__ void matrix_unrolling_kernel(const float *input, float *output,
  10.                                         const int Batch, const int Channel,
  11.                                         const int Height, const int Width,
  12.                                         const int K) {
  13.  
  14.     #define in_4d(i3, i2, i1, i0) input[(i3) * (Channel * Height * Width) + (i2) * (Height * Width) + (i1) * (Width) + i0]
  15.     #define out_3d(i1, i0) output[(i1) * (Batch * W_unroll) + i0]
  16.  
  17.     // Calculate output dimensions
  18.     const size_t Height_out = Height - K + 1;
  19.     const size_t Width_out = Width - K + 1;
  20.     const size_t W_unroll = Height_out * Width_out;
  21.     const size_t H_unroll = Channel * K * K;
  22.     const size_t W_total_unroll = Batch * W_unroll;
  23.  
  24.     // Calculate thread indices - note the COARSE_FACTOR multiplication
  25.     const size_t c_base = (blockIdx.x * blockDim.x + threadIdx.x) * COARSE_FACTOR;
  26.     const size_t hw_pos = blockIdx.y * blockDim.y + threadIdx.y;  
  27.     const size_t batch_idx = blockIdx.z * blockDim.z + threadIdx.z;
  28.  
  29.     // Extract height and width positions
  30.     const size_t h_out = hw_pos / Width_out;    
  31.     const size_t w_out = hw_pos % Width_out;    
  32.  
  33.     // First boundary check
  34.     if (h_out >= Height_out || w_out >= Width_out || batch_idx >= Batch) {
  35.         return;
  36.     }
  37.  
  38.     // Calculate shared position values
  39.     const size_t w_unroll = h_out * Width_out + w_out;
  40.     const size_t w_total_unroll = batch_idx * W_unroll + w_unroll;
  41.  
  42.     // Process COARSE_FACTOR channels in this thread
  43.     for (int cf = 0; cf < COARSE_FACTOR; cf++) {
  44.         const size_t c = c_base + cf;
  45.         if (c >= Channel) continue;  // Skip if we've exceeded channel count
  46.        
  47.         const size_t w_base = c * K * K;
  48.  
  49.         // Perform unrolling for this channel
  50.         for (int p = 0; p < K; p++) {
  51.             for (int q = 0; q < K; q++) {
  52.                 int h_unroll = w_base + p * K + q;
  53.                 out_3d(h_unroll, w_total_unroll) = in_4d(batch_idx, c, h_out + p, w_out + q);
  54.             }
  55.         }
  56.     }
  57.  
  58.     #undef in_4d
  59.     #undef out_3d
  60. }
  61.  
  62. // Tiled matrix multiplication kernel. Computes C = AB
  63. // You don't need to modify this kernel.
  64. __global__ void matrixMultiplyShared(const float *A, const float *B, float *C,
  65.                                      int numARows, int numAColumns,
  66.                                      int numBRows, int numBColumns,
  67.                                      int numCRows, int numCColumns)
  68. {
  69.     __shared__ float tileA[TILE_WIDTH][TILE_WIDTH];
  70.     __shared__ float tileB[TILE_WIDTH][TILE_WIDTH];
  71.  
  72.     int by = blockIdx.y, bx = blockIdx.x, ty = threadIdx.y, tx = threadIdx.x;
  73.  
  74.     int row = by * TILE_WIDTH + ty, col = bx * TILE_WIDTH + tx;
  75.     float val = 0;
  76.  
  77.     for (int tileId = 0; tileId < (numAColumns - 1) / TILE_WIDTH + 1; tileId++) {
  78.         if (row < numARows && tileId * TILE_WIDTH + tx < numAColumns) {
  79.             tileA[ty][tx] = A[(size_t) row * numAColumns + tileId * TILE_WIDTH + tx];
  80.         } else {
  81.             tileA[ty][tx] = 0;
  82.         }
  83.         if (col < numBColumns && tileId * TILE_WIDTH + ty < numBRows) {
  84.             tileB[ty][tx] = B[((size_t) tileId * TILE_WIDTH + ty) * numBColumns + col];
  85.         } else {
  86.             tileB[ty][tx] = 0;
  87.         }
  88.         __syncthreads();
  89.  
  90.         if (row < numCRows && col < numCColumns) {
  91.             for (int i = 0; i < TILE_WIDTH; i++) {
  92.                 val += tileA[ty][i] * tileB[i][tx];
  93.             }
  94.         }
  95.         __syncthreads();
  96.     }
  97.  
  98.     if (row < numCRows && col < numCColumns) {
  99.         C[row * numCColumns + col] = val;
  100.     }
  101. }
  102.  
  103. // Permutes the matmul result.
  104. // The output feature map after matmul is of shape Map_out x Batch x Height_out x Width_out,
  105. // and we need to permute it into Batch x Map_out x Height_out x Width_out.
  106. // You don't need to modify this kernel.
  107. __global__ void matrix_permute_kernel(const float *input, float *output, int Map_out,
  108.                                       int Batch, int image_size) {
  109.     int b = blockIdx.y;
  110.     int x = blockIdx.x * BLOCK_SIZE + threadIdx.x;
  111.     if (x < image_size) {
  112.         for (int m = 0; m < Map_out; m++) {
  113.             output[b * Map_out * image_size + m * image_size + x] =
  114.                     input[m * Batch * image_size + b * image_size + x];
  115.         }
  116.     }
  117. }
  118.  
  119. __host__ void GPUInterface::conv_forward_gpu_prolog(const float *host_output, const float *host_input, const float *host_mask, float **device_output_ptr, float **device_input_ptr, float **device_mask_ptr, const int Batch, const int Map_out, const int Channel, const int Height, const int Width, const int K)
  120. {
  121.     // TODO: Allocate memory and copy over the relevant data structures to the GPU
  122.  
  123.     // We pass double pointers for you to initialize the relevant device pointers,
  124.     //  which are passed to the other two functions.
  125.  
  126.     // Useful snippet for error checking
  127.     // cudaError_t error = cudaGetLastError();
  128.     // if(error != cudaSuccess)
  129.     // {
  130.     //     std::cout<<"CUDA error: "<<cudaGetErrorString(error)<<std::endl;
  131.     //     exit(-1);
  132.     // }
  133.  
  134.     //  allocating memory
  135.  
  136.     // Calculate sizes
  137.     const int Height_out = Height - K + 1;
  138.     const int Width_out = Width - K + 1;
  139.    
  140.     const int input_size = Batch * Channel * Height * Width * sizeof(float);
  141.     const int mask_size = Map_out * Channel * K * K * sizeof(float);
  142.     const int output_size = Batch * Map_out * Height_out * Width_out * sizeof(float);
  143.  
  144.     cudaMalloc((void**)device_input_ptr, input_size);
  145.     cudaMalloc((void**)device_mask_ptr, mask_size);
  146.     cudaMalloc((void**)device_output_ptr, output_size);
  147.  
  148.     cudaMemcpy(*device_input_ptr, host_input, input_size, cudaMemcpyHostToDevice);
  149.     cudaMemcpy(*device_mask_ptr, host_mask, mask_size, cudaMemcpyHostToDevice);
  150.  
  151. }
  152.  
  153.  
  154. __host__ void GPUInterface::conv_forward_gpu(float *device_output, const float *device_input, const float *device_mask, const int Batch, const int Map_out, const int Channel, const int Height, const int Width, const int K)
  155. {
  156.     const int Height_out = Height - K + 1;
  157.     const int Width_out = Width - K + 1;
  158.     const int Height_unrolled = Channel * K * K;
  159.     const int Width_unrolled = Batch * Height_out * Width_out;
  160.  
  161.     //allocating temping storage of unrolling matrix
  162.     float *unrolled_matrix;  // Pointer to device memory for storing the unrolled matrix
  163.     float *matmul_output;    // Pointer to device memory for storing the result of matrix multiplication
  164.     cudaMalloc((void**)&unrolled_matrix, (size_t) Batch * Channel * K * K * Height_out * Width_out * sizeof(float));
  165.     cudaMalloc((void**)&matmul_output, (Batch * Map_out * Height_out * Width_out) * sizeof(float));
  166.  
  167.     // TODO: Set the kernel dimensions and call the matrix unrolling kernel.
  168.     // dim3 gridDim((Channel * Width_unrolled + BLOCK_SIZE - 1) / BLOCK_SIZE, Batch, 1);
  169.     dim3 blockDim(4, 256, 1);  
  170.     dim3 gridDim(
  171.     (Channel + blockDim.x * COARSE_FACTOR - 1) / (blockDim.x * COARSE_FACTOR),                    // Maps dimension
  172.     (Height_out * Width_out + blockDim.y - 1) / blockDim.y,     // Combined Height/Width
  173.     ceil(1.0*Batch/blockDim.z));                                                      // Batch dimension
  174.  
  175.  
  176.     matrix_unrolling_kernel<<<gridDim, blockDim>>>(device_input, unrolled_matrix, Batch, Channel, Height, Width, K);
  177.  
  178.     // TODO: Set the kernel dimensions and call the matmul kernel
  179.     dim3 dimGrid((Width_unrolled - 1)/TILE_WIDTH + 1, (Map_out - 1)/TILE_WIDTH + 1, 1);
  180.     dim3 dimBlock(TILE_WIDTH, TILE_WIDTH, 1);
  181.     matrixMultiplyShared<<<dimGrid, dimBlock>>>(device_mask, unrolled_matrix, matmul_output, Map_out, Height_unrolled, Height_unrolled, Width_unrolled,
  182.     Map_out, Width_unrolled);
  183.  
  184.     // Permute the result of matrix multiplication
  185.     const int out_image_size = Height_out * Width_out;
  186.     dim3 permute_kernel_grid_dim((out_image_size - 1) / BLOCK_SIZE + 1, Batch, 1);
  187.     matrix_permute_kernel<<<permute_kernel_grid_dim, BLOCK_SIZE>>>(matmul_output, device_output, Map_out, Batch, out_image_size);
  188.  
  189.     cudaFree(matmul_output);
  190.     cudaFree(unrolled_matrix);
  191. }
  192.  
  193.  
  194. __host__ void GPUInterface::conv_forward_gpu_epilog(float *host_output, float *device_output, float *device_input, float *device_mask, const int Batch, const int Map_out, const int Channel, const int Height, const int Width, const int K)
  195. {
  196.  
  197.     // Calculate output size
  198.     const int Height_out = Height - K + 1;
  199.     const int Width_out = Width - K + 1;
  200.     const int output_size = Batch * Map_out * Height_out * Width_out * sizeof(float);
  201.  
  202.     // TODO: Copy the output back to host
  203.     cudaMemcpy(host_output, device_output, output_size, cudaMemcpyDeviceToHost);
  204.  
  205.     // TODO: Free device memory
  206.     cudaFree(device_output);
  207.     cudaFree(device_input);
  208.     cudaFree(device_mask);
  209. }
  210.  
  211.  
  212. __host__ void GPUInterface::get_device_properties()
  213. {
  214.     int deviceCount;
  215.     cudaGetDeviceCount(&deviceCount);
  216.  
  217.     for(int dev = 0; dev < deviceCount; dev++)
  218.     {
  219.         cudaDeviceProp deviceProp;
  220.         cudaGetDeviceProperties(&deviceProp, dev);
  221.  
  222.         std::cout<<"Device "<<dev<<" name: "<<deviceProp.name<<std::endl;
  223.         std::cout<<"Computational capabilities: "<<deviceProp.major<<"."<<deviceProp.minor<<std::endl;
  224.         std::cout<<"Max Global memory size: "<<deviceProp.totalGlobalMem<<std::endl;
  225.         std::cout<<"Max Constant memory size: "<<deviceProp.totalConstMem<<std::endl;
  226.         std::cout<<"Max Shared memory size per block: "<<deviceProp.sharedMemPerBlock<<std::endl;
  227.         std::cout<<"Max threads per block: "<<deviceProp.maxThreadsPerBlock<<std::endl;
  228.         std::cout<<"Max block dimensions: "<<deviceProp.maxThreadsDim[0]<<" x, "<<deviceProp.maxThreadsDim[1]<<" y, "<<deviceProp.maxThreadsDim[2]<<" z"<<std::endl;
  229.         std::cout<<"Max grid dimensions: "<<deviceProp.maxGridSize[0]<<" x, "<<deviceProp.maxGridSize[1]<<" y, "<<deviceProp.maxGridSize[2]<<" z"<<std::endl;
  230.         std::cout<<"Warp Size: "<<deviceProp.warpSize<<std::endl;
  231.     }
  232. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement