CUDA: оценка скорости чтения / записи в кэш L2 с помощью nvprof - PullRequest
Новая
       69

CUDA: оценка скорости чтения / записи в кэш L2 с помощью nvprof

0 голосов
/ 13 апреля 2020

Я пытаюсь оценить частоту обращений к кэшу L2 на основе результатов профилировщика (nvprof) и столкнулся с несколькими проблемами, которые мне не понятны.

1. Определение частоты обращений к кэш-памяти второго уровня (в архитектуре Pascal)

В архитектуре Pascal выглядит, что показатели частоты обращений к кэш-памяти второго уровня получают следующим образом:

l2_tex_read_hit_rate = sum(l2_subp*_read_tex_hit_sectors)/sum(l2_subp*_read_tex_sector_queries)

l2_tex_write_hit_rate = sum(l2_subp*_write_tex_hit_sectors)/sum(l2_subp*_write_tex_sector_queries)

Это было подтверждено проверкой следующего результата nvprof: 

// events                           Total
l2_subp0_read_tex_sector_queries    929
l2_subp1_read_tex_sector_queries    927
l2_subp0_write_tex_sector_queries   927
l2_subp1_write_tex_sector_queries   929
l2_subp0_read_tex_hit_sectors       129
l2_subp1_read_tex_hit_sectors       127
l2_subp0_write_tex_hit_sectors      425
l2_subp1_write_tex_hit_sectors      419

// metrics                          Avg
l2_tex_read_hit_rate                13.79% = (129+127)/(929+927)
l2_tex_write_hit_rate               45.47% = (425+419)/(927+929)
l2_tex_hit_rate                     29.63% = (129+127+425+419)/(929+927+927+929)

Таким образом, приведенное выше определение частоты посещений учитывает только отправленные запросы. из кеша текстур.

Поэтому мой вопрос таков: как общая частота попаданий в кэш L2 отличается от приведенного выше определения? Кеш текстур - это единственные компоненты, которые влияют на частоту попаданий в кэш L2?

2. Интерпретация событий пропущенных l2 и запросов (в архитектуре Pascal)

Я предположил, что они должны отличаться, поскольку кэш текстур - не единственный кэш самого низкого уровня. Чтобы проверить это, я попытался получить частоту попаданий в кэш L2 в более общем виде. Я думал, что смогу получить рейтинг попаданий, получив показатель промахов (потому что hit_rate + miss_rate = 1):

L2 read miss rate = sum(l2_subp*_read_sector_misses)/sum(l2_subp*_total_read_sector_queries)
L2 write miss rate = sum(l2_subp*_write_sector_misses)/sum(l2_subp*_total_write_sector_queries)

Однако число записей промахов было больше, чем запросов записи, где частота пропусков из моих определений будет больше 100%. 

// events                             Total
l2_subp0_read_sector_misses           800
l2_subp1_read_sector_misses           806
l2_subp0_write_sector_misses          1577
l2_subp1_write_sector_misses          1632

l2_subp0_total_read_sector_queries    977
l2_subp1_total_read_sector_queries    967
l2_subp0_total_write_sector_queries   932
l2_subp1_total_write_sector_queries   937

Так что этот подход может быть неправильным способом оценки общей частоты попаданий / промахов в кеше L2 , Как мне поступить с такими событиями (например, l2_subp*_write_sector_misses и l2_subp*_total_write_sector_queries), чтобы получить правильное?

3. Как оценить частоту попаданий WRITE в кэш L2 для архитектуры Kepler

Профилировщик для архитектуры Kepler предоставляет только l2_l1_read_hit_rate.

Почему он не дает l2_l1_write_hit_rate или l2_subp*_write_tex_hit_sectors? Есть ли способ получить этот показатель c, комбинируя события (как я делал в разделе 2)?

4. Краткое изложение вопроса

  1. (в архитектуре Pascal). Как общая частота попаданий в кэш L2 отличается от частоты, определенной для текстурного кэша? Или имеет смысл обсуждать общие вопросы?

  2. Если это имеет смысл, как я могу получить общую частоту обращений в кэш L2?

  3. ( В архитектуре Kepler) почему профилировщик не обеспечивает скорость попадания в кэш L2 для записи ? и Как я могу определить его по результатам профилирования?

Приложение

Ниже приведены коды, которые я использовал для сравнительного анализа

1. CudaTest.cu 

