concurrent

 avatar
user_5573880
c_cpp
a year ago
18 kB
8
Indexable
/* Copyright (c) 1993-2017, NVIDIA CORPORATION. All rights reserved.
 *
 * Redistribution and use in source and binary forms, with or without
 * modification, are permitted provided that the following conditions
 * are met:
 *  * Redistributions of source code must retain the above copyright
 *    notice, this list of conditions and the following disclaimer.
 *  * Redistributions in binary form must reproduce the above copyright
 *    notice, this list of conditions and the following disclaimer in the
 *    documentation and/or other materials provided with the distribution.
 *  * Neither the name of NVIDIA CORPORATION nor the names of its
 *    contributors may be used to endorse or promote products derived
 *    from this software without specific prior written permission.
 *
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
 * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
 * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
 * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
 * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
 * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
 * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
 * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
 * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 */

// System includes
#include <stdio.h>
#include <assert.h>
// #include <curand.h>
// #include <cublas_v2.h>

// CUDA runtime
#include <cuda_runtime.h>
#include <cuda_profiler_api.h>

// Define some error checking macros.
#define cudaErrCheck(stat)                       \
   {                                             \
      cudaErrCheck_((stat), __FILE__, __LINE__); \
   }
void cudaErrCheck_(cudaError_t stat, const char *file, int line)
{
   if (stat != cudaSuccess)
   {
      fprintf(stderr, "CUDA Error: %s %s %d\n", cudaGetErrorString(stat), file, line);
   }
}

// #define cublasErrCheck(stat) { cublasErrCheck_((stat), __FILE__, __LINE__); }
// void cublasErrCheck_(cublasStatus_t stat, const char *file, int line) {
//    if (stat != CUBLAS_STATUS_SUCCESS) {
//       fprintf(stderr, "cuBLAS Error: %d %s %d\n", stat, file, line);
//    }
// }

// #define curandErrCheck(stat) { curandErrCheck_((stat), __FILE__, __LINE__); }
// void curandErrCheck_(curandStatus_t stat, const char *file, int line) {
//    if (stat != CURAND_STATUS_SUCCESS) {
//       fprintf(stderr, "cuRand Error: %d %s %d\n", stat, file, line);
//    }
// }

#include <mma.h>
using namespace nvcuda;
// Must be multiples of 16 for wmma code to work
#define MATRIX_M 32
#define MATRIX_N 32
#define MATRIX_K 32

// The only dimensions currently supported by WMMA
const int WMMA_M = 16;
const int WMMA_N = 16;
const int WMMA_K = 16;

// cuda core: CUDA kernel for matrix multiplication
__global__ void matrixMulKernel(float *A, float *B, float *C, int N)
{
   // printf("----start--kernel---sp--\n");

   int ROW = blockIdx.y * blockDim.y + threadIdx.y;
   int COL = blockIdx.x * blockDim.x + threadIdx.x;

   float tmpSum = 0.0;
   if (ROW < N && COL < N)
   {
      for (int i = 0; i < N; i++)
      {
         tmpSum += A[ROW * N + i] * B[i * N + COL];
      }
   }
   C[ROW * N + COL] = tmpSum;
   // printf("---end---kernel---sp--\n");
}

// Function to initialize the matrix with some values
// void initializeMatrix(float *mat, int N)
// {
//    for (int i = 0; i < N * N; i++)
//    {
//       mat[i] = rand() % 100; // Assign a random value for simplicity
//    }
// }

