cuda基础----流

1：定义方法

　　创建一个cudaStream_t对象，并在启动内核和进行memcpy时将流对象作为参数传入。

2：运行时API实现

//创建了两个流，并在pinned memory中分配了float型的数组hostPtr
cudaStream_t stream[2];
for(int i = 0;i < 2;i++)
cudaStreamCreate(&stream[i]);

float* hostPtr;
cudaMallocHost((void**)&hostPtr,2*size);

//将每个流定义为由host到device的内存拷贝/kernel启动以及device到host的内存拷贝三个操作组成的序列,允许一个流的内存拷贝与另一个流的ｋｅｒｎｅｌ执行同时进行－－同步执行
for(int i =0 ;i < 2;i++)
cudamemcpyasync(inputdevptr+i*size/sizeof(float),hostptr+i*size/izeof(float),size,cudamemcpyhosttodevice,stream[i]);

for(int i = 0;i < 2;i++)
mykernel<<<100,512,0,stream[i]>>>(outputdevptr+i*size,inputptr+i*size,size);

for(int i = 0;i < 2;i++)
cudamemcpyasync(hostptr+i*size/sizeof(float),outputdevptr+i*size/sizeof(float),size,cudamemcpydevicetohost,strea[i]);

cudathreadsynchronize();//保证进行下一步操作前两个流都已经运行完毕
//cudaStreamSynchronize()用于主机与特定流的同步，同时又允许其他流继续在设备上执行

//流的释放
for(int i = 0;i  < 2;i++)
cudaStreamDestroy(stream[i]);

A：执行参数中没有流参数或使用０作为流参数时，不会创建流。此时为异步执行方式

Ｂ：以下函数会进行轮询，占大量ＣＰＵ资源，最好只在计时或分隔失败的ｋｅｒｎｅｌ启动或内存拷贝时调用这些函数。

　　cudaStreamQuery()-----查询一个流中的位置前的所有操作是否均已完成　　　　　

　　cudaStreamSynchronize()----显示强制ＣＵＤＡ运行时等待流中所有操作完成　　　　

　　cudaThreadSynchronize()---强制ＣＵＤＡ运行时进行等待知道调用位置之前所有流中的所有任务完成才能往下进行

　　cudaStreamSynchroniz0----等待直到所有与主机线程相关的命令执行完。

　　cudaStreamWaitEvent0----

　　cudaStreamDestroy()----一直等待，直到所在流中的任务都结束，才会销毁流，并返回主机线程；

　　cudaStreamAddCallback()---当设备端出现异常时，主机端自动返回

３：stream priorities

cudaStreamCreateWithPriority()-----设置优先级

cudaDeviceGetStreamPriorityRange()---获取优先级[height,low];

4：驱动ＡＰＩ实现

CUstream stream[2];
for(int i = 0;i < 2;i++)
cuStreamCreate(&stream[i],0);

float* hostPtr;
cuMemAllocHost((void**)&hostPtr,2*size);

for(int i = 0;i < 2;i++)
cuMemcpyHostToDAsync(inputDevptr+i*size,hostPtr+i*size,size,stream[i]);

for(int i = 0;i < 2;i++)
{
int offset = 0;
cuParamSeti(cuFunction,offset,outputDevPtr);
offset += sizeof(int);
cuParamSeti(cuFunction,offset,inputDevPtr);
offset += sizeof(int);
cuParamSeti(cuFunction,offset,size);
offset += sizeof(int);
cuParamSetSize(cuFunction,offset);

cuFuncSetBlockShape(cuFunction,512,1,1);
cuLaunchGridAsync(cuFunction,100,1,stream[i]);
}

for(int i = 0;i < 2;i++)
cuMemcpyDtoHAsync(hostPtr+i*size,outputDevPtr+i*size,size,stream[i]);

cuCtxSynchronize();

for(int i = 0;i < 2;i++)
cuStreamDestroy(&stream[i]);