计算双 FMA 和共享内存延迟的准确方法答案

【问题标题】：Accurate method to calculate double FMA and Shared memory latency计算双 FMA 和共享内存延迟的准确方法
【发布时间】：2015-01-11 17:38:21
【问题描述】：

我试图想出一种准确的方法来测量两个操作的延迟： 1) 双精度 FMA 操作的延迟。 2) 从共享内存加载双精度的延迟。我正在使用 K20x，想知道这段代码是否能提供准确的测量结果。

#include <cuda.h>

#include <stdlib.h>
#include <stdio.h>
#include <iostream>

using namespace std;

//Clock rate
#define MHZ 732e6
//number of streaming multiprocessors
#define SMS 14
// number of double precision units
#define DP_UNITS 16*4
//number of shared banks
#define SHARED_BANKS 32

#define ITER 100000
#define NEARONE 1.0000000000000004

__global__ void fma_latency_kernal(double *in, double *out){
  int tid = blockIdx.x*blockDim.x+threadIdx.x;
  double val = in[tid];
#pragma unroll 100
  for(int i=0; i<ITER; i++){
    val+=val*NEARONE;
  }
  out[tid]=val;
}

__global__ void shared_latency_kernel(double *in, double *out){
  volatile extern __shared__ double smem[];
  int tid = blockIdx.x*blockDim.x+threadIdx.x;
  smem[threadIdx.x]=in[tid];
#pragma unroll 32
  for(int i=0; i<ITER; i++){
    smem[threadIdx.x]=smem[(threadIdx.x+i)%32]*NEARONE;
  }
  out[tid]=smem[threadIdx.x];
}

int main (int argc , char **argv){

  float time;
  cudaEvent_t start, stop, start2, stop2;

  double *d_A, *d_B;
  cudaMalloc(&d_A, DP_UNITS*SMS*sizeof(float));
  cudaMalloc(&d_B, DP_UNITS*SMS*sizeof(float));

  cudaError_t err;

  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start, 0);

  fma_latency_kernal<<<SMS, DP_UNITS>>>(d_A, d_B);

  cudaEventRecord(stop, 0);
  cudaEventSynchronize(stop);
  cudaEventElapsedTime(&time, start, stop);
  time/=1000;
  err = cudaGetLastError();
  if(err!=cudaSuccess)
    printf("Error FMA: %s\n", cudaGetErrorString(err));
  printf("Latency of FMA = %3.1f clock cycles\n", (time/(double)ITER)*(double)MHZ);


  cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeFourByte);
  cudaEventCreate(&start2);
  cudaEventCreate(&stop2);
  cudaEventRecord(start2, 0);

  shared_latency_kernel<<<1, SHARED_BANKS, sizeof(double)>>>(d_A, d_B );

  cudaEventRecord(stop2, 0);
  cudaEventSynchronize(stop2);
  cudaEventElapsedTime(&time, start2, stop2);
  time/=1000;
  err = cudaGetLastError();
  if(err!=cudaSuccess)
    printf("Error Shared Memory: %s\n", cudaGetErrorString(err));

  printf("Latency of Shared Memory = %3.1f clock cycles\n", time/(double)ITER*(double)MHZ);

}

我在 K20x 上的结果如下： FMA 的延迟 = 16.4 个时钟周期共享内存的延迟 = 60.7 个时钟周期这对我来说似乎是合理的，但我不确定它有多准确。

【问题讨论】：

您的结果似乎在大致范围内，但有点高。您可能需要稍微改进您的方法。根据我的性能优化工作，我建议将 SM 超额订阅约 20 倍，也就是说，运行的线程数是物理并发运行的 20 倍。这减少了 GPU 中各种开销的影响，显示出稳定的性能。您可能对以前的微基准研究感兴趣：2010 paper、2014 poster
虽然您当前的代码似乎不会受到影响，但这里有一个小警告：GPU 上的指令缓存大小很小，我认为在 4KB 到 8KB 范围内。指令很大（通常包含 8 个字节）。没有分支预测。这意味着展开的循环变得如此之大以至于它们无法完全放入指令缓存中，当它们遇到循环关闭分支时，它们将经历强制性的 ICache 未命中。根据我的实验，这可能会导致大约 3% 的性能损失（这显然因代码上下文而异，并且可能因 GPU 架构而异）。
感谢您的提醒。我会尝试玩展开。我不确定如何在超额订阅 SM 时测量延迟。如果我开始向 SM 发送许多扭曲，它们将开始重叠执行指令。在这种情况下，您如何消除延迟？还是您建议我将共享内存设置为一次将执行限制为一个扭曲？

标签： cuda latency

【解决方案1】：

在我看来，您的延迟值非常高 - 几乎是我预期的两倍。要测量某物在 GPU 上需要多少个周期，您可以在内核函数的相关部分之前和之后插入 clock() 函数。 clock 函数以 int 形式返回当前周期，因此通过从第二个值中减去第一个值，您可以得到在调度第一条时钟指令和调度第二条时钟指令之间经过的周期数。

请注意，您从此方法获得的数字将包括时钟指令本身的额外时间；我相信默认情况下，一个线程会在每条时钟指令之前和之后立即阻塞几个周期，因此您可能想尝试一下它增加了多少个周期，以便您可以将它们减去。

【讨论】：