c++ - 在 CUDA 中增加每个线程的工作量的示例

Question

算法 :

我正在用 CUDA 编写程序，问题如下:

两个矩阵 A (n * 128) 和 B (m * 128)

我取 A 的第一行，并逐一计算该 vector 与 B 的所有行之间的距离。

我将每个距离的结果写在矩阵 C 的一行上，因此 C 的元素 C(i,j) 包含 A 的第 i 行和 B 的第 j 行之间的距离。

然后我继续处理 A 的下一行。

我是这样实现的:我有一个由 ( n * m ) 个块组成的网格，每个块有 128 个线程。 ( 1 * 128 )。

问题 :程序成功运行并获得预期结果，但执行时间仅比它的单线程 CPU 版本快 5 到 10 倍。所以我想知道 如何在减少之前增加每个线程的工作以提高性能 .

内核代码(原文:未优化)

 __global__ void EuclideanDistances( float *A, float *B , float *C , int n , int m)
{
    // SIZE is equal to 128
__shared__ float accumResult[SIZE];
float sA;
float sB;

    // MAPPING
int bx = blockIdx.x;  // n
int by = blockIdx.y;  // m
int ty = threadIdx.y; // 128
int tx = threadIdx.x; // 1


sA = A [bx * SIZE + ty];
sB = B [by * SIZE + ty];
__syncthreads();


accumResult[ty] = (sA - sB) * (sA - sB);
__syncthreads();


// Parallel tree-reduction
for (int stride = SIZE/2 ; stride > 0 ; stride >>= 1)
    if (ty < stride)
    {
        accumResult[ty] += accumResult [stride + ty];
          __syncthreads();
    }

    // Writing results to output matrix
if ((threadIdx.y == 0))
    C [bx * m + by] = accumResult[ty];
       __syncthreads();
}

更新

现在，我正在使用另一个映射:而不是采用 n 的网格来自 m块和块 128线程，我正在增加块内的线程数以减少块数。

新映射:
128块来自 8线程(总共 1024 个线程，这是最大大小)
n/8的网格来自 m/8块

不幸的是，它给出了错误的结果)。

优化内核代码(待更新)

__global__ void EuclideanDistances( float *A, float *B , float *C, int n , int m)
{
    __shared__ float accumResult[SIZE][8];
__shared__ float sA[SIZE][8];
__shared__ float sB[SIZE][8];

int bx = blockIdx.x;  // n / 8
int by = blockIdx.y;  // m / 8
int tx = threadIdx.x; // 8
int ty = threadIdx.y; // 128
int i = bx * tx * SIZE + ty;
int j = by * tx * SIZE + ty;

sA[ty][tx] = A [i];
sB[ty][tx] = B[j];
__syncthreads();


accumResult[ty][tx] = (sA[ty][tx] - sB[ty][tx]) * (sA[ty][tx] - sB[ty][tx]);
__syncthreads();

// Reduction
for (int stride = SIZE/2 ; stride > 0 ; stride>>=1)
    if (ty < stride)
    {
        accumResult[ty][tx] += accumResult [stride + ty][tx];
        __syncthreads();
    }

    C[bx *  m + by] = accumResult[0][tx];
}

