七种武器之二：用texture 来代替global memory reading

我们的算法包括多个kernel，而如果用两个global memory来存储input / output的话，实在是很浪费时间。

因为对于一个kernel，只有当计算量很大的时候，才能用computing 来隐藏global memory的latency。如果只是简单地取出一个数据来，做一个运算，再写回去，那么大部分的时间都是浪费在等待memory transfer了。

这时候可以用texture来做为input，当然写数据的时候还得写到global memory里的。

 

//declare texture and linear memory

texture<....> interSrcTex;

texture<....> interDstTex;

 

float *d_interSrcData, *d_interDstData;

 

//init

cudaMallocPitch((void**)&d_interSrcData, &pitch, PPL * sizeof(float), LPF);
cudaMallocPitch((void**)&d_interDstData, &pitch, PPL * sizeof(float), LPF);

cudaBindTexture2D(0, interSrcTex, d_interSrcData, floatTexDesc, PPL, LPF, pitch);
cudaBindTexture2D(0, interDstTex, d_interDstData, floatTexDesc, PPL, LPF, pitch);

 

__global__ void func1(texture<...> inputTex, float* outData)

{

......

}

 

__global__ void func2(texture<....> inputTex, float* outData)

{

......

}

 

//call the different kernels like this way

void main()

{

    func1(interSrcTex, d_interDstData);

    func2(interDstTex, d_interSrcData);

 

}

 

这样至少可以在read data上快一些