c - 在cuda中，加载到共享内存比加载到寄存器慢

Question

我不是经验丰富的 CUDA 程序员。我遇到了这样的问题。 我正在尝试将一个大矩阵 (10K*10K) 的图 block (32x32) 从全局内存加载到共享内存中，并在它发生时对其进行计时。我意识到如果我将它加载到私有(private)内存(寄存器)，它的加载速度比共享内存加载快 4-5 倍。

__global__ void speedtest( float *vel,int nx) {

int globalx = blockDim.x * blockIdx.x + threadIdx.x+pad;
int globalz = blockDim.y * blockIdx.y + threadIdx.y+pad;
int localx=threadIdx.x;
int localz=threadIdx.y;

float ptest;
__shared__ float stest[tile][tile];

//stest[localz][localx]=vel[globalz*nx+globalx]; //load to shared memory
ptest=vel[globalz*nx+globalx];  //load to private memory
__syncthreads();
}

我一一注释掉stest和ptest，用cudaeventrecord计算运行时间。 stest 加载时间为 3.2 毫秒，ptest 加载时间为 0.75 毫秒。我究竟做错了什么？时间应该非常相似吧？我错过了什么？

配置:Cuda 7.5，gtx 980，只有32位的变量和计算，没有特定用途，我只是玩玩。

我正在按要求发布示例代码

#include<stdio.h>
#include <math.h>
#define tile 32
#include <helper_cuda.h>
void makeittwo(float *array,int nz,int nx)
{
//this just assigns a number into the vector
int n2;
n2=nx*nz;
for (int i=0;i<n2;i++)
array[i]=2000;
}
__global__ void speedtest( float *vel,int nx,int nz) {

int globalx = blockDim.x * blockIdx.x + threadIdx.x;
int globalz = blockDim.y * blockIdx.y + threadIdx.y;
int localx=threadIdx.x;
int localz=threadIdx.y;

float ptest; //declarations
__shared__ float stest[tile][tile];

if (globalx<nx && globalz<nz){
stest[localz][localx]=vel[globalz*nx+globalx]; //shared variable
//ptest=vel[globalz*nx+globalx];                        //private variable

//comment out ptest and stest one by one to test them  
}
__syncthreads();

}       

int main(int argc,char *argv)
{
int nx,nz,N;
float *vel;

nz=10000;nx=10000; //matrix dimensions
N=nz*nx; //convert matrix into vector

checkCudaErrors(cudaMallocHost(&vel,sizeof(float)*N)); //using pinned memory
makeittwo(vel,nz,nx);

dim3 dimBlock(tile,tile);
dim3 dimGrid;

int blockx=dimBlock.x;
int blockz=dimBlock.y;

dimGrid.x = (nx + blockx - 1) / (blockx);
dimGrid.y = (nz + blockz - 1) / (blockz);

float *d_vel;
checkCudaErrors(cudaMalloc(&d_vel,sizeof(float)*(N))); //copying to device
checkCudaErrors(cudaMemcpy(d_vel, vel, sizeof(float)*(N), cudaMemcpyHostToDevice));

cudaEvent_t start,stop;
float elapsedTime;

cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);
speedtest<<<dimGrid,dimBlock>>>(d_vel,nx,nz); //calling the function
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime,start,stop);

printf("time=%3.3f ms\n",elapsedTime);
checkCudaErrors(cudaMemcpy(vel, d_vel, sizeof(float)*N, cudaMemcpyDeviceToHost)); 
//calling the matrix back  to check if all went well (this fails if out of bound calls are made)

cudaDeviceReset();

}

Answer 1

示例代码实际上并没有测量 OP 期望测量的内容，因为一些指令被编译器优化掉了。

在局部变量示例(ptest)中，加载不会影响内核外部的状态。在这种情况下，编译器可以自由地完全删除指令。这可以在 SASS 代码中看到。当 ptest=vel[globalz*nx+globalx]; 处于事件状态或两个语句(ptest 和 stest)都被删除时，SASS 代码是相同的。要检查 SASS 代码，您可以在目标文件上运行 cuobjdump --dump-sass。

显然，共享内存 示例中的指令并未优化，可以在 SASS 代码中检查。 (事实上​​ ，我希望指令也被删除。是否有遗漏的副作用？)

正如评论中已经讨论的那样，通过简单的计算 (ptest*=ptest) 和 写入全局内存，编译器无法删除指令，因为它改变了全局状态。

根据 OP 的评论，我认为对共享内存的加载操作的工作方式存在误解。实际上，数据是从全局内存加载到寄存器，然后存储在共享内存中。 生成的(相关)SASS 指令(针对 sm_30)如下所示

LD.E R2, [R6]; // load to register R2
STS [R0], R2; // store from register R2 to shared memory

以下乘法并存储到全局内存示例演示了另一种情况，其中编译器不会生成人们可能天真的期望的代码:

stest[localz][localx]=vel[globalz*nx+globalx]; // load to shared memory
stest[localz][localx]*=stest[localz][localx]; // multiply
vel[globalz*nx+globalx]=stest[localz][localx]; // save to global memory

SASS 代码显示变量仅在计算后存储在共享内存中(永远不会从共享内存中读取)。

LD.E R2, [R6]; // load to register
FMUL R0, R2, R2; // multiply
STS [R3], R0; // store the result in shared memory
ST.E [R6], R0; // store the result in global memory

我不是 SASS 代码的真正专家，如果我错了或遗漏了任何重要内容，请纠正我。