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

Question

我不是经验丰富的 CUDA 程序员。我遇到了这样的问题。 我正在尝试将一个大矩阵 (10K*10K) 的瓦片 (32x32) 从全局内存加载到共享内存中，并在它发生时对其进行计时。我意识到如果我将它加载到私有内存（寄存器），它的加载速度比共享内存加载快 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，只有32bit变量和计算，没有具体用途，我只是在玩。

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

#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 代码中检查。 （实际上，我本来希望这些说明也被删除。是否有遗漏的副作用？）

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

从 OP 的 cmets 来看，我认为对共享内存的加载操作的工作方式存在误解。实际上，数据是从全局内存加载到寄存器中，然后存储在共享内存中。 生成的（相关）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 代码专家，如果我错了或遗漏了任何重要的内容，请纠正我。