主机代码(分配 + 内核调用)

    int main()
{
     int m = 20000; //MatrixA size : m * SIZE
     int n = 4000;  //MatrixB size : n * SIZE

     srand((unsigned)time(0));

     // Host Allocations
     float *matrixA = (float *) malloc (n * SIZE * sizeof(float));
     for(int i=0; i < n * SIZE; i++)
         matrixA[i] = (float) (rand()%100)+1;

     float *matrixB = (float *) malloc (m * SIZE * sizeof(float));
     for(int i=0; i < m * SIZE; i++)
         matrixB[i] = (float) (rand()%100)+1;

     float *results_kernel1 = (float *) malloc (n * m * sizeof(float));
     float *results_kernel2 = (float *) malloc (n * m * sizeof(float));


     //Device Allocation
     float *d_matrixA;
     float *d_matrixB;
     cudaMalloc((void **)&d_matrixA, n * SIZE * sizeof(float));
     cudaMalloc((void **)&d_matrixB, m * SIZE * sizeof(float));
     cudaMemcpy(d_matrixA , matrixA , n * SIZE * sizeof(float) , cudaMemcpyHostToDevice);
     cudaMemcpy(d_matrixB , matrixB , m * SIZE * sizeof(float) , cudaMemcpyHostToDevice);

     float *d_results_kernel1;
     float *d_results_kernel2;
     cudaMalloc((void **)&d_results_kernel1 , n * m * sizeof(float));
     cudaMalloc((void **)&d_results_kernel2 , n * m * sizeof(float));

     dim3 threads1 (1 , 128);
     dim3 blocks1  (n , m);
     EuclideanDistances1 <<<blocks1 , threads1>>> (d_matrixA , d_matrixB , d_results_kernel1 , n , m);
     cudaDeviceSynchronize();
     cudaMemcpy(results_kernel1 , d_results_kernel1 , n * m *sizeof(float) , cudaMemcpyDeviceToHost);
     cudaFree(d_results_kernel1);

     dim3 threads2 (8 , 128);   // 1024 threads per block (maximum)
     dim3 blocks2  (ceil((float)n/8) , ceil((float)m/8));
     EuclideanDistances2 <<<blocks2 , threads2>>> (d_matrixA , d_matrixB , d_results_kernel2 , n , m);
     cudaDeviceSynchronize();
     cudaMemcpy(results_kernel2 , d_results_kernel2 , n * m *sizeof(float) , cudaMemcpyDeviceToHost);
     cudaFree(d_results_kernel2);

     // Visualising and comparing results
     for (int i = 0 ; i < 50 ; i++)
         std::cout << "kernel1 : " << results_kernel1[i] << "  |  kernel2 : " << results_kernel2[i] << std::endl;

     free(matrixA);
     free(matrixB);
     free(results_kernel1);
     free(results_kernel2);

     return 0;
}

PS :我有 CUDA 6.0 和 NVIDIA GTX 650(计算能力 3.0)

Answer 1

您的问题似乎有两个组成部分:

为什么我的第二个内核不工作？

如何让我的代码运行得更快？

Why isn't my second kernel working?

你有几个问题:

初始计算中的索引问题i , j以及用于存储 C 的索引值(value)。

violation of usage的 _syncthreads()在条件块内

第 1 项是使代码正常工作的关键要素。

How do I make my code run faster?

这涉及更多。首先，您尝试“增加每个线程的工作量”并没有做任何类似的事情，它只是增加了每个块的线程数(从 128 到 8*128)。每个线程执行的工作量大致相同。此外，在尝试使用 2D 线程块的过程中，我相信发生了一些不好的事情:

各种合并和共享内存库冲突加载和存储模式被打破。

由于每个块所需的共享内存量，有效占用率下降。

第二个内核的净效果是执行时间大约翻倍。所以这不是我们想要的。

但是，增加每个线程的工作量可能是一个好主意，同时使用共享内存，并尝试保留良好的(全局、共享)内存访问模式，并允许增加占用率。

