<divclass="textblock"><p>Modern scientific computing typically leverages GPU-powered parallel processing cores to speed up large-scale applications. This chapters discusses how to implement heterogeneous decomposition algorithms using CPU-GPU collaborative tasking.</p>
<p>Taskflow enables concurrent CPU-GPU tasking by leveraging <ahref="https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__GRAPH.html">CUDA Graph</a>. The tasking interface is referred to as <em>cudaFlow</em>. A cudaFlow is a graph object of type <aclass="el" href="classtf_1_1cudaFlow.html" title="methods for building a CUDA task dependency graph. ">tf::cudaFlow</a> created at runtime similar to dynamic tasking. It manages a task node in a taskflow and associates it with a CUDA Graph. To create a cudaFlow, emplace a callable with an argument of type <aclass="el" href="classtf_1_1cudaFlow.html" title="methods for building a CUDA task dependency graph. ">tf::cudaFlow</a>. The following example implements the canonical saxpy (A·X Plus Y) task graph.</p>
<divclass="fragment"><divclass="line"> 1: #include <taskflow/taskflow.hpp></div><divclass="line"> 2: </div><divclass="line"> 3: <spanclass="comment">// saxpy (single-precision A·X Plus Y) kernel</span></div><divclass="line"> 4: __global__ <spanclass="keywordtype">void</span> saxpy(<spanclass="keywordtype">int</span> n, <spanclass="keywordtype">float</span> a, <spanclass="keywordtype">float</span> *x, <spanclass="keywordtype">float</span> *y) {</div><divclass="line"> 5: <spanclass="keywordtype">int</span> i = blockIdx.x*blockDim.x + threadIdx.x;</div><divclass="line"> 6: <spanclass="keywordflow">if</span> (i < n) {</div><divclass="line"> 7: y[i] = a*x[i] + y[i];</div><divclass="line"> 8: }</div><divclass="line"> 9: }</div><divclass="line">10:</div><divclass="line">11: <spanclass="comment">// main function begins</span></div><divclass="line">12: <spanclass="keywordtype">int</span> main() {</div><divclass="line">13:</div><divclass="line">14: <aclass="code" href="classtf_1_1Taskflow.html">tf::Taskflow</a> taskflow;</div><divclass="line">15: <aclass="code" href="classtf_1_1Executor.html">tf::Executor</a> executor;</div><divclass="line">16: </div><divclass="line">17: <spanclass="keyword">const</span><spanclass="keywordtype">unsigned</span> N = 1<<20; <spanclass="comment">// size of the vector</span></div><divclass="line">18:</div><divclass="line">19: <aclass="codeRef" doxygen="/home/tsung-wei/Code/taskflow/doxygen/cppreference-doxygen-web.tag.xml:http://en.cppreference.com/w/" href="http://en.cppreference.com/w/cpp/container/vector.html">std::vector<float></a> hx(N, 1.0f); <spanclass="comment">// x vector at host</span></div><divclass="line">20: <aclass="codeRef" doxygen="/home/tsung-wei/Code/taskflow/doxygen/cppreference-doxygen-web.tag.xml:http://en.cppreference.com/w/" href="http://en.cppreference.com/w/cpp/container/vector.html">std::vector<float></a> hy(N, 2.0f); <spanclass="comment">// y vector at host</span></div><divclass="line">21:</div><divclass="line">22: <spanclass="keywordtype">float</span> *dx{<spanclass="keyword">nullptr</span>}; <spanclass="comment">// x vector at device</span></div><divclass="line">23: <spanclass="keywordtype">float</span> *dy{<spanclass="keyword">nullptr</span>}; <spanclass="comment">// y vector at device</span></div><divclass="line">24: </div><divclass="line">25: <aclass="code" href="classtf_1_1Task.html">tf::Task</a> allocate_x = taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>(</div><divclass="line">26: [&](){ cudaMalloc(&dx, N*<spanclass="keyword">sizeof</span>(<spanclass="keywordtype">float</span>));}</div><divclass="line">27: );</div><divclass="line">28:</div><divclass="line">29: <aclass="code" href="classtf_1_1Task.html">tf::Task</a> allocate_y = taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>(</div><divclass="line">30: [&](){ cudaMalloc(&dy, N*<spanclass="keyword">sizeof</span>(<spanclass="keywordtype">float</span>));}</div><divclass="line">31: );</div><divclass="line">32:</div><divclass="line">33: <aclass="code" href="classtf_1_1Task.html">tf::Task</a> cudaflow = taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([&](<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf) {</div><divclass="line">34: <spanclass="comment">// create data transfer tasks</span></div><divclass="line">35: <aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> h2d_x = cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(dx, hx.data(), N); <spanclass="comment">// host-to-device x data transfer</span></div><divclass="line">36: <aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> h2d_y = cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(dy, hy.data(), N); <spanclass="comment">// host-to-device y data transfer</span></div><divclass="line">37: <aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> d2h_x = cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(hx.data(), dx, N); <spanclass="comment">// device-to-host x data transfer</span></div><divclass="line">38: <aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> d2h_y = cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(hy.data(), dy, N); <spanclass="comment">// device-to-host y data transfer</span></div><divclass="line">39:</div><divclass="line">40: <spanclass="comment">// launch saxpy<<<(N+255)/256, 256, 0>>>(N, 2.0f, dx, dy)</span></div><divclass="line">41: <aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> kernel = cf.<aclass="code" href="classtf_1_1cudaFlow.html#adb731be71bdd436dfb5e36e6213a9a17">kernel</a>((N+255)/256, 256, 0, saxpy, N, 2.0f, dx, dy);</div><divclass="line">42:</div><divclass="line">43: kernel.<aclass="code" href="classtf_1_1cudaTask.html#a4a9ca1a34bac47e4c9b04eb4fb2f7775">succeed</a>(h2d_x, h2d_y)</div><divclass="line">44: .<aclass="code" href="classtf_1_1cudaTask.html#abdd68287ec4dff4216af34d1db44d1b4">precede</a>(d2h_x, d2h_y);</div><divclass="line">45: });</div><divclass="line">46: cudaflow.<aclass="code" href="classtf_1_1Task.html#a331b1b726555072e7c7d10941257f664">succeed</a>(allocate_x, allocate_y); <spanclass="comment">// overlap data allocations</span></div><divclass="line">47: </div><divclass="line">48: executor.<aclass="code" href="classtf_1_1Executor.html#a81f35d5b0a20ac0646447eb80d97c0aa">run</a>(taskflow).wait();</div><divclass="line">49:</div><divclass="line">50: taskflow.<aclass="code" href="classtf_1_1Taskflow.html#ac433018262e44b12c4cc9f0c4748d758">dump</a>(<aclass="codeRef" doxygen="/home/tsung-wei/Code/taskflow/doxygen/cppreference-doxygen-web.tag.xml:http://en.cppreference.com/w/" href="http://en.cppreference.com/w/cpp/io/basic_ostream.html">std::cout</a>); <spanclass="comment">// dump the taskflow</span></div><divclass="line">51: }</div></div><!-- fragment --><divclass="image">
<li>Lines 3-9 define a saxpy kernel using CUDA </li>
<li>Lines 19-20 declare two host vectors, <code>hx</code> and <code>hy</code></li>
<li>Lines 22-23 declare two device vector pointers, <code>dx</code> and <code>dy</code></li>
<li>Lines 25-31 declare two tasks to allocate memory for <code>dx</code> and <code>dy</code> on device, each of <code>N*sizeof(float)</code> bytes </li>
<li>Lines 33-45 create a cudaFlow to capture kernel work in a graph (two host-to-device data transfer tasks, one saxpy kernel task, and two device-to-host data transfer tasks) </li>
<li>Lines 46-48 define the task dependency between host tasks and the cudaFlow tasks and execute the taskflow</li>
</ul>
<p>Taskflow does not expend unnecessary efforts on kernel programming but focus on tasking CUDA operations with CPU work. We give users full privileges to craft a CUDA kernel that is commensurate with their domain knowledge. Users focus on developing high-performance kernels using a native CUDA toolkit, while leaving difficult task parallelism to Taskflow.</p>
<p>Use <ahref="https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html">nvcc</a> (at least v10) to compile a cudaFlow program:</p>
<divclass="fragment"><divclass="line">~$ nvcc my_cudaflow.cu -I path/to/include/taskflow -O2 -o my_cudaflow</div><divclass="line">~$ ./my_cudaflow</div></div><!-- fragment --><p>Our source autonomously enables cudaFlow when detecting a CUDA compiler.</p>
<p>By default, the executor spawns one worker per GPU. We dedicate a worker set to each heterogeneous domain, for example, host domain and CUDA domain. If your systems has 4 CPU cores and 2 GPUs, the default number of workers spawned by the executor is 4+2, where 4 workers run CPU tasks and 2 workers run GPU tasks (cudaFlow). You can construct an executor with different numbers of GPU workers.</p>
<divclass="fragment"><divclass="line"><aclass="code" href="classtf_1_1Executor.html">tf::Executor</a> executor(17, 8); <spanclass="comment">// 17 CPU workers and 8 GPU workers</span></div></div><!-- fragment --><p>The above executor spawns 17 and 8 workers for running CPU and GPU tasks, respectively. These workers coordinate with each other to balance the load in a work-stealing loop highly optimized for performance.</p>
<p>You can run a cudaFlow on multiple GPUs by explicitly associating a cudaFlow or a kernel task with a CUDA device. A CUDA device is an integer number in the range of <code>[0, N)</code> representing the identifier of a GPU, where <code>N</code> is the number of GPUs in a system. The code below creates a cudaFlow that runs on the GPU device 2 through <code>my_stream</code>.</p>
<divclass="fragment"><divclass="line">taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([](<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf) {</div><divclass="line"> cf.<aclass="code" href="classtf_1_1cudaFlow.html#ad8c0664e4dc3748f043eaa31b69c11cc">device</a>(2);</div><divclass="line"> cf.stream(my_stream); <spanclass="comment">// by default, a cudaFlow runs on a per-worker stream managed by the executor</span></div><divclass="line"><spanclass="comment">// adding more cudaTasks below (all tasks are placed on GPU 2 unless specified explicitly)</span></div><divclass="line">});</div></div><!-- fragment --><p>You can place a kernel on a device explicitly through the method <aclass="el" href="classtf_1_1cudaFlow.html#a4a839dbaa01237a440edfebe8faf4e5b" title="creates a kernel task on a device ">tf::cudaFlow::kernel_on</a> that takes the device identifier in the first argument.</p>
<divclass="fragment"><divclass="line"> 1: #include <taskflow/taskflow.hpp></div><divclass="line"> 2: </div><divclass="line"> 3: <spanclass="comment">// saxpy (single-precision A·X Plus Y) kernel</span></div><divclass="line"> 4: __global__ <spanclass="keywordtype">void</span> saxpy(<spanclass="keywordtype">int</span> n, <spanclass="keywordtype">int</span> a, <spanclass="keywordtype">int</span> *x, <spanclass="keywordtype">int</span> *y, <spanclass="keywordtype">int</span> *z) {</div><divclass="line"> 5: <spanclass="keywordtype">int</span> i = blockIdx.x*blockDim.x + threadIdx.x;</div><divclass="line"> 6: <spanclass="keywordflow">if</span> (i < n) {</div><divclass="line"> 7: z[i] = a*x[i] + y[i];</div><divclass="line"> 8: }</div><divclass="line"> 9: }</div><divclass="line">10:</div><divclass="line">11: <spanclass="keywordtype">int</span> main() {</div><divclass="line">12:</div><divclass="line">13: <spanclass="keyword">const</span><spanclass="keywordtype">unsigned</span> N = 1<<20;</div><divclass="line">14: </div><divclass="line">15: <spanclass="keywordtype">int</span>* dx {<spanclass="keyword">nullptr</span>};</div><divclass="line">16: <spanclass="keywordtype">int</span>* dy {<spanclass="keyword">nullptr</span>};</div><divclass="line">17: <spanclass="keywordtype">int</span>* z1 {<spanclass="keyword">nullptr</span>};</div><divclass="line">18: <spanclass="keywordtype">int</span>* z2 {<spanclass="keyword">nullptr</span>};</div><divclass="line">19: </div><divclass="line">20: cudaMallocManaged(&dx, N*<spanclass="keyword">sizeof</span>(<spanclass="keywordtype">int</span>)); <spanclass="comment">// create a unified memory block for x</span></div><divclass="line">21: cudaMallocManaged(&dy, N*<spanclass="keyword">sizeof</span>(<spanclass="keywordtype">int</span>)); <spanclass="comment">// create a unified memory block for y</span></div><divclass="line">22: cudaMallocManaged(&z1, N*<spanclass="keyword">sizeof</span>(<spanclass="keywordtype">int</span>)); <spanclass="comment">// result of saxpy task 1</span></div><divclass="line">23: cudaMallocManaged(&z2, N*<spanclass="keyword">sizeof</span>(<spanclass="keywordtype">int</span>)); <spanclass="comment">// result of saxpy task 2</span></div><divclass="line">24: </div><divclass="line">25: <spanclass="keywordflow">for</span>(<spanclass="keywordtype">unsigned</span> i=0; i<N; ++i) {</div><divclass="line">26: dx[i] = 1;</div><divclass="line">27: dy[i] = 2;</div><divclass="line">28: }</div><divclass="line">29:</div><divclass="line">30: <aclass="code" href="classtf_1_1Taskflow.html">tf::Taskflow</a> taskflow;</div><divclass="line">31: <aclass="code" href="classtf_1_1Executor.html">tf::Executor</a> executor;</div><divclass="line">32: </div><divclass="line">33: taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([&](<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf){</div><divclass="line">34: <spanclass="comment">// launch the cudaFlow on GPU 0</span></div><divclass="line">35: cf.<aclass="code" href="classtf_1_1cudaFlow.html#ad8c0664e4dc3748f043eaa31b69c11cc">device</a>(0);</div><divclass="line">36:</div><divclass="line">37: <spanclass="comment">// launch the first saxpy kernel on GPU 1</span></div><divclass="line">38: cf.<aclass="code" href="classtf_1_1cudaFlow.html#a4a839dbaa01237a440edfebe8faf4e5b">kernel_on</a>(1, (N+255)/256, 256, 0, saxpy, N, 2, dx, dy, z1);</div><divclass="line">39:</div><divclass="line">40: <spanclass="comment">// launch the second saxpy kernel on GPU 3</span></div><divclass="line">41: cf.<aclass="code" href="classtf_1_1cudaFlow.html#a4a839dbaa01237a440edfebe8faf4e5b">kernel_on</a>(3, (N+255)/256, 256, 0, saxpy, N, 2, dx, dy, z2);</div><divclass="line">42: });</div><divclass="line">43:</div><divclass="line">44: executor.<aclass="code" href="classtf_1_1Executor.html#a81f35d5b0a20ac0646447eb80d97c0aa">run</a>(taskflow).wait();</div><divclass="line">45:</div><divclass="line">46: cudaFree(dx);</div><divclass="line">47: cudaFree(dy);</div><divclass="line">48: </div><divclass="line">49: <spanclass="comment">// verify the solution; max_error should be zero</span></div><divclass="line">50: <spanclass="keywordtype">int</span> max_error = 0;</div><divclass="line">51: <spanclass="keywordflow">for</span> (<spanclass="keywordtype">size_t</span> i = 0; i < N; i++) {</div><divclass="line">52: max_error = <aclass="codeRef" doxygen="/home/tsung-wei/Code/taskflow/doxygen/cppreference-doxygen-web.tag.xml:http://en.cppreference.com/w/" href="http://en.cppreference.com/w/cpp/algorithm/max.html">std::max</a>(max_error, abs(z1[i]-4));</div><divclass="line">53: max_error = <aclass="codeRef" doxygen="/home/tsung-wei/Code/taskflow/doxygen/cppreference-doxygen-web.tag.xml:http://en.cppreference.com/w/" href="http://en.cppreference.com/w/cpp/algorithm/max.html">std::max</a>(max_error, abs(z2[i]-4));</div><divclass="line">54: }</div><divclass="line">55: <aclass="codeRef" doxygen="/home/tsung-wei/Code/taskflow/doxygen/cppreference-doxygen-web.tag.xml:http://en.cppreference.com/w/" href="http://en.cppreference.com/w/cpp/io/basic_ostream.html">std::cout</a> << <spanclass="stringliteral">"saxpy finished with max error: "</span> << max_error << <spanclass="charliteral">'\n'</span>;</div><divclass="line">56: }</div></div><!-- fragment --><p>Debrief:</p>
<ul>
<li>Lines 3-9 define a CUDA saxpy kernel that stores the result to z </li>
<li>Lines 15-23 declare four unified memory blocks accessible from any processor </li>
<li>Lines 25-28 initialize <code>dx</code> and <code>dy</code> blocks by CPU </li>
<li>Lines 33-42 create a cudaFlow task </li>
<li>Lines 34-35 associate the cudaFlow on GPU 0 </li>
<li>Lines 37-38 create a kernel task to launch the first saxpy on GPU 1 and store the result in <code>z1</code></li>
<li>Lines 40-41 create a kernel task to launch the second saxpy on GPU 3 and store the result in <code>z2</code></li>
<li>Lines 44-55 run the taskflow and verify the result (<code>max_error</code> should be zero)</li>
</ul>
<p>Running the program gives the following <ahref="https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html">nvidia-smi</a> snapshot in a system of 4 GPUs:</p>
<divclass="fragment"><divclass="line">+-----------------------------------------------------------------------------+</div><divclass="line">| NVIDIA-SMI 430.50 Driver Version: 430.50 CUDA Version: 10.1 |</div><divclass="line">|-------------------------------+----------------------+----------------------+</div><divclass="line">| GPU Name Persistence-M| Bus-Id Disp.A | Volatile Uncorr. ECC |</div><divclass="line">| Fan Temp Perf Pwr:Usage/Cap| Memory-Usage | GPU-Util Compute M. |</div><divclass="line">|===============================+======================+======================|</div><divclass="line">| 0 GeForce RTX 208... Off | 00000000:18:00.0 Off | N/A |</div><divclass="line">| 32% 35C P2 68W / 250W | 163MiB / 11019MiB | 0% Default |</div><divclass="line">+-------------------------------+----------------------+----------------------+</div><divclass="line">| 1 GeForce RTX 208... Off | 00000000:3B:00.0 Off | N/A |</div><divclass="line">| 33% 43C P2 247W / 250W | 293MiB / 11019MiB | 100% Default |</div><divclass="line">+-------------------------------+----------------------+----------------------+</div><divclass="line">| 2 GeForce RTX 208... Off | 00000000:86:00.0 Off | N/A |</div><divclass="line">| 32% 37C P0 72W / 250W | 10MiB / 11019MiB | 0% Default |</div><divclass="line">+-------------------------------+----------------------+----------------------+</div><divclass="line">| 3 GeForce RTX 208... Off | 00000000:AF:00.0 Off | N/A |</div><divclass="line">| 31% 43C P2 245W / 250W | 293MiB / 11019MiB | 100% Default |</div><divclass="line">+-------------------------------+----------------------+----------------------+</div><divclass="line"></div><divclass="line">+-----------------------------------------------------------------------------+</div><divclass="line">| Processes: GPU Memory |</div><divclass="line">| GPU PID Type Process name Usage |</div><divclass="line">|=============================================================================|</div><divclass="line">| 0 53869 C ./a.out 153MiB |</div><divclass="line">| 1 53869 C ./a.out 155MiB |</div><divclass="line">| 3 53869 C ./a.out 155MiB |</div><divclass="line">+-----------------------------------------------------------------------------+</div></div><!-- fragment --><p>Even if cudaFlow provides interface for device placement, it is your responsibility to ensure correct memory access. For example, you may not allocate a memory block on GPU 2 using <code>cudaMalloc</code> and access it from a kernel on GPU 1. A safe practice is to allocate unified memory blocks using <code>cudaMallocManaged</code> and let the CUDA runtime perform automatic memory migration between processors (as demonstrated in the code example above).</p>
<p>As the same example, you may create two cudaFlows for the two kernels on two GPUs, respectively. The overhead of creating a kernel on the same device as a cudaFlow is much less than the different one.</p>
<p><aclass="el" href="classtf_1_1cudaFlow.html" title="methods for building a CUDA task dependency graph. ">cudaFlow</a> provides a set of methods for users to manipulate device memory data. There are two categories, raw data and typed data. Raw data operations are methods with prefix <code>mem</code>, such as <code>memcpy</code> and <code>memset</code>, that take action on a device memory area in <em>bytes</em>. Typed data operations such as <code>copy</code>, <code>fill</code>, and <code>zero</code>, take <em>logical count</em> of elements. For instance, the following three methods have the same result of zeroing <code>sizeof(int)*count</code> bytes of the device memory area pointed by <code>target</code>.</p>
<divclass="fragment"><divclass="line"><spanclass="keywordtype">int</span>* target;</div><divclass="line">cudaMalloc(&target, count*<spanclass="keyword">sizeof</span>(<spanclass="keywordtype">int</span>));</div><divclass="line"></div><divclass="line">taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([&](<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf){</div><divclass="line"><aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> memset_target = cf.<aclass="code" href="classtf_1_1cudaFlow.html#a079ca65da35301e5aafd45878a19e9d2">memset</a>(target, 0, <spanclass="keyword">sizeof</span>(<spanclass="keywordtype">int</span>) * count);</div><divclass="line"><aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> same_as_above = cf.<aclass="code" href="classtf_1_1cudaFlow.html#aee1fa4aff12a41737ea585fa2e106a35">fill</a>(target, 0, count);</div><divclass="line"><aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> same_as_above_again = cf.<aclass="code" href="classtf_1_1cudaFlow.html#a91c1739bb9a2832f306f3d12693a0994">zero</a>(target, count);</div><divclass="line">});</div></div><!-- fragment --><p>The method <aclass="el" href="classtf_1_1cudaFlow.html#aee1fa4aff12a41737ea585fa2e106a35" title="creates a fill task that fills a typed memory block with a value ">cudaFlow::fill</a> is a more powerful version of <aclass="el" href="classtf_1_1cudaFlow.html#a079ca65da35301e5aafd45878a19e9d2" title="creates a memset task ">cudaFlow::memset</a>. It can fill a memory area with any value of type <code>T</code>, given that <code>sizeof(T)</code> is 1, 2, or 4 bytes. For example, the following code sets each element in the array <code>target</code> to 1234.</p>
<divclass="fragment"><divclass="line">taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([&](<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf){</div><divclass="line"> cf.<aclass="code" href="classtf_1_1cudaFlow.html#aee1fa4aff12a41737ea585fa2e106a35">fill</a>(target, 1234, count);</div><divclass="line">});</div></div><!-- fragment --><p>Similar concept applies to <aclass="el" href="classtf_1_1cudaFlow.html#ad37637606f0643f360e9eda1f9a6e559" title="creates a memcpy task ">cudaFlow::memcpy</a> and <aclass="el" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f" title="creates a copy task ">cudaFlow::copy</a> as well.</p>
<p>You can create a cudaFlow once and launch it multiple times using cudaFlow::repeat or cudaFlow::predicate, given that the graph parameters remain <em>unchanged</em> across all iterations.</p>
<divclass="fragment"><divclass="line">taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([&] (<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf) {</div><divclass="line"><spanclass="comment">// construct the GPU task dependency graph ...</span></div><divclass="line"></div><divclass="line"><spanclass="comment">// launch the cudaFlow 10 times</span></div><divclass="line"> cf.repeat(10);</div><divclass="line"></div><divclass="line"><spanclass="comment">// equivalently</span></div><divclass="line"> cf.predicate([n=10] () <spanclass="keyword">mutable</span> { <spanclass="keywordflow">return</span> n-- == 0; });</div><divclass="line">});</div></div><!-- fragment --><p>The executor iterate the execution of the cudaFlow until the predicate evaluates to <code>true</code>.</p>
<h1><aclass="anchor" id="C6_Granularity"></a>
Granularity</h1>
<p>Creating a cudaFlow has certain overhead, which means fined-grained tasking such as one GPU operation per cudaFlow may not give you any performance gain. You should aggregate as many GPU operations as possible in a cudaFlow to launch the entire graph once instead of separate calls. For example, the following code creates the saxpy task graph at a very fine-grained level using one cudaFlow per GPU operation.</p>
<divclass="fragment"><divclass="line"><aclass="code" href="classtf_1_1Task.html">tf::Task</a> h2d_x = taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([&](<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf) {</div><divclass="line"> cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(dx, hx.data(), N);</div><divclass="line">};</div><divclass="line"></div><divclass="line"><aclass="code" href="classtf_1_1Task.html">tf::Task</a> h2d_y = taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([&](<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf) {</div><divclass="line"> cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(dy, hy.data(), N);</div><divclass="line">};</div><divclass="line"></div><divclass="line"><aclass="code" href="classtf_1_1Task.html">tf::Task</a> d2h_x = taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([&](<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf) {</div><divclass="line"> cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(hx.data(), dx, N);</div><divclass="line">};</div><divclass="line"></div><divclass="line"><aclass="code" href="classtf_1_1Task.html">tf::Task</a> d2h_y = taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([&](<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf) {</div><divclass="line"> cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(hy.data(), dy, N);</div><divclass="line">};</div><divclass="line"></div><divclass="line"><aclass="code" href="classtf_1_1Task.html">tf::Task</a> kernel = taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([&](<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf) {</div><divclass="line"> cf.<aclass="code" href="classtf_1_1cudaFlow.html#adb731be71bdd436dfb5e36e6213a9a17">kernel</a>((N+255)/256, 256, 0, saxpy, N, 2.0f, dx, dy);</div><divclass="line">};</div><divclass="line"></div><divclass="line">kernel.<aclass="code" href="classtf_1_1Task.html#a331b1b726555072e7c7d10941257f664">succeed</a>(h2d_x, h2d_y)</div><divclass="line"> .<aclass="code" href="classtf_1_1Task.html#a8c78c453295a553c1c016e4062da8588">precede</a>(d2h_x, d2h_y);</div></div><!-- fragment --><p>The following code aggregates the five GPU operations using one cudaFlow to deliver much better performance.</p>
<divclass="fragment"><divclass="line"><aclass="code" href="classtf_1_1Task.html">tf::Task</a> cudaflow = taskflow.<aclass="code" href="classtf_1_1FlowBuilder.html#a796e29175380f70246cf2a5639adc437">emplace</a>([&](<aclass="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>& cf) {</div><divclass="line"><aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> h2d_x = cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(dx, hx.data(), N);</div><divclass="line"><aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> h2d_y = cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(dy, hy.data(), N);</div><divclass="line"><aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> d2h_x = cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(hx.data(), dx, N);</div><divclass="line"><aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> d2h_y = cf.<aclass="code" href="classtf_1_1cudaFlow.html#af03e04771b655f9e629eb4c22e19b19f">copy</a>(hy.data(), dy, N);</div><divclass="line"><aclass="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> kernel = cf.<aclass="code" href="classtf_1_1cudaFlow.html#adb731be71bdd436dfb5e36e6213a9a17">kernel</a>((N+255)/256, 256, 0, saxpy, N, 2.0f, dx, dy);</div><divclass="line"> kernel.<aclass="code" href="classtf_1_1cudaTask.html#a4a9ca1a34bac47e4c9b04eb4fb2f7775">succeed</a>(h2d_x, h2d_y)</div><divclass="line"> .<aclass="code" href="classtf_1_1cudaTask.html#abdd68287ec4dff4216af34d1db44d1b4">precede</a>(d2h_x, d2h_y);</div><divclass="line">});</div></div><!-- fragment --><p>We encourage users to study and understand the parallel structure of their applications, in order to come up with the best granularity of task decomposition. A refined task graph can have significant performance difference from the raw counterpart. </p>
</div></div><!-- contents -->
</div><!-- doc-content -->
<!-- start footer part -->
<divid="nav-path" class="navpath"><!-- id is needed for treeview function! -->