cuda可以让不互相依赖的任务进行并行操作，主要任务包括如下：

• 主机端的计算
• 设备端的计算
• 主机到设备端的内存数据传输
• 设备到主机端的内存数据传输
• 设备内的内存之间的数据传输
• 设备到设备之间的数据传输

在主机和设备间可以并行执行的操作有：

• 内核函数执行
• 单个设备内存间的数据拷贝
• 主机到设备的内存拷贝，内存块的大小，小于64KB。
• 以async为前缀的内存拷贝函数
• memory set function calls

并行执行的特性，可以查询设备的属性，查看是否支持。

• cudaDeviceProp.asyncEngineCount > 0    // 支持cpu与gpu间的数据拷贝
• cudaDeviceProp.concurrentKernels > 0     // 支持多个核函数并行执行
• cudaDeviceProp.asyncEngineCount > 2    //支持多个设备间的数据拷贝

必须注意的是在这些拷贝的过程中，主机端的内存必须是非分页内存(page-locked)。cuda使用streams来管理并行操作。

例子：

#include <cstdio>
#include <cstddef>
#include <vector>
#include <cuda_runtime.h>

const int STREAM_COUNT = 4;

__global__ void incKernel(int *in, int *out, int N, int value)
{
    for(int i=blockDim.x *blockIdx.x + threadIdx.x; i<N; i+=blockDim.x * gridDim.x)
    {
        out[i] = in[i] + value;
    }
}

int main()
{
    int deviceNum = 0;
    cudaGetDeviceCount(&deviceNum);

    for(int i=0; i<deviceNum; i++)
    {
        cudaDeviceProp deviceProp;
        cudaGetDeviceProperties(&deviceProp, i);

        if(deviceProp.asyncEngineCount > 0)
        {
            printf("cuda device support memory copy concurrently with kernel exectuion.\n");
        }

        if(deviceProp.concurrentKernels > 0)
        {
            printf("cuda device support mutiple kernel.\n");
        }
    }

    int *h_data_in[STREAM_COUNT];
    int *h_data_out[STREAM_COUNT];
    int *d_data_in[STREAM_COUNT];
    int *d_data_out[STREAM_COUNT];
    cudaStream_t stream[STREAM_COUNT];

    int memsize = 1024 * sizeof(int);

    int data[1024] = {0};

    for(int i=0; i<STREAM_COUNT; i++)
    {
        cudaHostAlloc(&h_data_in[i], memsize, cudaHostAllocDefault);
        memcpy(h_data_in[i], &data[0], memsize);

        cudaHostAlloc(&h_data_out[i], memsize, cudaHostAllocDefault);
        memcpy(h_data_out[i], &data[0], memsize);

        cudaMalloc(&d_data_in[i], memsize);
        cudaMemset(d_data_in[i], 0, memsize);

        cudaMalloc(&d_data_out[i], memsize);
        cudaMemset(d_data_out[i], 0, memsize);

        cudaStreamCreate(&stream[i]);
    }

    int block = 1024 / 32;
    for(int i=0; i<STREAM_COUNT; i++)
    {
        incKernel<<<1, block, 0, stream[i]>>>(d_data_in[i], d_data_out[i], 1024, i*10);
        cudaMemcpyAsync(d_data_in[i], h_data_in[i], memsize, cudaMemcpyHostToDevice, stream[i]);
        cudaMemcpyAsync(h_data_out[i], d_data_out[i], memsize, cudaMemcpyDeviceToHost, stream[i]);
    }

    cudaDeviceSynchronize();

    for(int i=0; i<STREAM_COUNT; i++)
    {
        for(int j=0; j<1024; j++)
        {
            printf("data_out[%d][%d]:%d\n", i, j, h_data_out[i][j]);
        }
    }

    for(int i=0; i<STREAM_COUNT; i++)
    {
        cudaFree(d_data_in[i]);
        cudaFree(d_data_out[i]);
        cudaFreeHost(h_data_in[i]);
        cudaFreeHost(h_data_out[i]);

        cudaStreamDestroy(stream[i]);
    }

    return 0;
}