CUDA-流

 页锁定内存：

    cudaHostAlloc()分配页锁定内存,页锁定内存也称为固定内存或不可分页内存，它有一个重要的属性：操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中。

流：

    cuda流用于任务的并行。任务并行性是指并行执行两个或多个不同的任务，而不是在大量数据上执行同一个任务的数据并行性。比如处理同一副图，你用一个流处理左边半张图片，再用第二个流处理右边半张图片，这两个流中的代码同时执行，加快了处理速度。

示例：

#include <stdio.h>
#include <cuda_runtime.h>

#define N (1024*1024)
#define DATA_SIZE (N*20)

__global__ void add(int *a,int *b,int *c){
 int idx=threadIdx.x+blockIdx.x*blockDim.x;
 if(idx<N){
  int idx1=(idx+1)%256;
  int idx2=(idx+2)%256;
  float as=(a[idx]+a[idx1]+a[idx2])/3.0f;
  float bs=(b[idx]+b[idx1]+b[idx2])/3.0f;
  c[idx]=(as+bs)/2;
 }
}

int main(){

 cudaDeviceProp prop;
 int whichDevice;
 cudaGetDevice(&whichDevice);
 cudaGetDeviceProperties(&prop,whichDevice);
 if(!prop.deviceOverlap){
   printf("Device not overlap....");
   return 0;
 }

 int *a,*b,*c;
 int *a1,*b1,*c1;
 int *a_host,*b_host,*c_host;
 cudaEvent_t start,end;
 float elapsedTime;
 cudaEventCreate(&start);
 cudaEventCreate(&end);
 cudaEventRecord(start,0);

 cudaStream_t stream0,stream1;
 cudaStreamCreate(&stream0);
cudaStreamCreate(&stream1);

 cudaMalloc((void **)&a,N*sizeof(int));
 cudaMalloc((void **)&b,N*sizeof(int));
 cudaMalloc((void **)&c,N*sizeof(int));

 cudaMalloc((void **)&a1,N*sizeof(int));
 cudaMalloc((void **)&b1,N*sizeof(int));
 cudaMalloc((void **)&c1,N*sizeof(int));


 cudaHostAlloc((void **)&a_host,DATA_SIZE*sizeof(int),cudaHostAllocDefault);
 cudaHostAlloc((void **)&b_host,DATA_SIZE*sizeof(int),cudaHostAllocDefault);
 cudaHostAlloc((void **)&c_host,DATA_SIZE*sizeof(int),cudaHostAllocDefault);

 for(int i=0;i<DATA_SIZE;i++){
  a_host[i]=i;
  b_host[i]=i;
 }

 for(int i=0;i<DATA_SIZE;i+=N*2){
  cudaMemcpyAsync(a,a_host+i,N*sizeof(int),cudaMemcpyHostToDevice,stream0);
  cudaMemcpyAsync(a1,a_host+i+N,N*sizeof(int),cudaMemcpyHostToDevice,stream1);


  cudaMemcpyAsync(b,b_host+i,N*sizeof(int),cudaMemcpyHostToDevice,stream0);
  cudaMemcpyAsync(b1,b_host+i+N,N*sizeof(int),cudaMemcpyHostToDevice,stream1);

  add<<<N/256,256,0,stream0>>>(a,b,c);
  add<<<N/256,256,0,stream1>>>(a1,b1,c1);

  cudaMemcpyAsync(c_host+i,c,N*sizeof(int),cudaMemcpyDeviceToHost,stream0);
  cudaMemcpyAsync(c_host+i+N,c,N*sizeof(int),cudaMemcpyDeviceToHost,stream1);

 }

 cudaStreamSynchronize(stream0);
 cudaStreamSynchronize(stream1);
 cudaEventRecord(end,0);
 cudaEventSynchronize(end);

 cudaEventElapsedTime(&elapsedTime,start,end);
 printf("tie===%3.1f ms\n",elapsedTime);

 cudaFreeHost(a_host);
 cudaFreeHost(b_host);
 cudaFreeHost(c_host);
 cudaFree(a);
 cudaFree(b);
 cudaFree(c);
 cudaFree(a1);
 cudaFree(b1);
 cudaFree(c1);

 cudaStreamDestroy(stream0);
 cudaStreamDestroy(stream1);

return 0;
}

 注：硬件在处理内存复制和核函数执行时分别采用了不同的引擎，因此我们需要知道，将操作放入流中队列中的顺序将影响着CUDA驱动程序调度这些操作以及执行的方式。 因此在将操作放入流的队列时应该采用宽度优先方式，而非深度优先方式。

 

 

参考：《GPU高性能编程CUDA实战》