比较GPU（CUDA）中加法和除法运算的时间要求

Question

我需要找出在 GPU 中除法运算比加法运算快多少。我写了一段代码，将两个长度为 1000000 的随机向量相加，并测量了核函数的运行时间。这是我的代码

         __global__ void addition(float *a, float *b, float *c, int n)
            {

                int id = blockIdx.x*blockDim.x+threadIdx.x;


                if (id < n)
                    c[id] = a[id] + b[id];
            }
        int main( int argc, char* argv[] )
        {
           srand(time(NULL));

            int n = 1000000;
            clock_t t1,t2;

            float *h_a;
            float *h_b;
            float *h_c;
            float *d_a;
            float *d_b;
            float *d_c;

            size_t bytes = n*sizeof(float);


            h_a = (float*)malloc(bytes);
            h_b = (float*)malloc(bytes);
            h_c = (float*)malloc(bytes);

            cudaMalloc(&d_a, bytes);
            cudaMalloc(&d_b, bytes);
            cudaMalloc(&d_c, bytes);

            int i;

            for( i = 0; i < n; i++ ) {
                h_a[i] =(float)rand();
                h_b[i] =(float)rand();
            }


            cudaMemcpy( d_a, h_a, bytes, cudaMemcpyHostToDevice);
            cudaMemcpy( d_b, h_b, bytes, cudaMemcpyHostToDevice);

         int blockSize, gridSize;

              blockSize =1;
              gridSize =1;

                t1 = clock();
                addition<<<gridSize, blockSize>>>(d_a, d_b, d_d, n);
                t2 = clock();

            cudaMemcpy( h_c, d_c, bytes, cudaMemcpyDeviceToHost );

        float time = t2-t1;

            printf("Time is: %f\n", time);

            cudaFree(d_a);
            cudaFree(d_b);
            cudaFree(d_c);

            free(h_a);
            free(h_b);
            free(h_c);

     return 0;
    }

在这段代码中，我只考虑了每个块 1 个线程。我通过在内核函数中将“+”替换为“/”来编写相同的除法代码。

两个代码的运行时间几乎相同，这意味着除法和加法需要相同的时间。真的有可能吗？据我所知，除法比加法等其他操作更复杂、更慢，那么我的代码不正确吗？

Answer 1

如果您想通过计算整个内核运行时间来衡量特定操作的速度，您需要使该操作占内核的主要部分运行.

在你上面的内核代码中，有两个问题。

每个线程只会做一次你要测量的操作，但同时会访问全局内存3次，这占用了内核的大部分运行时间，还有其他开销（计算 id 并比较 id < n）；
启动内核后没有 cudaDeviceSynchronize();，您测量的是内核启动时间而不是内核运行时间。

要解决问题 1，您可以使用展开的 for 循环重复目标操作并最小化全局内存访问。例如在下面的内核中，一个有 1024 个线程的块将执行 'div' 1,024,000 x n 次并且全局内存访问仅一次。

__global__ void div(float* result, float a, float b, int n) {
  float r = 1;

  for (int i = 0; i < n; i++) {
#pragma unroll
    for (int j = 0; j < 500; j++) {
      r /= a;
      r /= b;
    }
  }
  if (threadIdx.x == 0) {
    result[blockIdx.x] = r;
  }
}

对于问题 2，不要忘记在内核在主机端启动后添加 cudaDeviceSynchronize()。

int grid = 1000;
int block = 1024;
int num_repeat = 100;
double t1;

thrust::device_vector<float> result(grid, 0.0);
float* raw_result = thrust::raw_pointer_cast(result.data());

t1 = omp_get_wtime();
div<<<grid, block>>>(raw_result, 1.25, 0.8, num_repeat);
cudaDeviceSynchronize();
t1 = omp_get_wtime() - t1;

使用上面的代码，您可以使用大的 grid dim 和 block dim 来充分利用 GPU，或者仅使用 1 thread/block 和 1 block/grid 来测量速度。他们会给出相似的结果。

在 K40c 上，加法比 div 快约 21 倍。

add: time: 0.072437s  speed: 1413.64Gflops  result[0]: 100001
div: time: 1.52071s  speed: 67.337Gflops  result[0]: 1
speed ratio (add/div): 20.9935

其实更简单的方法是使用设备定时功能clock64()和启动设置<<<1, 1>>>。它给出了~19x的速度比。

__global__ void test(float* result, float a, float b) {
  float r_add = 0, r_div = 1;

  long long t1 = clock64();
#pragma unroll
  for (int i = 0; i < 500; i++) {
    r_add += a;
    r_add += b;
  }
  long long t2 = clock64();
#pragma unroll
  for (int i = 0; i < 500; i++) {
    r_div /= a;
    r_div /= b;
  }
  long long t3 = clock64();
  long long t4 = clock64();

  result[0] = (t2 - t1) - (t4 - t3);
  result[1] = (t3 - t2) - (t4 - t3);
  result[2] = float(t3 - t2 - (t4 - t3)) / float(t2 - t1 - (t4 - t3));
  result[3] = t4 - t3;
  result[4] = r_add;
  result[5] = r_div;
}

比较GPU（CUDA）中加法和除法运算的时间要求

Comparing the time requirements of addition and division operation in GPU (CUDA)

parallel-processing

time

cuda

gpu

gpgpu