// Performs an MxNxK GEMM (C=alpha*A*B + beta*C) assuming:
//  1) Matrices are packed in memory.
//  2) M, N and K are multiples of 16.
//  3) Neither A nor B are transposed.
// Note: This is NOT a high performance example but is for demonstration purposes only
//       For a high performance code please use the GEMM provided in cuBLAS.
__global__ void wmma_example(half *a, half *b, float *c, int M, int N, int K, float alpha, float beta)
{
   // Leading dimensions. Packed with no transpositions.
   int lda = M;
   int ldb = K;
   int ldc = M;

   // Tile using a 2D grid
   int warpM = (blockIdx.x * blockDim.x + threadIdx.x) / warpSize;
   int warpN = (blockIdx.y * blockDim.y + threadIdx.y);
   
   // Declare the fragments
   wmma::fragment<wmma::matrix_a, WMMA_M, WMMA_N, WMMA_K, half, wmma::col_major> a_frag;
   wmma::fragment<wmma::matrix_b, WMMA_M, WMMA_N, WMMA_K, half, wmma::col_major> b_frag;
   wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> acc_frag;
   wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> c_frag;

   wmma::fill_fragment(acc_frag, 0.0f);

   // Loop over k
   // printf("----start--forloop-- threadIdx.x: %d, threadIdx.y: %d\n", threadIdx.x, threadIdx.y);
   for (int i = 0; i < K; i += WMMA_K)
   {
      // printf("----start--target---sp--\n");

      int aRow = warpM * WMMA_M;
      int aCol = i;

      int bRow = i;
      int bCol = warpN * WMMA_N;

      // printf("aRow: %d, aCol: %d, bRow: %d, bCol: %d, i: %d\n", aRow, aCol, bRow, bCol, i);
      // printf("---end---target---sp--\n");

      // Bounds checking
      if (aRow < M && aCol < K && bRow < K && bCol < N)
      {
         // Load the inputs
         // printf("----start--target---tensor--\n");
         wmma::load_matrix_sync(a_frag, a + aRow + aCol * lda, lda);

         wmma::load_matrix_sync(b_frag, b + bRow + bCol * ldb, ldb);

         // Perform the matrix multiplication
         wmma::mma_sync(acc_frag, a_frag, b_frag, acc_frag);
         // printf("---end---target---tensor--\n");
      }
   }

   // Load in the current value of c, scale it by beta, and add this our result scaled by alpha
   int cRow = warpM * WMMA_M;
   int cCol = warpN * WMMA_N;

   if (cRow < M && cCol < N)
   {
      wmma::load_matrix_sync(c_frag, c + cRow + cCol * ldc, ldc, wmma::mem_col_major);

#pragma unroll
      for (int i = 0; i < c_frag.num_elements; i++)
      {
         c_frag.x[i] = alpha * acc_frag.x[i] + beta * c_frag.x[i];
      }

      // Store the output
      wmma::store_matrix_sync(c + cRow + cCol * ldc, c_frag, ldc, wmma::mem_col_major);
   }
}

__global__ void convertFp32ToFp16(half *out, float *in, int n)
{
   // count cycles
   int idx = blockDim.x * blockIdx.x + threadIdx.x;
   if (idx < n)
   {
      out[idx] = in[idx];
   }
}

