【问题标题】:Low DRAM utilization with shared memory共享内存的 DRAM 利用率低
【发布时间】:2013-07-21 18:04:07
【问题描述】:

我正在 GPU 上实现一个简单的几何布朗运动。 我的代码运行良好,即给出了正确的值。我担心的是关于我得到的加速,我期待更多。 到目前为止,我有 2 个实现,一个只访问全局内存,速度提高了大约 3 倍,第二个是使用共享内存,速度提高了大约 2.3 倍。

我的问题是在使用 Nvidia Visual Profiler 分析应用程序后提出的。根据它,我的加载/存储效率为 100%,但 DRAM 利用率非常低(约 10%),并且由于非合并访问,全局内存重放率几乎为 50%。

有一次我看到我一直试图使用共享内存来避免全局内存访问,但令我惊讶的是 DRAM 变低了 (4.5%) 并且全局内存重放到了 46.3%

我注意到我的内核启动中的占用率很低,因为我几乎使用了每个块的所有可用共享内存,但我不知道这是否可以解释第二种方法的性能较差。

您能否就性能方面可能发生的情况提供一些建议,以及我可以在哪里/可以寻找什么来尝试改进它?

CUDA_IMPLEMENTATION.CU

#define BLOCK_SIZE  64

#define SHMEM_ROWS  7       //The same as c_numTimeSteps = numTimeSteps
#define SHMEM_COLS  BLOCK_SIZE

__constant__ double c_c1;
__constant__ double c_c2;
__constant__ int c_numTimeSteps;
__constant__ int c_numPaths;
__constant__ double c_timeNodes[2000];

__global__
void kernelSharedMem(double *rv, double *pb)
{
    __shared__ double sh_rv[SHMEM_ROWS*SHMEM_COLS];
    __shared__ double sh_pb[(SHMEM_ROWS+1)*SHMEM_COLS];

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

    //The idea of this outter loop is to have tiles along the rows
    for(int tb = 0; tb < c_numTimeSteps; tb += SHMEM_ROWS)
    {
        //Copy values into shared memory
        for(int is = tb, isSh = 0;
            is < tb+SHMEM_ROWS && is < c_numTimeSteps;
            is++, isSh++)
        {
            sh_rv[isSh*SHMEM_COLS+threadIdx.x] = 
                rv[is*c_numPaths+p];
        }

        sh_pb[threadIdx.x] = pb[tb*numPaths+p];

        __syncthreads();

        //Main computation in SHARED MEMORY
        for(int isSh = 0; isSh < SHMEM_ROWS; isSh++)
        {
            double dt = c_timeNodes[isSh];
            double sdt = sqrt(dt) * c_c1;
            double mdt = c_c2 * dt;

            sh_pb[(isSh+1)*SHMEM_COLS+threadIdx.x] =
                sh_pb[isSh*SHMEM_COLS+threadIdx.x] *
                exp(mdt + sdt * rv[isSh*SHMEM_COLS+threadIdx.x]);

        }

        __syncthreads();

        for(int is = tb, isSh = 0;
            is < tb+SHMEM_ROWS && is < c_numTimeSteps;
            is++, isSh++)
        {
            pb[(is+1)*c_numPaths+p] = 
                sh_pb[(isSh+1)*SHMEM_COLS+threadIdx.x];
        }

    }

}

__global__
void kernelGlobalMem(double *rv, double *pb)
{
    int p = blockDim.x * blockIdx.x + threadIdx.x;

    for(int i = 0; i < c_numTimeSteps; i++)
    {
        double dt = c_timeNodes[i];
        double sdt = sqrt(dt) * c_c1;
        double mdt = c_c2 * dt;

        pb[(i+1)*c_numPaths+p] = 
            pb[i*c_numPaths+p] *
            exp(mdt + sdt * rv[i*c_numPaths+p]);

    }

}