接下来是沿着这些路线进行的工作。以下代码修复了您的第二个内核、计时基础设施、完整数据验证以及 2 个新内核。第一个新内核(#3)就是我所说的“naive”内核。它只是为每个输出点分配一个线程，每个线程循环遍历必要的 vector ，计算其各自的结果。不使用共享内存，甚至不注重合并或任何其他优化。然而，通过对线程块配置 (16,16) -> (8,32) 线程的调整，我从@talonmies 的回答(现已删除)中观察到，该内核的执行速度比您的“快速”内核快得多(3 倍)。在进一步思考 (8,32) 观察之后，我得出结论，下一次优化尝试应该集中在:

消除使用并行归约来计算 vector 距离(即允许相邻线程使用直接 for 循环来循环遍历 vector )

从缓存中获得最大 yield

共享内存的高效使用

坚持完美全局合并/完美使用共享内存以进行所有读取和写入

第 4 项在评论中提示了“我可以转置矩阵吗？”的问题。有了这个权限，就可以重新组织数据以方便上面的第 4 项。上面的第 2 项在我的“快速”内核(#4)中通过将 B vector 加载到共享内存中来解决，同时允许缓存主要集中于缓存 A vector ，希望减少缓存抖动(A 是 2 中较小的一个 vector 数组，大约 2MB - fermi L2 是 768K，Kepler L2 是 1.5MB)。通过以转置形式传递 A，并有效地从共享内存“转置”片上 B，可以使用直接的 for 循环来计算 vector 距离，同时允许相邻线程完美合并读取和写入，以及“高效”使用共享内存(即非库冲突加载和广播读取)。

对于我的特定时间，(Quadro5000 cc2.0 GPU、CUDA 6、RHEL 5.5)我看到你的“快速”内核需要大约 2 秒，我的“天真的”内核需要大约 0.7 秒，我的“快速”内核需要大约 0.2秒，尽管有转置(A，C)数据。

编辑:我做了一项额外的优化，即让每个块一次计算多个 ( CHKSIZE ) B vector 。您可以将 CHKSIZE 设置为 1 以查看之前的结果(~0.2 秒)。我发现 4 的 CHKSIZE 有很好的改进。这是一种试图利用 A 的数据重用的攻击。通过 CHKSIZE 为 4 的额外优化，内核 4 的内核时间下降到大约 0.1 秒。

以下是代码和示例运行:

$ cat t460.cu 
#include <stdio.h>
#include <stdlib.h>
#include <iostream>

// both M and N must be evenly divisible by SIZE, M must be evenly divisible by CHKSIZE
#define SIZE 128
#define N 4000
#define M 20000
#define CHKSIZE 4

 __global__ void EuclideanDistances1( float *A, float *B , float *C , int n , int m)
{
    // SIZE is equal to 128
__shared__ float accumResult[SIZE];
float sA;
float sB;

    // MAPPING
int bx = blockIdx.x;  // n
int by = blockIdx.y;  // m
int ty = threadIdx.y; // 128
//int tx = threadIdx.x; // 1

sA = A [bx * SIZE + ty];
sB = B [by * SIZE + ty];
__syncthreads();

accumResult[ty] = (sA - sB) * (sA - sB);
__syncthreads();

// Parallel tree-reduction
for (int stride = SIZE/2 ; stride > 0 ; stride >>= 1){
    if (ty < stride)
    {
        accumResult[ty] += accumResult [stride + ty];
    }
          __syncthreads();
  }

    // Writing results to output matrix
if ((ty == 0))
    C [bx * m + by] = accumResult[ty];
       __syncthreads();
}

__global__ void EuclideanDistances2( float *A, float *B , float *C, int n , int m)
{
__shared__ float accumResult[SIZE][8];
__shared__ float sA[SIZE][8];
__shared__ float sB[SIZE][8];

int bx = blockIdx.x;  // n / 8
int by = blockIdx.y;  // m
int tx = threadIdx.x; // 8
int ty = threadIdx.y; // 128
int i = ((bx*8) + tx) * SIZE + ty;
int j = by * SIZE + ty;

sA[ty][tx] = A[i];
sB[ty][tx] = B[j];
__syncthreads();

accumResult[ty][tx] = (sA[ty][tx] - sB[ty][tx]) * (sA[ty][tx] - sB[ty][tx]);
__syncthreads();

// Reduction
for (int stride = SIZE/2 ; stride > 0 ; stride>>=1){
    if (ty < stride)
    {
        accumResult[ty][tx] += accumResult [stride + ty][tx];
    }
    __syncthreads();
  }

if (ty == 0)
    C[((bx*8)+tx) *  m + by] = accumResult[0][tx];
}
//naive kernel
__global__ void EuclideanDistances3( float *A, float *B , float *C, int n , int m){
  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  int idy = threadIdx.y+blockDim.y*blockIdx.y;
  float result = 0.0f;

  if ((idx < n) && (idy < m)){
    for (int i = 0; i < SIZE; i++){
      float temp = A[(idx*SIZE)+i] - B[(idy*SIZE)+i];
      result += temp * temp;}
    C[(idx*m) + idy] = result;
  }
}
//optimized kernel
__global__ void EuclideanDistances4( const float *A, const float *B , float *C, const int n , const int m){
  // n, A,  4000 this kernel assumes A is column-major A(SIZE, n)
  // m, B, 20000 this kernel assumes B is row-major    B(m, SIZE)
  // this kernel assumes C is column-major             C(m,n)
  // this kernel assumes number of threads per threadblock == SIZE
  // CHKSIZE is the number of B vectors that will be compute per block
  __shared__ float my_sB[CHKSIZE*SIZE];  // enough shared storage for CHKSIZE vectors of B
  int bx  = blockIdx.x; // one block per CHKSIZE rows of B (the larger input matrix)
  while ((bx*CHKSIZE) < m){ // not used, this while loop could be used to extend a block to multiple chunks
    int tx  = threadIdx.x;
    for (int i = 0; i < CHKSIZE; i++)  // load vectors of B into shared memory
      my_sB[(i*SIZE)+tx] = B[(((bx*CHKSIZE)+i)*SIZE)+tx];
    __syncthreads();
    while (tx < n){  //loop across all vectors in A
      float result[CHKSIZE];
      for (int i = 0; i < CHKSIZE; i++)
        result[i] = 0.0f;
      for (int i = 0; i < SIZE; i++){
        float Atemp = A[(n*i)+tx];
        for (int j = 0; j < CHKSIZE; j++){ // compute all CHKSIZE B vectors with read of A
          float temp = Atemp - my_sB[i + (j*SIZE)];
          result[j] += temp * temp;}}
      for (int i = 0; i < CHKSIZE; i++) // store CHKSIZE results
        C[((i+(bx*CHKSIZE))*n)+ tx] = result[i];
      tx += blockDim.x;  } // continue looping across vectors in A
    __syncthreads(); // necessary to prevent warps from racing ahead, if block looping is used
    bx += gridDim.x;}
}

float comp_euclid_sq(const float *rA, const float *rB, const int size){

  float result = 0.0f;
  float temp;
  for (int i = 0; i < size; i++){
    temp = (rA[i] - rB[i]);
    result += temp * temp;}
  return result;
}

int main()
{
     float et1=0.0f, et2=0.0f, et3=0.0f, et4=0.0f;
     cudaEvent_t start1, start2, start3,start4, stop1, stop2, stop3, stop4;
     cudaEventCreate(&start1);
     cudaEventCreate(&start2);
     cudaEventCreate(&start3);
     cudaEventCreate(&start4);
     cudaEventCreate(&stop1);
     cudaEventCreate(&stop2);
     cudaEventCreate(&stop3);
     cudaEventCreate(&stop4);

     int n = N;  //MatrixA size : n * SIZE
     int m = M; //MatrixB size : m * SIZE

     srand((unsigned)time(0));

     // Host Allocations
     float *matrixA = (float *) malloc (n * SIZE * sizeof(float));
     for(int i=0; i < n * SIZE; i++)
         matrixA[i] = (float) (rand()%100)+1;

     float *matrixB = (float *) malloc (m * SIZE * sizeof(float));
     for(int i=0; i < m * SIZE; i++)
         matrixB[i] = (float) (rand()%100)+1;

     float *results_kernel = (float *) malloc (n * m * sizeof(float));
     float *cpu_results_kernel = (float *) malloc (n * m * sizeof(float));
     for (int i = 0; i< n*m; i++)
       cpu_results_kernel[i] = comp_euclid_sq(matrixA + ((i/m)*SIZE), matrixB + (i%m)*SIZE, SIZE);

     //Device Allocation
     float *d_matrixA;
     float *d_matrixB;
     cudaMalloc((void **)&d_matrixA, n * SIZE * sizeof(float));
     cudaMalloc((void **)&d_matrixB, m * SIZE * sizeof(float));
     cudaMemcpy(d_matrixA , matrixA , n * SIZE * sizeof(float) , cudaMemcpyHostToDevice);
     cudaMemcpy(d_matrixB , matrixB , m * SIZE * sizeof(float) , cudaMemcpyHostToDevice);

     float *d_results_kernel;
     cudaMalloc((void **)&d_results_kernel , n * m * sizeof(float));


     dim3 threads1 (1 , SIZE);
     dim3 blocks1  (n , m);
     cudaEventRecord(start1);
     EuclideanDistances1 <<<blocks1 , threads1>>> (d_matrixA , d_matrixB , d_results_kernel , n , m);
     cudaEventRecord(stop1);
     cudaMemcpy(results_kernel , d_results_kernel , n * m *sizeof(float) , cudaMemcpyDeviceToHost);
     for (int i = 0; i< n*m; i++) {
       if (results_kernel[i] != cpu_results_kernel[i])  {printf("cpu/kernel1 mismatch at %d, cpu: %f, kernel1: %f\n", i, cpu_results_kernel[i], results_kernel[i]); return 1;}}
     cudaMemset(d_results_kernel, 0, n*m*sizeof(float));
     cudaEventSynchronize(stop1);
     cudaEventElapsedTime(&et1, start1, stop1);

     dim3 threads2 (8 , SIZE);   // 1024 threads per block (maximum)
     dim3 blocks2  (n/8 , m); // assumes n evenly divisible by 8
     cudaEventRecord(start2);
     EuclideanDistances2 <<<blocks2 , threads2>>> (d_matrixA , d_matrixB , d_results_kernel , n , m);
     cudaEventRecord(stop2);
     cudaMemcpy(results_kernel , d_results_kernel , n * m *sizeof(float) , cudaMemcpyDeviceToHost);
     for (int i = 0; i< n*m; i++) {
       if (results_kernel[i] != cpu_results_kernel[i])  {printf("cpu/kernel2 mismatch at %d, cpu: %f, kernel1: %f\n", i, cpu_results_kernel[i], results_kernel[i]); return 1;}}
     cudaMemset(d_results_kernel, 0, n*m*sizeof(float));
     cudaEventSynchronize(stop2);
     cudaEventElapsedTime(&et2, start2, stop2);

     cudaFuncSetCacheConfig(EuclideanDistances3, cudaFuncCachePreferL1);
     dim3 threads3 (8, 32);   // 1024 threads per block (maximum)
     dim3 blocks3  (n/threads3.x , m/threads3.y); // assumes evenly divisible
     cudaEventRecord(start3);
     EuclideanDistances3 <<<blocks3 , threads3>>> (d_matrixA , d_matrixB , d_results_kernel , n , m);
     cudaEventRecord(stop3);
     cudaMemcpy(results_kernel , d_results_kernel , n * m *sizeof(float) , cudaMemcpyDeviceToHost);
     for (int i = 0; i< n*m; i++) {
       if (results_kernel[i] != cpu_results_kernel[i])  {printf("cpu/kernel3 mismatch at %d, cpu: %f, kernel3: %f\n", i, cpu_results_kernel[i], results_kernel[i]); return 1;}}
     cudaMemset(d_results_kernel, 0, n*m*sizeof(float));
     cudaEventSynchronize(stop3);
     cudaEventElapsedTime(&et3, start3, stop3);

     // transpose matrix A
     float *matrixA_T = (float *) malloc (n * SIZE * sizeof(float));
       for (int i = 0; i < n; i++)
         for (int j = 0; j < SIZE; j++)
           matrixA_T[(j*n)+i] = matrixA[(i*SIZE)+j];
     cudaMemcpy(d_matrixA , matrixA_T , n * SIZE * sizeof(float) , cudaMemcpyHostToDevice);

     cudaFuncSetCacheConfig(EuclideanDistances4, cudaFuncCachePreferL1);
     dim3 threads4(SIZE); // one thread per vector element
     dim3 blocks4(m/CHKSIZE);
     cudaEventRecord(start4);
     EuclideanDistances4 <<<blocks4 , threads4>>> (d_matrixA , d_matrixB , d_results_kernel , n , m);
     cudaEventRecord(stop4);
     cudaMemcpy(results_kernel , d_results_kernel , n * m *sizeof(float) , cudaMemcpyDeviceToHost);
     // test for correct transposed result C(m,n)
     for (int i = 0; i< n; i++)
       for (int j = 0; j < m; j++)
         if (results_kernel[(j*n)+i] != cpu_results_kernel[(i*m)+j])  {printf("cpu/kernel4 mismatch at %d,%d, cpu: %f, kernel4: %f\n", i,j, cpu_results_kernel[(i*m)+j], results_kernel[(j*n)+i]); return 1;}
     cudaEventSynchronize(stop4);
     cudaEventElapsedTime(&et4, start4, stop4);
     cudaFree(d_results_kernel);

     printf("Success!\n");
     printf("kernel1 : %.fms, kernel2 : %.fms, kernel3 : %.fms, kernel4 : %.fms\n", et1, et2, et3, et4);

     free(matrixA);
     free(matrixB);
     free(results_kernel);

     return 0;
}

$ nvcc -O3 -arch=sm_20 -o t460 t460.cu
$ ./t460
Success!
kernel1 : 2213ms, kernel2 : 4660ms, kernel3 : 691ms, kernel4 : 99ms
$

希望这会让你有更多关于工作的想法。当然，您可能会在 cc3.0 设备上获得不同的计时。

是否可以进一步优化？大概。我要考虑的第一个目标是弄清楚如何利用 vector A 上的数据重用机会。( vector B 的数据重用已经在内核 4 中通过将其加载到共享内存中来处理。可能有是使用一些共享内存来存储 A 的一部分以使代码运行得更快的方法。)

我想我还应该提到，按照您提供的代码的指导，此代码正在计算 euclidean distance 的平方。 .对内核的一个微不足道的修改可以使它计算实际欧几里德距离( C[...] = sqrtf(...); ) 然而，我所包含的验证假设结果是“在范围内”，以便在 float 中完美存储整数数量。 .您的测试用例满足此要求，但否则需要修改验证代码(如果使用了 sqrtf)。