CUDA-全局内存读取-实验（缓存+非缓存-Fermi架构-sm2.1）

特别声明： 设备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;
}