特别声明: 设备GT540M, 计算能力2.1.代码附在后面;
缓存加载:
(1)Fermi架构,默认情况是启用L1缓存,即采用128字节内存事务。
采用不同的偏移量,以实现非对齐访问。命令行为:“nvprof --metircs gld_efficiency test.exe N” (N为偏移量)。采用批处理,计算0-255的偏移量的全局内存加载效率,统计结果如下:偏移量每隔32,跳变一次。
非缓存加载(L2缓存)
(1)Fermi架构,编译命令:-Xptxas -dlcm=cg 禁用L1缓存,即采用32字节内存事务。偏移量每隔8,跳变一次。
代码如下:
#include"iostream"
#include"cuda_runtime.h"
#include"device_launch_parameters.h"
using namespace std;
__global__ void fun1(float* datain,float* dataout,int n,int offset)
{
int idx=threadIdx.x+blockIdx.x*blockDim.x;
int k=idx+offset;
if(k<n)
datain[idx]=datain[k]+dataout[k];
}
int main(int argc,char* argv[])
{
int offset = atoi(argv[1]);
cout << offset << endl;
const int N=512*15000;
float* h_out,*h_in,*d_in,*d_out;
cudaMallocHost((void**)&h_in,N*sizeof(float));
cudaMallocHost((void**)&h_out,N*sizeof(float));
cudaMalloc((void**)&d_in,N*sizeof(float));
cudaMalloc((void**)&d_out,N*sizeof(float));
for(int i=0;i<N;i++)
{
h_in[i]=i;
}
cudaMemcpy(d_in,h_in,N*sizeof(float),cudaMemcpyHostToDevice);
fun1<<<15000,512>>>(d_in,d_out,N,offset);
cudaMemcpy(d_out,h_out,N*sizeof(float),cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();
return 0;
}