extern "C" void computePathGpu(vector<vector<double>>* rv,
                                vector<vector<double>>* pb,
                                int numTimeSteps, int numPaths,
                                vector<double> timeNodes,
                                double c1, double c2)
{

    cudaMemcpyToSymbol(c_c1, &c1, sizeof(double));
    cudaMemcpyToSymbol(c_c2, &c2, sizeof(double));
    cudaMemcpyToSymbol(c_numTimeSteps, &numTimeSteps, sizeof(int));
    cudaMemcpyToSymbol(c_numPaths, &numPaths, sizeof(int));
    cudaMemcpyToSymbol(c_timeNodes, &(timeNodes[0]), sizeof(double)*numTimeSteps);

    double *d_rv;
    double *d_pb;

    cudaMalloc((void**)&d_rv, sizeof(double)*numTimeSteps*numPaths);
    cudaMalloc((void**)&d_pb, sizeof(double)*(numTimeSteps+1)*numPaths);

    vector<vector<double>>::iterator itRV;
    vector<vector<double>>::iterator itPB;

    double *dst = d_rv;
    for(itRV = rv->begin(); itRV != rv->end(); ++itRV)
    {
        double *src = &((*itRV)[0]);
        size_t s = itRV->size();
        cudaMemcpy(dst, src, sizeof(double)*s, cudaMemcpyHostToDevice);
        dst += s;
    }

    cudaMemcpy(d_pb, &((*(pb->begin()))[0]),
        sizeof(double)*(pb->begin())->size(), cudaMemcpyHostToDevice);

    dim3 block(BLOCK_SIZE);
    dim3  grid((numPaths+BLOCK_SIZE-1)/BLOCK_SIZE);

    kernelGlobalMem<<<grid, block>>>(d_rv, d_pb);
    //kernelSharedMem<<<grid, block>>>(d_rv, d_pb);
    cudaDeviceSynchronize();

    dst = d_pb;
    for(itPB = ++(pb->begin()); itPB != pb->end(); ++itPB)
    {
        double *src = &((*itPB)[0]);
        size_t s = itPB->size();
        dst += s;
        cudaMemcpy(src, dst, sizeof(double)*s, cudaMemcpyDeviceToHost);
    }

    cudaFree(d_pb);
    cudaFree(d_rv);

}

MAIN.CPP

extern "C" void computeOnGPU(vector<vector<double>>* rv,
                                vector<vector<double>>* pb,
                                int numTimeSteps, int numPaths,
                                vector<double> timeNodes,
                                double c1, double c2);

int main(){

    int numTimeSteps = 7;
    int numPaths = 2000000;

    vector<vector<double>> rv(numTimeSteps, vector<double>(numPaths));
    //Fill rv

    vector<double> timeNodes(numTimeSteps);
    //Fill timeNodes

    vector<vector<double>> pb(numTimeSteps, vector<double>(numPaths, 0));

    computeOnGPU(&rv, &pb, numTimeSteps, numPaths, timeNodes, 0.2, 0.123);

}

