GPU で単純な幾何学的ブラウン運動を実装しています。私のコードはうまく機能します。つまり、正しい値が得られます。私の懸念は、私が得ているスピードアップに関してです。私はもう少し期待していました。これまでのところ、グローバル メモリのみにアクセスする 2 つの実装があり、約 3 倍の速度が向上し、2 つ目は共有メモリを使用して約 2.3 倍の速度が向上します。
私の質問は、Nvidia Visual Profiler でアプリケーションをプロファイリングした後に来ました。それによると、ロード/ストアの効率は 100% ですが、DRAM の使用率は非常に低く (約 10%)、アクセスが結合されていないため、グローバル メモリのリプレイはほぼ 50% です。
共有メモリを使用して常にグローバル メモリ アクセスを回避しようとしていることがわかりましたが、DRAM が低下し (4.5%)、グローバル メモリ リプレイが 46.3% になったことに驚きました。
ブロックごとに使用可能なすべての共有メモリをほとんど使用しているため、カーネル起動時の占有率が低いことに気付きましたが、これが 2 番目のアプローチのパフォーマンスの低下を説明できるかどうかはわかりません。
パフォーマンスに関して何が起こっているのか、そしてそれを改善するためにどこで/何を探すことができるかについて、アドバイスをいただけますか?
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);
}
メイン.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);
}