int main(int argc, char *argv[])
{
   // Initialize the problem
   //! concurrent
   int nkernels = 2;            // number of concurrent kernels
   int nstreams = nkernels + 1; // use one more stream than concurrent kernel
   // int nbytes = nkernels * sizeof(float_t);  // number of data bytes
   // float kernel_time = 10; // time the kernel should run in ms
   // float elapsed_time;     // timing variables
   int cuda_device = 0;

   cudaDeviceProp deviceProp;
   cudaErrCheck(cudaGetDevice(&cuda_device));

   cudaErrCheck(cudaGetDeviceProperties(&deviceProp, cuda_device));

   if ((deviceProp.concurrentKernels == 0))
   {
      printf("> GPU does not support concurrent kernel execution\n");
      printf("  CUDA kernel runs will be serialized\n");
   }
   else
   {
      printf("concurrent kernel: %d\n", deviceProp.concurrentKernels);
   }

   // printf("> Detected Compute SM %d.%d hardware with %d multi-processors\n",
   //   deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount);
   //--------------------

   // cuda core
   int N = MATRIX_N; // Define the size of the matrix
   size_t size = N * N * sizeof(float);
   float *h_A, *h_B, *h_C; // host copies of A, B, C
   float *d_A, *d_B, *d_C; // device copies of A, B, C

   // float *a_fp32;
   // float *b_fp32;
   half *a_fp16;
   half *b_fp16;
   // printf("WMMA Example2\n");

   float *c;
   //    float *c_cublas;
   float *c_wmma;

   //    float *c_host_cublas;
   float *c_host_wmma;
   // printf("WMMA Example3\n");

   // cuda core: Allocate space for host copies and setup values
   cudaErrCheck(cudaMallocHost((void **)&h_A, size));
   cudaErrCheck(cudaMallocHost((void **)&h_B, size));
   cudaErrCheck(cudaMallocHost((void **)&h_C, size));
   //    h_A = (float *)malloc(size);
   //    h_B = (float *)malloc(size);
   //    h_C = (float *)malloc(size);

   // Allocate space for device copies
   cudaErrCheck(cudaMalloc((void **)&d_A, size));
   cudaErrCheck(cudaMalloc((void **)&d_B, size));
   cudaErrCheck(cudaMalloc((void **)&d_C, size));
   //    cudaMalloc((void **)&d_A, size);
   //    cudaMalloc((void **)&d_B, size);
   //    cudaMalloc((void **)&d_C, size);

   // Initialize matrices A and B with random values
   for (int i = 0; i < N * N; i++)
   {
      h_A[i] = (float)rand() / (float)RAND_MAX * 100.0; // Assign a random float value between 0 and 100
      h_B[i] = (float)rand() / (float)RAND_MAX * 100.0; // Assign a random float value between 0 and 100
      h_C[i] = 0;
   }
   // initializeMatrix(h_A, N);
   // initializeMatrix(h_B, N);

   // Copy inputs to device
   cudaErrCheck(cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice));
   cudaErrCheck(cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice));
   cudaErrCheck(cudaMemcpy(d_C, h_C, size, cudaMemcpyHostToDevice));
   //    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
   //    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
   
   // stream create
   // allocate and initialize an array of stream handles

   //! concurrent
   cudaStream_t *streams =
       (cudaStream_t *)malloc(nstreams * sizeof(cudaStream_t));

   for (int i = 0; i < nstreams; i++)
   {
      cudaErrCheck(cudaStreamCreate(&(streams[i])));
   }

   //    curandGenerator_t gen;
   //    cublasHandle_t cublasHandle;

   cudaEvent_t startWMMA;
   cudaEvent_t stopWMMA;

   //    cudaEvent_t startcublas;
   //    cudaEvent_t stopcublas;

   cudaErrCheck(cudaEventCreate(&startWMMA));
   cudaErrCheck(cudaEventCreate(&stopWMMA));

   //    cudaErrCheck(cudaEventCreate(&startcublas));
   //    cudaErrCheck(cudaEventCreate(&stopcublas));

   //    cublasErrCheck(cublasCreate(&cublasHandle));

   // Use tensor cores
   //    cublasErrCheck(cublasSetMathMode(cublasHandle, CUBLAS_TENSOR_OP_MATH));

   // cudaErrCheck(cudaMalloc((void **)&a_fp32, MATRIX_M * MATRIX_K * sizeof(float)));
   // cudaErrCheck(cudaMalloc((void **)&b_fp32, MATRIX_K * MATRIX_N * sizeof(float)));
   cudaErrCheck(cudaMalloc((void **)&a_fp16, MATRIX_M * MATRIX_K * sizeof(half)));
   cudaErrCheck(cudaMalloc((void **)&b_fp16, MATRIX_K * MATRIX_N * sizeof(half)));

   cudaErrCheck(cudaMalloc((void **)&c, MATRIX_M * MATRIX_N * sizeof(float)));
   //    cudaErrCheck(cudaMalloc((void**)&c_cublas, MATRIX_M * MATRIX_N * sizeof(float)));
   cudaErrCheck(cudaMalloc((void **)&c_wmma, MATRIX_M * MATRIX_N * sizeof(float)));

   //    c_host_cublas = (float*)malloc(MATRIX_M * MATRIX_N * sizeof(float));
   c_host_wmma = (float *)malloc(MATRIX_M * MATRIX_N * sizeof(float));

   //    curandErrCheck(curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT));
   //    curandErrCheck(curandSetPseudoRandomGeneratorSeed(gen, 1337ULL));

   //    curandErrCheck(curandGenerateUniform(gen, a_fp32, MATRIX_M * MATRIX_K));
   //    curandErrCheck(curandGenerateUniform(gen, b_fp32, MATRIX_K * MATRIX_N));

   // curand doesn't currently support fp16 so we generate in fp32 and convert to fp16.
   printf("Converting to fp16...\n");
   // printf("Current cycle time:");

   //* sequential
   // convertFp32ToFp16<<<(MATRIX_M * MATRIX_K + 31) / 32, 32, 0, 0>>>(a_fp16, h_A, MATRIX_M * MATRIX_K);
   // convertFp32ToFp16<<<(MATRIX_K * MATRIX_N + 31) / 32, 32, 0, 0>>>(b_fp16, h_B, MATRIX_K * MATRIX_N);
   
   //! concurrent
   convertFp32ToFp16<<<(MATRIX_M * MATRIX_K + 31) / 32, 32, 0, streams[0]>>>(a_fp16, h_A, MATRIX_M * MATRIX_K);
   convertFp32ToFp16<<<(MATRIX_K * MATRIX_N + 31) / 32, 32, 0, streams[1]>>>(b_fp16, h_B, MATRIX_K * MATRIX_N);
   
   // printf("Current cycle time done:");

   // printf("Current cycle time: %f\n", getCycleTime());
   
   printf("Converting to fp16... DONE\n");

   //    curandErrCheck(curandGenerateUniform(gen, c, MATRIX_M * MATRIX_N));

   //    curandErrCheck(curandDestroyGenerator(gen));

   //    cudaErrCheck(cudaMemcpy(c_cublas, c, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToDevice));
   cudaErrCheck(cudaMemcpy(c_wmma, c, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToDevice));

   float alpha = 2.0f;
   float beta = 2.0f;

   printf("\nM = %d, N = %d, K = %d. alpha = %f, beta = %f\n\n", MATRIX_M, MATRIX_N, MATRIX_K, alpha, beta);


   // First: using WMMA
   dim3 gridDim;
   dim3 blockDim;

   // blockDim.x must be a multple of warpSize
   // 128x4 means we have 16 warps and a block computes a 64x64 output tile
   blockDim.x = 32;
   blockDim.y = 1;
   // blockDim.x = 64;
   // blockDim.y = 1;
   // blockDim.x = 128;
   // blockDim.y = 4;

   gridDim.x = (MATRIX_M + (WMMA_M * blockDim.x / 32 - 1)) / (WMMA_M * blockDim.x / 32);
   gridDim.y = (MATRIX_N + WMMA_N * blockDim.y - 1) / (WMMA_N * blockDim.y);
   printf("gridDim.x = %d, gridDim.y = %d\n", gridDim.x, gridDim.y);
   printf("Running with wmma...\n");
   cudaErrCheck(cudaEventRecord(startWMMA));

   //* sequential
   // wmma_example<<<gridDim, blockDim, 0, 0>>>(a_fp16, b_fp16, c_wmma, MATRIX_M, MATRIX_N, MATRIX_K, alpha, beta);
   // matrixMulKernel<<<gridDim, blockDim, 0, 0>>>(d_A, d_B, d_C, N);

   //! concurrent
   wmma_example<<<gridDim, blockDim, 0, streams[0]>>>(a_fp16, b_fp16, d_C, MATRIX_M, MATRIX_N, MATRIX_K, alpha, beta);
   matrixMulKernel<<<gridDim, blockDim, 0, streams[1]>>>(d_A, d_B, d_C, N);
   
   printf("Running with wmma...done\n");
   cudaErrCheck(cudaEventRecord(stopWMMA));
   cudaErrCheck(cudaEventSynchronize(stopWMMA));

   //! concurrent
   cudaStreamSynchronize(streams[0]);
   cudaStreamSynchronize(streams[1]);

   // Now using cuBLAS
   //    printf("Running with cuBLAS...\n");
   // Warm up cuBLAS run starts
   //    cublasErrCheck(cublasGemmEx(cublasHandle, CUBLAS_OP_N, CUBLAS_OP_N,
   //                 MATRIX_M, MATRIX_N, MATRIX_K,
   //                 &alpha,
   //                 a_fp16, CUDA_R_16F, MATRIX_M,
   //                 b_fp16, CUDA_R_16F, MATRIX_K,
   //                 &beta,
   //                 c_cublas, CUDA_R_32F, MATRIX_M,
   //                 CUDA_R_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP));
   // Warm up cuBLAS run ends

   // reset the c_cublas buffer
   //    cudaErrCheck(cudaMemcpy(c_cublas, c, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToDevice));

   //    cudaErrCheck(cudaEventRecord(startcublas));
   //    cublasErrCheck(cublasGemmEx(cublasHandle, CUBLAS_OP_N, CUBLAS_OP_N,
   //                 MATRIX_M, MATRIX_N, MATRIX_K,
   //                 &alpha,
   //                 a_fp16, CUDA_R_16F, MATRIX_M,
   //                 b_fp16, CUDA_R_16F, MATRIX_K,
   //                 &beta,
   //                 c_cublas, CUDA_R_32F, MATRIX_M,
   //                 CUDA_R_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP));
   //    cudaErrCheck(cudaEventRecord(stopcublas));
   //    cudaErrCheck(cudaEventSynchronize(stopcublas));

   // Error checking
   printf("\nChecking results...\n");

   //* sequential
   cudaErrCheck(cudaMemcpy(c_host_wmma, c_wmma, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToHost));
   cudaErrCheck(cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost));

   //! concurrent
   cudaErrCheck(cudaMemcpyAsync(c_host_wmma, c_wmma, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToHost, streams[nstreams - 1]));
   cudaErrCheck(cudaMemcpyAsync(h_C, d_C, size, cudaMemcpyDeviceToHost, streams[nstreams - 1]));
   
   //    cudaErrCheck(cudaMemcpy(c_host_cublas, c_cublas, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToHost));

   // 0.01% relative tolerance. 1e-5 absolute tolerance.
   //    int errors = 0;
   //    for (int i = 0; i < MATRIX_M * MATRIX_N; i++) {
   //       float v1 = c_host_wmma[i];
   //       float v2 = c_host_cublas[i];
   //       float diff  = fabs(v1 - v2);
   //       float relative_err = diff / v2;
   //       float eps = 1e-4;
   //       if ((relative_err >= eps)) {
   //          errors++;
   //          if (errors < 10) printf("%f %f\n", v1, v2);
   //       }
   //    }

   //    if (errors > 0) {
   //       printf("WMMA does not agree with cuBLAS! %d errors!\n", errors);
   //    }
   //    else {
   //       printf("Results verified: cublas and WMMA agree.\n\n");
   //       float wmmaTime;
   //       float cublasTime;
   //       cudaErrCheck(cudaEventElapsedTime(&wmmaTime, startWMMA, stopWMMA));
   //       cudaErrCheck(cudaEventElapsedTime(&cublasTime, startcublas, stopcublas));
   //       printf("wmma took %fms\n", wmmaTime);
   //       printf("cublas took %fms\n", cublasTime);

   //       printf("\nFor a faster code using wmma you should check out the cudaTensorCoreGemm sample in the CUDA Toolkit.\nThis code was written as a demo only!\n\n");
   //    }
   float milliseconds = 0;
   cudaEventElapsedTime(&milliseconds, startWMMA, stopWMMA);
   printf("Kernel execution time: %f ms\n", milliseconds);

   //! concurrent
   cudaStreamDestroy(streams[0]);
   cudaStreamDestroy(streams[1]);

   cudaErrCheck(cudaEventDestroy(startWMMA));
   cudaErrCheck(cudaEventDestroy(stopWMMA));

   //    cudaErrCheck(cudaEventDestroy(startcublas));
   //    cudaErrCheck(cudaEventDestroy(stopcublas));

   // cudaErrCheck(cudaFree(a_fp32));
   // cudaErrCheck(cudaFree(b_fp32));
   cudaErrCheck(cudaFree(a_fp16));
   cudaErrCheck(cudaFree(b_fp16));

   cudaErrCheck(cudaFree(c));
   //    cudaErrCheck(cudaFree(c_cublas));
   cudaErrCheck(cudaFree(c_wmma));

   // cuda core
   cudaFree(d_A);
   cudaFree(d_B);
   cudaFree(d_C);
   free(h_A);
   free(h_B);
   free(h_C);

   //    free(c_host_cublas);
   free(c_host_wmma);

   cudaErrCheck(cudaDeviceReset());
   return 0;
}
Editor is loading...
Leave a Comment