【问题讨论】:

    标签: cuda


    【解决方案1】:

    正如其他人指出的那样,共享内存版本根本不会改变全局内存访问模式,并且内核中线程之间实际上没有数据重用。所以合并问题并没有解决,您实际上所做的只是添加共享内存访问和几个同步点作为开销。

    但是看看内核到底在做什么。内核以双精度工作,这在消费卡上很慢,并且在计算循环中具有相当合理的操作数量,这很好。如果无法访问编译器,我猜大约一半的总时间​​是浮点计算在exp 调用中,一半在sqrt 调用中。这可能不应该是消费者 GPU 上的内存绑定内核。但是大约一半的双精度操作只是每个线程计算 same sqrt(dt) 值。这是对周期的巨大浪费。为什么不让内核在“无量纲”sqrt(dt) 域中迭代。这意味着您在主机上预先计算(最多)2000 个sqrt(dt) 值并将它们存储在常量内存中。然后内核循环可以写成这样的:

    double pb0 = pb[p];
    for(int i = 0; i < c_numTimeSteps; i++)
    {
        double sdt = c_stimeNodes[i]; // sqrt(dt)
        double mdt = c_c2 * sdt * sdt;
        sdt *= c_c1;
    
        double pb1 = pb0 *  exp(mdt + sdt * rv[p]);
    
        p += c_numPaths;
        pb[p] = pb1;
        pb0 = pb1;
    }
    

    [免责声明:早上 5 点写在拉普兰中部的 ipad 中。使用风险自负]

    这样做会将 sqrt 替换为乘法,这大大减少了操作。 请注意,我还冒昧地将索引计算简化为每个循环添加一个整数。编译器非常聪明,但你可以让它的工作变得像你想要的那样简单或困难。我怀疑像上面这样的循环会比你现在的要快得多。

    【讨论】:

      【解决方案2】:

      在我的 Tesla M2090 上分析您的代码后,我认为我们应该重新排序这些答案提供的所有这些建议。

      1. 尝试减少内存复制时间。 97% 的时间花在 memcopy 上,包括 H2D 和 D2H。由于你使用的是pageable memcpy,所以速度是2.5G/s~3G/s。您可以使用pinned mem cpy 将速度提高一倍。可以应用零拷贝和其他Mem optimization 技术来进一步提高内存拷贝速度。

      2. 将 sqrt() 移出内核。您可以在 CPU 上执行 7 次 sqrt(),而不是在 GPU 上执行 7 x 2,000,000 次。但是,由于您的内核很小(占computePathGpu() 总时间的 3%),因此不会产生太大影响。

      3. 减少全局内存访问。在您的代码中,您只需读取一次rv,读取一次pb,然后写入一次pb。但是,在调用 kenel 之前,只有 pb 的第一行包含有用的数据。所以整个pb的读取可以通过使用寄存器来消除。解决方案在代码中提供。

      4. 关于非合并mem访问,可以找讨论here。您的案例属于“顺序但未对齐的访问模式”。使用cudaMallocPitch() 的解决方案如下所述,并在以下代码中提供。

      注意:您提到您的 DRAM 利用率较低(大约 10%),但在我的设备上进行分析是可以的(55.8%)。可能是我的设备有点旧(M2090 CC2.0)

      #include <vector>
      
      using namespace std;
      
      #define BLOCK_SIZE  64
      #define BLOCK_SIZE_OPT  256
      
      __constant__ double c_c1;
      __constant__ double c_c2;
      __constant__ int c_numTimeSteps;
      __constant__ int c_numPaths;
      __constant__ double c_timeNodes[2000];
      
      __global__ void kernelGlobalMem(double *rv, double *pb)
      {
          int p = blockDim.x * blockIdx.x + threadIdx.x;
      
          for (int i = 0; i < c_numTimeSteps; i++)
          {
              double dt = c_timeNodes[i];
              double sdt = sqrt(dt) * c_c1;
              double mdt = c_c2 * dt;
      
              pb[(i + 1) * c_numPaths + p] =
                      pb[i * c_numPaths + p] *
                              exp(mdt + sdt * rv[i * c_numPaths + p]);
      
          }
      
      }
      
      __global__ void kernelGlobalMemOpt(double *rv, double *pb, const size_t ld_rv, const size_t ld_pb)
      {
          int p = blockDim.x * blockIdx.x + threadIdx.x;
      
          double pb0 = pb[p];
          for (int i = 0; i < c_numTimeSteps; i++)
          {
              double dt = c_timeNodes[i];
              double sdt = dt * c_c1;
              double mdt = c_c2 * dt * dt;
      
              pb0 *= exp(mdt + sdt * rv[i * ld_rv + p]);
              pb[(i + 1) * ld_pb + p] = pb0;
          }
      }
      
      void computePathGpu(vector<vector<double> >* rv,
              vector<vector<double> >* pb,
              int numTimeSteps, int numPaths,
              vector<double> timeNodes,
              double c1, double c2)
      {
      
          cudaMemcpyToSymbol(c_c1, &c1, sizeof(double));
          cudaMemcpyToSymbol(c_c2, &c2, sizeof(double));
          cudaMemcpyToSymbol(c_numTimeSteps, &numTimeSteps, sizeof(int));
          cudaMemcpyToSymbol(c_numPaths, &numPaths, sizeof(int));
          cudaMemcpyToSymbol(c_timeNodes, &(timeNodes[0]), sizeof(double) * numTimeSteps);
      
          double *d_rv;
          double *d_pb;
      
          cudaMalloc((void**) &d_rv, sizeof(double) * numTimeSteps * numPaths);
          cudaMalloc((void**) &d_pb, sizeof(double) * (numTimeSteps + 1) * numPaths);
      
          vector<vector<double> >::iterator itRV;
          vector<vector<double> >::iterator itPB;
      
          double *dst = d_rv;
          for (itRV = rv->begin(); itRV != rv->end(); ++itRV)
          {
              double *src = &((*itRV)[0]);
              size_t s = itRV->size();
              cudaMemcpy(dst, src, sizeof(double) * s, cudaMemcpyHostToDevice);
              dst += s;
          }
      
          cudaMemcpy(d_pb, &((*(pb->begin()))[0]),
                  sizeof(double) * (pb->begin())->size(), cudaMemcpyHostToDevice);
      
          dim3 block(BLOCK_SIZE);
          dim3 grid((numPaths + BLOCK_SIZE - 1) / BLOCK_SIZE);
      
          kernelGlobalMem<<<grid, block>>>(d_rv, d_pb);
          //kernelSharedMem<<<grid, block>>>(d_rv, d_pb);
          cudaDeviceSynchronize();
      
          dst = d_pb;
          for (itPB = ++(pb->begin()); itPB != pb->end(); ++itPB)
          {
              double *src = &((*itPB)[0]);
              size_t s = itPB->size();
              dst += s;
              cudaMemcpy(src, dst, sizeof(double) * s, cudaMemcpyDeviceToHost);
          }
      
          cudaFree(d_pb);
          cudaFree(d_rv);
      
      }
      
      void computePathGpuOpt(vector<vector<double> >* rv,
              vector<vector<double> >* pb,
              int numTimeSteps, int numPaths,
              vector<double> timeNodes,
              double c1, double c2)
      {
          for(int i=0;i<timeNodes.size();i++)
          {
              timeNodes[i]=sqrt(timeNodes[i]);
          }
      
          cudaMemcpyToSymbol(c_c1, &c1, sizeof(double));
          cudaMemcpyToSymbol(c_c2, &c2, sizeof(double));
          cudaMemcpyToSymbol(c_numTimeSteps, &numTimeSteps, sizeof(int));
          cudaMemcpyToSymbol(c_numPaths, &numPaths, sizeof(int));
          cudaMemcpyToSymbol(c_timeNodes, &(timeNodes[0]), sizeof(double) * numTimeSteps);
      
          double *d_rv;
          double *d_pb;
          size_t ld_rv, ld_pb;
      
          cudaMallocPitch((void **) &d_rv, &ld_rv, sizeof(double) * numPaths, numTimeSteps);
          cudaMallocPitch((void **) &d_pb, &ld_pb, sizeof(double) * numPaths, numTimeSteps + 1);
          ld_rv /= sizeof(double);
          ld_pb /= sizeof(double);
      
      //  cudaMalloc((void**) &d_rv, sizeof(double) * numTimeSteps * numPaths);
      //  cudaMalloc((void**) &d_pb, sizeof(double) * (numTimeSteps + 1) * numPaths);
      
          vector<vector<double> >::iterator itRV;
          vector<vector<double> >::iterator itPB;
      
          double *dst = d_rv;
          for (itRV = rv->begin(); itRV != rv->end(); ++itRV)
          {
              double *src = &((*itRV)[0]);
              size_t s = itRV->size();
              cudaMemcpy(dst, src, sizeof(double) * s, cudaMemcpyHostToDevice);
              dst += ld_rv;
          }
      
          cudaMemcpy(d_pb, &((*(pb->begin()))[0]),
                  sizeof(double) * (pb->begin())->size(), cudaMemcpyHostToDevice);
      
          dim3 block(BLOCK_SIZE_OPT);
          dim3 grid((numPaths + BLOCK_SIZE_OPT - 1) / BLOCK_SIZE_OPT);
      
          kernelGlobalMemOpt<<<grid, block>>>(d_rv, d_pb, ld_rv, ld_pb);
          //kernelSharedMem<<<grid, block>>>(d_rv, d_pb);
          cudaDeviceSynchronize();
      
          dst = d_pb;
          for (itPB = ++(pb->begin()); itPB != pb->end(); ++itPB)
          {
              double *src = &((*itPB)[0]);
              size_t s = itPB->size();
              dst += ld_pb;
              cudaMemcpy(src, dst, sizeof(double) * s, cudaMemcpyDeviceToHost);
          }
      
          cudaFree(d_pb);
          cudaFree(d_rv);
      
      }
      
      int main()
      {
      
          int numTimeSteps = 7;
          int numPaths = 2000000;
      
          vector<vector<double> > rv(numTimeSteps, vector<double>(numPaths));
          vector<double> timeNodes(numTimeSteps);
          vector<vector<double> > pb(numTimeSteps, vector<double>(numPaths, 0));
          vector<vector<double> > pbOpt(numTimeSteps, vector<double>(numPaths, 0));
          computePathGpu(&rv, &pb, numTimeSteps, numPaths, timeNodes, 0.2, 0.123);
          computePathGpuOpt(&rv, &pbOpt, numTimeSteps, numPaths, timeNodes, 0.2, 0.123);
      }
      

      您的每个 cuda 线程为所有时间步计算一条路径。根据您的 GlobalMem 代码,您不会在路径之间共享任何数据。所以不需要共享内存。

      对于 nvprof 检测到的非合并访问问题,这是因为您的数据 pb 和 rv 没有很好地对齐。 pb 和 rv 可以看作是大小为 [time steps x #paths] 的矩阵。由于您的#path 不是缓存行的倍数,因此从第二行开始,即时间步长,所有全局内存访问都是非合并的。如果您的 CUDA 设备较旧,则会导致 50% 的内存重放。较新的设备不会受到这种非合并访问的影响。

      解决方案很简单。您只需将填充字节添加到行的每一端,以便每一行都可以从合并的 DRAM 地址开始。这可以由cudaMallocPitch()自动完成

      还有一个问题。在您的代码中,您只需读取一次 rv,读取一次 pb 并写入一次 pb。 但是,在调用 kenel 之前,您的 pb 不包含任何有用的数据。所以使用寄存器可以消除对pb的读取,除了解决非合并访问问题之外,您还可以额外提高50%的速度。

      【讨论】:

      • 我看到路径之间不共享,但是每个路径都使用之前的值 p1 = p0 * exp(),然后 p2 = p1 * exp() 等等。因此,我更多地考虑避免在给定时间步长访问每条路径的全局内存(希望这是有道理的)。我最近的 GPU 是 Quadro K4000。我有点理解 DRAM 总线的宽度,但我会读一下,试试cudaMalloc2D() 看看它是怎么回事。关于pb的最后一部分,你能再解释一下吗?在内核之前的 pb 中唯一有用的值是所有路径的第一行。
      • @BRabbit27 :对不起,这是 cudaMallocPitch()。我添加了一个指向 cuda doc 的链接。 talonmies 的回答展示了如何避免从全局内存中读取无用的内容。后面我会给出完整的代码。
      • 可以肯定的是,DRAM 总线宽度 = 内存总线宽度 (deviceQuery) = 192 位。至少这是 Quadro K4000 的,对吧?
      • @BRabbit27,你不需要知道号码。只需使用 cudaMallocPitch(),给它请求的width,它会返回实际的pitch,其中pitch >= width,而pitch 是总线宽度/缓存线的倍数。
      • 谢谢,这仅用于教育目的。我现在正在尝试使用 cudaMallocPitch()。
      【解决方案3】:

      kernelGlobalMem 你正在做3 * c_numTimeSteps 读/写rvpb

      kernelSharedMem 你正在做3 * c_numTimeSteps + c_numTimeSteps / SHMEM_ROWS 读/写rvpb

      kernelSharedMem 更复杂,内存模式看起来很相似。

      kernelGlobalMem 绝对比kernelSharedMem 好。

      【讨论】:

        猜你喜欢
        • 1970-01-01
        • 2016-04-25
        • 2018-06-21
        • 2019-09-26
        • 1970-01-01
        • 1970-01-01
        • 2018-11-30
        • 2014-07-30
        • 1970-01-01
        相关资源
        最近更新 更多