企业官网建设流程全解析

网站上的广告位是怎么做的,哪个公司做网站专业,wordpress 私有文章,页面永久升级一、CUDA中的事件 大家可能在别的开发语言中都学习过事件这个概念#xff0c;其实在CUDA中事件这个概念与它们都类似。不过#xff0c;在CUDA中事件更贴近于其字面本身的意义#xff0c;它是类似一种标志#xff0c;用来密切监视设备进度即同步工具。同时可以通过让应用程序…一、CUDA中的事件大家可能在别的开发语言中都学习过事件这个概念其实在CUDA中事件这个概念与它们都类似。不过在CUDA中事件更贴近于其字面本身的意义它是类似一种标志用来密切监视设备进度即同步工具。同时可以通过让应用程序在程序中的任何点异步记录事件并查询这些事件何时完成来执行准确的计时。当事件之前的所有任务或可选地给定流中的所有命令都已完成时事件即已完成。在所有流中的所有先前任务和命令完成之后流零指空或默认流中的事件也会完成。也就是说在CUDA编程中事件既可以作为同步工具又可以作为精确测量执行时间的工具。二、CUDA事件主要应用场景通过上面的说明大家可以基本明白CUDA事件的定义那么对事件来说其应用场景是什么呢性能测量利用事件既可以进行内核时间的测量还可以进行流线的操作计算重叠等cudaEventRecord(start,0);for(inti0;i2;i){cudaMemcpyAsync(inputDevi*size,inputHosti*size,size,cudaMemcpyHostToDevice,stream[i]);MyKernel100,512,0,stream[i](outputDevi*size,inputDevi*size,size);cudaMemcpyAsync(outputHosti*size,outputDevi*size,size,cudaMemcpyDeviceToHost,stream[i]);}cudaEventRecord(stop,0);cudaEventSynchronize(stop);floatelapsedTime;cudaEventElapsedTime(elapsedTime,start,stop);流的同步主要用于控制流间的依赖关系如多流的同步不同阶段间的计算同步以及多GPU的数据依赖和Graphic另外还可以进行动态工作流的控制处理。__global__voidfoo1(char*A){*A0x1;}__global__voidfoo2(char*B){printf(%d\n,*B);// *B *A 0x1 assuming foo2 waits for foo1// to complete before launching}cudaMemcpyAsync(B,input,size,stream1);// Aliases are allowed at// operation boundariesfoo11,1,0,stream1(A);// allowing foo1 to access A.cudaEventRecord(event,stream1);cudaStreamWaitEvent(stream2,event);foo21,1,0,stream2(B);cudaStreamWaitEvent(stream3,event);cudaMemcpyAsync(output,B,size,stream3);// Both launches of foo2 andcudaMemcpy (which both read)// wait for foo1 (which writes) to complete before proceeding注上面代码来自CUDA官网三、流和流同步虽然事件可以进行流同步但与流同步还是有一些不同之处主要有事件的同步机制更灵活可以多流间控制。而流同步一般是整个流内部事件控制比流同步更精确上文也提到了事件可以进行点的控制而流同步一般是整个流事件与流同步相比开销更低事件应用比流同步要广泛除了同步外还可以进行计时等处理技术概念间互相对比可以更好的加强学习理解的深刻性。四、例程针对上面的说明可以看下面的例程#includecuda_runtime.h#includedevice_launch_parameters.h#includestdio.h#includestdlib.h__global__voidkernelFunc(float*data,intnum,floatfactor){intidxblockIdx.x*blockDim.xthreadIdx.x;if(idxnum){data[idx]data[idx]*factoridx*0.001f;}}intmain(){constintN120;// 100wconstintnumStreams4;constintchunkSizeN/numStreams;constsize_tchunkByteschunkSize*sizeof(float);constsize_ttotalBytesN*sizeof(float);printf(mul stream sync test...\n);float*hDataNULL;cudaMallocHost(hData,totalBytes);for(inti0;iN;i){hData[i](float)rand()/RAND_MAX;}float*dDataNULL;cudaMalloc(dData,totalBytes);cudaStream_tstreams[numStreams];for(inti0;inumStreams;i){cudaStreamCreate(streams[i]);}//1 cacl timecudaEvent_tstartEvent,stopEvent;cudaEvent_tkernelEvents[numStreams];cudaEventCreate(startEvent);cudaEventCreate(stopEvent);for(inti0;inumStreams;i){cudaEventCreateWithFlags(kernelEvents[i],cudaEventDisableTiming);}cudaEventRecord(startEvent,0);intthreadsPerBlock256;intblocksPerChunk(chunkSizethreadsPerBlock-1)/threadsPerBlock;for(inti0;inumStreams;i){intoffseti*chunkSize;cudaMemcpyAsync(dData[offset],hData[offset],chunkBytes,cudaMemcpyHostToDevice,streams[i]);kernelFuncblocksPerChunk,threadsPerBlock,0,streams[i](dData[offset],chunkSize,(float)(i1)*0.5f);cudaGetLastError();cudaEventRecord(kernelEvents[i],streams[i]);cudaMemcpyAsync(hData[offset],dData[offset],chunkBytes,cudaMemcpyDeviceToHost,streams[i]);}for(inti0;inumStreams;i){cudaStreamSynchronize(streams[i]);}// record finish timepointcudaEventRecord(stopEvent,0);cudaEventSynchronize(stopEvent);floattotalTime0.f;cudaEventElapsedTime(totalTime,startEvent,stopEvent);printf(sum time: %.3f ms\n,totalTime);//2 stream dependedprintf(\n stream sync :\n);cudaEvent_tsyncEvent;cudaEventCreate(syncEvent);kernelFuncblocksPerChunk,threadsPerBlock,0,streams[0](dData,chunkSize,2.0f);cudaEventRecord(syncEvent,streams[0]);for(inti1;inumStreams;i){cudaStreamWaitEvent(streams[i],syncEvent,0);kernelFuncblocksPerChunk,threadsPerBlock,0,streams[i](dData[i*chunkSize],chunkSize,1.5f);}for(inti0;inumStreams;i){cudaStreamSynchronize(streams[i]);}printf(sync finish \n);printf(\n event query:\n);cudaEvent_tqueryEvent;cudaEventCreate(queryEvent);kernelFuncblocksPerChunk,threadsPerBlock(dData,chunkSize,1.0f);cudaEventRecord(queryEvent,0);intmaxChecks100;intcheckCount0;while(cudaEventQuery(queryEvent)cudaErrorNotReady){checkCount;if(checkCountmaxChecks){intdummy0;for(intj0;j1000;j){dummyj;}}else{cudaEventSynchronize(queryEvent);break;}}printf(event query count: %d\n,checkCount);cudaEventDestroy(syncEvent);cudaEventDestroy(queryEvent);cudaEventDestroy(startEvent);cudaEventDestroy(stopEvent);for(inti0;inumStreams;i){cudaEventDestroy(kernelEvents[i]);cudaStreamDestroy(streams[i]);}cudaFree(dData);cudaFreeHost(hData);cudaDeviceReset();printf(\n all finish\n);return0;}上面代码的功能主要用于多流间的同步同时对整体的操作时间进行计算。有前面代码的基础应该不难。如果有什么不太清楚的上机运行增加一些打印日志就明白了。五、总结在前面CUDA流的学习中对CUDA事件进行了顺带的说明。本文则展开事件对其具体的内容和应用进行阐述。通过实际的例程让大家可以更清楚的明白事件的运行机制。