![]() To my knowledge, there is no launch batching in the CUDA driver for any platform other than Windows with the default WDDM driver. I could imagine that higher overhead may be observed for interfaces of reduced width (I think some people use interface sizes down to x1 in crypto-currency mining rigs) or when PCIe gen2 is used (the latter condition shouldn’t be encountered with any system built in the past five years). Note: The 5 usec lower limit to hardware-based launch overhead applies when a PCIe gen3 x16 interface is used. I know that this is not always possible, and that there are real-life use cases that result in extremely short kernel run times (and run times are further shortened with each new generation of GPUs). In general, one should partition work such that kernel run times are in the multi-millisecond range even for the fastest GPU, to avoid undue impact of launch overhead on application-level performance. I don’t know the answer to that but it seems possible. For that reason I recommend CPUs with a base core frequency >= 3.5 GHz, and system memory that uses the fastest speed grade the CPU supports and with as many DDR4 channels as possible.Īn interesting question is whether kernel launch overhead was negatively impacted by the operating system workarounds for recently identified CPU vulnerabilities (Meltdown, Spectre). Note that all overhead beyond a ~5 usec hardware baseline is caused by software, and that this is single-threaded work, meaning the additional software overhead is primarily influenced by single-core CPU performance and therefore core frequency, secondarily by host system memory performance. The observation that use of the CUDA profiler adds about 2 usec per kernel launch seems very plausible given that the profiler needs to insert a hook into the launch mechanism in order to log data about launches. For realistic kernels with arguments, launch overhead should be expected to be around 7 to 8 usec. kernels that do nothing and are not being passed any arguments. If you then run on a host system with a fast CPU with high single-thread performance you should be able to get close to 5 usec when launching null kernels, i.e. If you need / want low launch overhead, either use Windows with a TCC driver ( not possible with a consumer GPU like GTX 745) or use Linux. Batching often causes the overhead of specific launches to fluctuate from close to the lower limit imposed by hardware (around 5 usec) to much higher values (e.g. I am referring to the average because with WDDM the CUDA driver tries to batch launches in order to reduce the average launch overhead. But it is often detrimental to performance as well as other aspects such as GPU memory allocation.Īverage launch overhead of around 25 usec seems perfectly normal in a WDDM scenario. ![]() This has benefits from Microsoft’s perspective, such as increased system stability compared to the previous Windows XP driver model. With WDDM (introduced with Windows 7), and even more so with WDDM 2.x under Windows 10, it is the Windows operating system that controls most aspects of GPU operation. What operating system are you on? Number one reason for high launch overhead is the use of Windows with a WDDM driver. My initial idea is timing execution of an empty kernel. My question is: how to measure the time of launching a new thread block. NVIDIA Developer Forums – 15 Nov 16 Overhead of launching a new thread blockĪs we know, when a thread block terminates, global block scheduler picks and schedules a new thread block to an SM. The GPU and so does not include cudaDeviceSynchronize() etc. LongY reported about half a microsecond (420 tics) but that was internal to I think means about 30% of my elapse time is disappearing in just As expected the time between kernel>Īnd following gpuErrchk( cudaDeviceSynchronize() ) falls but only to about 30% So I tried an experiment: in the existing code I insertĪ conditional return as the first line, which is always true. I have been puzzled by the performance of my GPU code. ![]()
0 Comments
Leave a Reply. |
AuthorWrite something about yourself. No need to be fancy, just an overview. ArchivesCategories |