#include "Acts/Utilities/Platforms/CUDA/CudaVector.cu"
#include <Eigen/Dense>
#include <boost/test/unit_test.hpp>
#include <cuda_profiler_api.h>

template<typename AFloat, int row, int col>
__global__ void MatrixLoadStore(const Eigen::Matrix<AFloat,row,col>* input,
                                Eigen::Matrix<AFloat,row,col>* output){

  for (int i=0; i<col; i++){
    output[blockIdx.x](threadIdx.x,i) = input[blockIdx.x](threadIdx.x,i);
  }
}

namespace Acts{
namespace Test{

BOOST_AUTO_TEST_SUITE(Utilities)
BOOST_AUTO_TEST_CASE( CUDAOBJ_TEST ){

  const int vecDim = 100;  // Vector Dimension                                                           
  const int nVec   = 128; // Number of vectors                                                           

  dim3 gridSize(1,1,1);
  dim3 blockSize(vecDim,1,1);
  int  bufSize;

  bufSize   = gridSize.x * blockSize.x;
  Eigen::Matrix<float, vecDim, nVec>  iMat_cpu[bufSize];
  for (int i=0; i< bufSize; i++){
    iMat_cpu[i] = Eigen::Matrix<float,vecDim,nVec>::Random();
  }

  cudaProfilerStart();

  CudaVector<Eigen::Matrix<float,vecDim,nVec>> iMat_cuda(bufSize, iMat_cpu, bufSize, 0);
  CudaVector<Eigen::Matrix<float,vecDim,nVec>> oMat_cuda(bufSize);
  MatrixLoadStore<float, vecDim, nVec><<< gridSize, blockSize >>>(iMat_cuda.Get(),oMat_cuda.Get());
  Eigen::Matrix<float,vecDim,nVec>* oMat_cpu = oMat_cuda.GetHost();

  cudaProfilerStop();

  BOOST_REQUIRE( iMat_cpu[0] == oMat_cpu[0] );
}
BOOST_AUTO_TEST_SUITE_END()

}
}

2. CudaVector.cu 

pragma once

#include <iostream>
#include <memory>
#include "cuda.h"
#include "cuda_runtime.h"
#include "CudaUtils.cu"

namespace Acts{

template<typename Var_t>
class CudaVector{

public:

  CudaVector(size_t size){
    fSize = size;
    cudaErrChk( cudaMalloc((Var_t**)&fDevPtr, fSize*sizeof(Var_t)) );
  }

  CudaVector(size_t size, Var_t* vector){
    fSize = size;
    cudaErrChk( cudaMalloc((Var_t**)&fDevPtr, fSize*sizeof(Var_t)) );
    CopyH2D(vector, fSize, 0);
  }

  CudaVector(size_t size, Var_t* vector, size_t len, size_t offset){
    fSize = size;
    cudaErrChk( cudaMalloc((Var_t**)&fDevPtr, fSize*sizeof(Var_t)) );
    CopyH2D(vector, len, offset);
  }

  ~CudaVector(){
    cudaFree(fDevPtr);
  }

  size_t GetSize(){return fSize;}

  Var_t* Get(size_t offset=0) { return fDevPtr+offset; }

  Var_t* GetHost() {
    Var_t* fHostPtr = new Var_t[fSize];
    cudaErrChk( cudaMemcpy(fHostPtr, fDevPtr, fSize*sizeof(Var_t), cudaMemcpyDeviceToHost) );
    return fHostPtr;
  }

  void CopyH2D(Var_t* vector, size_t len, size_t offset){
    cudaErrChk( cudaMemcpy(fDevPtr+offset, vector, len*sizeof(Var_t), cudaMemcpyHostToDevice) );
  }
  void CopyH2D(Var_t* vector, size_t len, size_t offset, cudaStream_t* stream){
    cudaErrChk( cudaMemcpyAsync(fDevPtr+offset, vector, len*sizeof(Var_t), cudaMemcpyHostToDevice, *stream) );
  }

private:
  Var_t* fDevPtr;
  size_t fSize;
};
}

3. CudaUtils.cu 

#pragma once

#include <cuda.h>
#include <cuda_runtime.h>
#include <iostream>

#define cudaErrChk(ans) { cudaAssert((ans), __FILE__, __LINE__); }
inline void cudaAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
  if (code != cudaSuccess)
    {
      fprintf(stderr,"CUDAassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
    }
}
Добро пожаловать на сайт PullRequest, где вы можете задавать вопросы и получать ответы от других членов сообщества.

Похожие темы

...