diff --git a/docs/cudax/stf.rst b/docs/cudax/stf.rst index a2b4607c0d7..e3a66acf963 100644 --- a/docs/cudax/stf.rst +++ b/docs/cudax/stf.rst @@ -92,43 +92,43 @@ CUDASTF is part of the CUDA Experimental library of the CCCL project. It is not Using CUDASTF ^^^^^^^^^^^^^ -CUDASTF is a header-only C++ library which only require to include its -main header. CUDASTF API is part of the ``cuda::experimental::stf`` C++ -namespace, and we will assume for brevity that we are using this -workspace in the rest of this document. +CUDASTF is a header-only C++ library that requires only its main header. +The CUDASTF API is part of the ``cuda::experimental::stf`` C++ namespace; +for brevity, the rest of this document imports that namespace. .. code:: cpp #include - using cuda::experimental::stf; + using namespace cuda::experimental::stf; Compiling ^^^^^^^^^ CUDASTF requires a compiler conforming to the C++17 standard or later. -Although there is no need to link against CUDASTF itself, the library -internally utilizes the CUDA library. +Although there is no CUDASTF library to link, applications use the CUDA Runtime +and Driver APIs. CMake is the recommended integration method. For a manual +source-tree build, add both public CCCL include roots: .. code:: bash - # Compilation flags - nvcc -std=c++17 --expt-relaxed-constexpr --extended-lambda -I$(cudastf_path) - # Linking flags - nvcc -lcuda + export CCCL_ROOT=/path/to/cccl + nvcc -std=c++17 --expt-relaxed-constexpr --extended-lambda \ + -I"${CCCL_ROOT}/cudax/include" -I"${CCCL_ROOT}/libcudacxx/include" \ + example.cu -lcuda -lcudart It is also possible to use CUDASTF without ``nvcc``. This is for example -useful when calling existing CUDA libraries such as CUBLAS which do not +useful when calling existing CUDA libraries such as cuBLAS, which do not require authoring custom kernels. Note that CUDASTF APIs intended to automatically generate CUDA kernels such as ``parallel_for`` or ``launch`` are disabled when compiling without nvcc. .. code:: bash - # Compilation flags - g++ -I$(cudastf_path) - # Linking flags - g++ -lcuda -lcudart + export CCCL_ROOT=/path/to/cccl + export CUDA_HOME=/usr/local/cuda + g++ -std=c++17 -I"${CCCL_ROOT}/cudax/include" -I"${CCCL_ROOT}/libcudacxx/include" \ + -I"${CUDA_HOME}/include" example.cpp -L"${CUDA_HOME}/lib64" -lcuda -lcudart Using CUDASTF within a CMake project ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ @@ -137,7 +137,9 @@ As part of the CCCL project, CUDASTF uses CMake for its build and installation infrastructure, and is the recommended way of building applications that use CUDASTF. -This is facilitated by the CMake Package Manager as illustrated in this simple example which is available `here `_, and which is described in the next paragraph. +This is facilitated by the CMake Package Manager, as illustrated by the +`standalone CUDASTF example `_ +described in the next section. A simple example ^^^^^^^^^^^^^^^^ @@ -146,45 +148,9 @@ The following example illustrates the use of CUDASTF to implement the well-known AXPY kernel, which computes ``Y = Y + alpha * X`` where ``X`` and ``Y`` are two vectors, and ``alpha`` is a scalar_view value. -.. code:: cpp - - #include - - using namespace cuda::experimental::stf; - - template - __global__ void axpy(T a, slice x, slice y) { - int tid = blockIdx.x * blockDim.x + threadIdx.x; - int nthreads = gridDim.x * blockDim.x; - - for (int ind = tid; ind < x.size(); ind += nthreads) { - y(ind) += a * x(ind); - } - } - - int main(int argc, char** argv) { - context ctx; - - const size_t N = 16; - double X[N], Y[N]; - - for (size_t ind = 0; ind < N; ind++) { - X[ind] = sin((double)ind); - Y[ind] = col((double)ind); - } - - auto lX = ctx.logical_data(X); - auto lY = ctx.logical_data(Y); - - double alpha = 3.14; - - /* Compute Y = Y + alpha X */ - ctx.task(lX.read(), lY.rw())->*[&](cudaStream_t s, auto sX, auto sY) { - axpy<<<16, 128, 0, s>>>(alpha, sX, sY); - }; - - ctx.finalize(); - } +.. literalinclude:: ../../cudax/examples/stf/01-axpy.cu + :language: cpp + :caption: AXPY expressed as a CUDASTF task The code is organized into several steps, which will be described in more detail in the following sections: @@ -207,19 +173,22 @@ consumption and system instability. .. code:: bash - mkdir -p build - cd build - cmake .. --preset cudax - cd cudax - ninja cudax.examples.stf -j4 + cmake --preset cudax + cmake --build --preset cudax --target cudax.example.stf.01-axpy -j 4 + +To launch the example from a non-devcontainer build, run: + +.. code:: bash + + ./build/cudax/bin/cudax.example.stf.01-axpy -To launch examples, simply run binaries under the `bin/` -subdirectory in the current directory. For instance, to launch the `01-axpy` -example: +Inside a development container, ``CCCL_BUILD_INFIX`` adds a directory between +``build`` and ``cudax``. The configured example can also be located and run +portably through CTest: .. code:: bash - ./bin/cudax.cpp17.example.stf.01-axpy + ctest --preset cudax -R '^cudax.example.stf.01-axpy$' --output-on-failure Backends and contexts ------------------------------- @@ -1069,8 +1038,8 @@ Box shape There are situations where the desired index space does not correspond to the shape of a logical data object. For those cases, CUDASTF also -provides the template class ``box`` (located in -the header ``cudastf/utility/dimensions.h``) that allows user code to +provides the template class ``box``, available +through the main ```` header, which allows user code to define multidimensional shapes with explicit bounds. The template parameter represents the dimension of the shape. @@ -1929,13 +1898,13 @@ patterns. Generating visualizations of task graphs ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ -Let us consider the ``examples/01-axpy.cu`` example which we compile as -usual with ``make build/examples/01-axpy``. +Consider the ``cudax/examples/stf/01-axpy.cu`` example built in the +getting-started section. .. code:: bash # Run the application with CUDASTF_DOT_FILE set to the filename - CUDASTF_DOT_FILE=axpy.dot build/examples/01-axpy + CUDASTF_DOT_FILE=axpy.dot ./build/cudax/bin/cudax.example.stf.01-axpy # Generate the visualization from this dot file ## PDF format @@ -1988,7 +1957,7 @@ visualization. .. code:: bash - CUDASTF_DOT_FILE=heat.dot build/examples/heat_mgpu 1000 8 4 + CUDASTF_DOT_FILE=heat.dot ./build/cudax/bin/cudax.example.stf.heat_mgpu 1000 8 4 dot -Tpng heat.dot -o heat.png .. image:: stf/images/dot-output-heat.png @@ -1997,9 +1966,9 @@ For advanced users, it is also possible to display internally generated asynchronous operations by setting the ``CUDASTF_DOT_IGNORE_PREREQS`` environment variable to 0. -.. code:: c++ +.. code:: bash - CUDASTF_DOT_IGNORE_PREREQS=0 CUDASTF_DOT_FILE=axpy-with-events.dot build/examples/01-axpy + CUDASTF_DOT_IGNORE_PREREQS=0 CUDASTF_DOT_FILE=axpy-with-events.dot ./build/cudax/bin/cudax.example.stf.01-axpy dot -Tpng axpy-with-events.dot -o axpy-with-events.png .. image:: stf/images/dot-output-axpy-events.png @@ -2107,23 +2076,18 @@ name the generated kernel “updateA” : auto lA = ctx.logical_data(A); ctx.parallel_for(lA.shape(), lA.write()).set_symbol("updateA")->*[] __device__ (size_t i, auto sA) { - A(i) = 2*i + 1; + sA(i) = 2*i + 1; }; -Example with miniWeather -~~~~~~~~~~~~~~~~~~~~~~~~ - -Kernel tuning should always be performed on optimized code : - -.. code:: bash - - make build/examples/miniweather +Example command +~~~~~~~~~~~~~~~ -The following command will analyse the performance of kernels : +Kernel tuning should always be performed on an optimized build. The following +command analyzes a CUDASTF application whose generated kernels have symbols: .. code:: bash - ncu --section=ComputeWorkloadAnalysis --print-nvtx-rename=kernel --nvtx -o output build/examples/miniWeather + ncu --section=ComputeWorkloadAnalysis --print-nvtx-rename=kernel --nvtx -o output ./your_stf_application Note that ``--print-nvtx-rename=kernel --nvtx`` is used to name kernels accordingly to ``NVTX`` traces (which are enabled by the ``set_symbol`` @@ -2148,8 +2112,8 @@ The file generated by ``ncu`` can be opened using ``ncu-ui`` : ncu-ui output.ncu-rep -In this case, we can see that the kernel are named accordingly to the -symbols set in the tasks of the miniWeather examples : |image1| +The generated report displays kernel names derived from the symbols attached to +the corresponding CUDASTF constructs: |image1| .. |image1| image:: stf/images/ncu-ui.png @@ -2161,13 +2125,13 @@ This section gives a brief overview of the CUDASTF API. Using CUDASTF ^^^^^^^^^^^^^ -STF is a C++ header-only library which API is defined in the `cuda::experimental::stf` namespace. +STF is a C++ header-only library whose API is defined in the ``cuda::experimental::stf`` namespace. .. code-block:: cpp - #include + #include - using cuda::experimental::stf; + using namespace cuda::experimental::stf; Creating a Context ^^^^^^^^^^^^^^^^^^ @@ -2200,10 +2164,6 @@ Creating a Logical Data Purpose: Encapsulates data structures (e.g. arrays, slices) to be shared and accessed by tasks. Logical data represents the abstraction of data in the model. -.. code-block:: cpp - - // Create a logical data from an existing piece of data - auto ctx.logical_data(data view [data_place = data_place::current_device()]); Examples: @@ -2218,25 +2178,25 @@ Examples: .. code-block:: cpp - auto data_handle = ctx.logical_data(slice(addr {n})); + auto data_handle = ctx.logical_data(make_slice(addr, n)); - Describing a contiguous matrix of size (m, n) .. code-block:: cpp - auto data_handle = ctx.logical_data(slice(addr, {m, n})); + auto data_handle = ctx.logical_data(make_slice(addr, ::std::tuple{m, n}, m)); - Describing a matrix of size (m, n) with a stride of ld elements .. code-block:: cpp - auto data_handle = ctx.logical_data(slice(addr, {m, n}, {ld})); + auto data_handle = ctx.logical_data(make_slice(addr, ::std::tuple{m, n}, ld)); - Create a logical data from a shape .. code-block:: cpp - auto ctx.logical_data(shape); + auto data_handle = ctx.logical_data(shape); Examples: @@ -2251,7 +2211,7 @@ Tasks Data Dependency ~~~~~~~~~~~~~~~ -Purpose: Define how a logical data should be used in a task construct (and derivated constructs such as `parallel_for`, `launch`, `host_launch`). +Purpose: Define how logical data is used in a task construct and derived constructs such as ``parallel_for``, ``launch``, and ``host_launch``. Syntax: diff --git a/docs/cudax/stf/custom_data_interface.rst b/docs/cudax/stf/custom_data_interface.rst index fb8735094cd..2f549585b71 100644 --- a/docs/cudax/stf/custom_data_interface.rst +++ b/docs/cudax/stf/custom_data_interface.rst @@ -1,330 +1,40 @@ .. _stf_custom_data_interface: -CUDASTF offers an extensible API that allows users to implement their -own data interface. - -Let us for example go through the different steps to implement a data -interface for a very simple simple implementation of a matrix class. - -For the sake of simplicity, we here only consider the CUDA stream -backend, but adding support for the CUDA graph backend simply require -some extra steps which use the CUDA graph API. - -Implementation of the ``matrix`` class -====================================== - -For the sake of simplicity, we consider a very simple representation of -matrix, only defined by the dimensions m and n, and by the base address -of the matrix which we assume to be contiguous. - -.. code:: c++ - - template - class matrix { - public: - matrix(size_t m, size_t n, T* base) : m(m), n(n), base(base) {} - __host__ __device__ T& operator()(size_t i, size_t j) { return base[i + j * m]; } - __host__ __device__ const T& operator()(size_t i, size_t j) const { return base[i + j * m]; } - size_t m, n; - T* base; - }; - -Defining the shape of a matrix -============================== - -The first step consists in defining what is the *shape* of a matrix. The -shape of a matrix should be a class that defines all parameters which -are the same for all data instances, ``m`` and ``n``. On the other hand, -the base address should not be part of this shape class, because each -data instance will have its own base address. - -To define what is the shape of a matrix, we need to specialize the -``cudastf::shape_of`` trait class. - -.. code:: c++ - - template - class cudastf::shape_of> { - public: - /** - * @brief The default constructor. - * - * All `shape_of` specializations must define this constructor. - */ - shape_of() = default; - - explicit shape_of(size_t m, size_t n) : m(m), n(n) {} - - /** - * @name Copies a shape. - * - * All `shape_of` specializations must define this constructor. - */ - shape_of(const shape_of&) = default; - - /** - * @brief Extracts the shape from a matrix - * - * @param M matrix to get the shape from - * - * All `shape_of` specializations must define this constructor. - */ - shape_of(const matrix& M) : shape_of>(M.m, M.n) {} - - /// Mandatory method : defined the total number of elements in the shape - size_t size() const { return m * n; } - - size_t m; - size_t n; - }; - -We here see that ``shape_of>`` contains two ``size_t`` fields -``m`` and ``n``. - -In addition, we need to define a default constructor and a copy -constructors. - -To implement the ``.shape()`` member of the ``logical_data`` class, we -need to define a constructor which takes a const reference to a matrix. - -Finally, if the ``ctx.parallel_for`` construct is needed, we must define -a ``size_t size() const`` method which computes the total number of -elements in a shape. - -Hash of a matrix -================ - -For internal needs, such as using (unordered) maps of data instances, -CUDASTF need to have specialized forms of the ``std::hash`` trait class. - -The ``()`` operator of this class should compute a unique identifier -associated to the description of the data instance. This typically means -computing a hash of the matrix sizes, and of the base address. Note that -this hash *does not* depend on the actual content of the matrix. - -In code snippet, we are using the ``cudastf::hash_combine`` helper which -updates a hash value with another value. This function is available from -the ``cudastf/utility/hash.h`` header. - -.. code:: c++ - - template - struct std::hash> { - std::size_t operator()(matrix const& m) const noexcept { - // Combine hashes from the base address and sizes - return cudastf::hash_all(m.m, m.n, m.base); - } - }; - -Defining a data interface -========================= - -We can now implement the actual data interface for a matrix class, which -defines the basic operations that CUDASTF need to perform on a matrix. - -The ``matrix_stream_interface`` class inherits from the -``data_interface`` class, but to implement a data interface using APIs -based on CUDA streams, ``matrix_stream_interface`` inherits from -``stream_data_interface_simple>`` which contains pure virtual -functions that need to be implemented. - -.. code:: c++ - - template - class matrix_stream_interface : public stream_data_interface_simple> { - public: - using base = stream_data_interface_simple>; - using base::shape_t; - - /// Initialize from an existing matrix - matrix_stream_interface(matrix m) : base(std::move(m)) {} - - /// Initialize from a shape of matrix - matrix_stream_interface(shape_t s) : base(s) {} - - /// Copy the content of an instance to another instance - /// - /// This implementation assumes that we have registered memory if one of the data place is the host - void stream_data_copy(const data_place& dst_memory_node, instance_id_t dst_instance_id, - const data_place& src_memory_node, instance_id_t src_instance_id, cudaStream_t stream) override { - assert(src_memory_node != dst_memory_node); - - cudaMemcpyKind kind = cudaMemcpyDeviceToDevice; - if (src_memory_node == data_place::host) { - kind = cudaMemcpyHostToDevice; - } - - if (dst_memory_node == data_place::host) { - kind = cudaMemcpyDeviceToHost; - } - - const matrix& src_instance = this->instance(src_instance_id); - const matrix& dst_instance = this->instance(dst_instance_id); - - size_t sz = src_instance.m * src_instance.n * sizeof(T); - - cuda_safe_call(cudaMemcpyAsync((void*) dst_instance.base, (void*) src_instance.base, sz, kind, stream)); - } - - /// allocate an instance on a specific data place - /// - /// setting *s to a negative value informs CUDASTF that the allocation - /// failed, and that a memory reclaiming mechanism need to be performed. - void stream_data_allocate(backend_ctx_untyped& ctx, const data_place& memory_node, instance_id_t instance_id, ssize_t& s, - void** extra_args, cudaStream_t stream) override { - matrix& instance = this->instance(instance_id); - size_t sz = instance.m * instance.n * sizeof(T); - - T* base_ptr; - - if (memory_node == data_place::host) { - // Fallback to a synchronous method as there is no asynchronous host allocation API - cuda_safe_call(cudaStreamSynchronize(stream)); - cuda_safe_call(cudaHostAlloc(&base_ptr, sz, cudaHostAllocMapped)); - } else { - cuda_safe_call(cudaMallocAsync(&base_ptr, sz, stream)); - } - - // By filling a positive number, we notify that the allocation was successful - *s = sz; - - instance.base = base_ptr; - } - - /// deallocate an instance - void stream_data_deallocate(backend_ctx_untyped& ctx, const data_place& memory_node, instance_id_t instance_id, void* extra_args, - cudaStream_t stream) override { - matrix& instance = this->instance(instance_id); - if (memory_node == data_place::host) { - // Fallback to a synchronous method as there is no asynchronous host deallocation API - cuda_safe_call(cudaStreamSynchronize(stream)); - cuda_safe_call(cudaFreeHost(instance.base)); - } else { - cuda_safe_call(cudaFreeAsync(instance.base, stream)); - } - } - - /// Register the host memory associated to an instance of matrix - /// - /// Note that this pin_host_memory method is not mandatory, but then it is - /// the responsibility of the user to only passed memory that is already - /// registered, and the allocation method on the host must allocate - /// registered memory too. Otherwise, copy methods need to be synchronous. - bool pin_host_memory(instance_id_t instance_id) override { - matrix& instance = this->instance(instance_id); - if (!instance.base) { - return false; - } - - cuda_safe_call(pin_memory(instance.base, instance.m * instance.n * sizeof(T))); - - return true; - } - - /// Unregister memory pinned by pin_host_memory - void unpin_host_memory(instance_id_t instance_id) override { - matrix& instance = this->instance(instance_id); - unpin_memory(instance.base); - } - }; - -``matrix_stream_interface`` must meet the following requirements so that -they can be used in the CUDA stream backend : - It must provide -constructors which take either a matrix, or a shape of matrix as -arguments. - It must implement the ``stream_data_copy``, -``stream_data_allocate`` and ``stream_data_deallocate`` virtual methods, -which respectively define how to copy an instance into another instance, -how to allocate an instance, and how to deallocate an instance. - It may -implement the ``pin_host_memory`` and ``unpin_host_memory`` virtual -methods which respectively register and unregister the memory associated -to an instance allocated on the host. These two methods are not -mandatory, but it is the responsibility of the user to either only pass -and allocate registered host buffers, or to ensure that the copy method -does not require such memory pinning. Similarly, accessing an instance -located in host memory from a device typically requires to access -registered memory. - -Associating a data interface with the CUDA stream backend -========================================================= - -To ensure that we can initialize a logical data from a matrix, or from -the shape of a matrix with ``stream_ctx::logical_data``, we then need to -specialize the ``cudastf::streamed_interface_of`` trait class. - -The resulting class must simply define a type named ``type`` which is -the type of the data interface for the CUDA stream backend. - -.. code:: c++ - - template - class cudastf::streamed_interface_of> { - public: - using type = matrix_stream_interface; - }; - -Once we have defined this trait class, it is for example possible to -initialize a logical data from a matrix, or from a matrix shape : - -.. code:: c++ - - std::vector v(m * n, 0); - matrix M(m, n, &v[0]); - - // Initialize from a matrix - auto lM = ctx.logical_data(M); - - // Initialize from a shape - auto lM2 = ctx.logical_data(shape_of>(m, n)); - -Example of code using the ``matrix`` data interface -=================================================== - -We can now use the ``matrix`` class in CUDASTF, and access it from -tasks. In this code, we first initialize a matrix on the host, we then -apply a task which will update its content on the current device. We -finally check that the content is correct, by the means of the -write-back mechanism that automatically updates the reference data -instance of a logical data when calling ``ctx.sync()``. - -.. code:: c++ - - template - __global__ void kernel(matrix M) { - int tid_x = blockIdx.x * blockDim.x + threadIdx.x; - int nthreads_x = gridDim.x * blockDim.x; - - int tid_y = blockIdx.y * blockDim.y + threadIdx.y; - int nthreads_y = gridDim.y * blockDim.y; - - for (int x = tid_x; x < M.m; x += nthreads_x) - for (int y = tid_y; y < M.n; y += nthreads_y) { - M(x, y) += -x + 7 * y; - } - } - - int main() { - stream_ctx ctx; - - const size_t m = 8; - const size_t n = 10; - std::vector v(m * n); - - for (size_t j = 0; j < n; j++) - for (size_t i = 0; i < m; i++) { - v[i + j * m] = 17 * i + 23 * j; - } - - matrix M(m, n, &v[0]); - - auto lM = ctx.logical_data(M); - - // M(i,j) += -i + 7*i - ctx.task(lM.rw())->*[](cudaStream_t s, auto dM) { kernel<<>>(dM); }; - - ctx.sync(); - - for (size_t j = 0; j < n; j++) - for (size_t i = 0; i < m; i++) { - assert(v[i + j * m] == (17 * i + 23 * j) + (-i + 7*i)); - } - } +Custom data interfaces +====================== + +CUDASTF can manage user-defined data types through custom data interfaces. A +data interface describes the shape of the data and implements the operations +CUDASTF needs to create, copy, and destroy data instances on different data +places. + +A custom type used with the stream backend typically provides: + +* A specialization of ``cuda::experimental::stf::shape_of`` describing the + logical dimensions of the data. Shapes used by ``parallel_for`` also provide + an ``index_to_coords`` mapping. +* A class derived from ``stream_data_interface_simple`` implementing + ``stream_data_copy``, ``stream_data_allocate``, and + ``stream_data_deallocate``. It may also implement ``pin_host_memory`` and + ``unpin_host_memory``. +* A specialization of + ``cuda::experimental::stf::streamed_interface_of`` associating the user + type with its stream data interface. +* A specialization of ``cuda::experimental::stf::hash`` so instances can be + identified by their layout and storage, independently of their contents. + +The example below implements those customization points for a contiguous +two-dimensional ``matrix``. It creates logical data from a host matrix, +updates it first with a regular CUDASTF task and then with ``parallel_for``, and +calls ``finalize()`` to write the result back to the original host allocation. + +This is the canonical example built by the ``cudax.example.stf.custom_data_interface`` +CMake target, so the documentation and compiled source remain synchronized. + +.. literalinclude:: ../../../cudax/examples/stf/custom_data_interface.cu + :language: cpp + :caption: Complete custom data interface example + +The example implements only the stream backend. Supporting ``graph_ctx`` also +requires a graph data interface that performs the corresponding allocation, +copy, and deallocation operations with CUDA Graph APIs. diff --git a/docs/cudax/stf/lower_level_api.rst b/docs/cudax/stf/lower_level_api.rst index 77508e2879a..1567c596c51 100644 --- a/docs/cudax/stf/lower_level_api.rst +++ b/docs/cudax/stf/lower_level_api.rst @@ -11,10 +11,9 @@ interface for creating tasks, which is described below. .. code:: cpp - #include "cudastf/stf.h" - #include "cudastf/__stf/stream/stream_ctx.h" + #include - using namespace cudastf; + using namespace cuda::experimental::stf; template __global__ void axpy(int n, T a, T* x, T* y) { @@ -26,7 +25,7 @@ interface for creating tasks, which is described below. } } - int main(int argc, char** argv) { + int main() { stream_ctx ctx; const size_t N = 16; @@ -45,12 +44,12 @@ interface for creating tasks, which is described below. /* Compute Y = Y + alpha X */ auto t = ctx.task(lX.read(), lY.rw()); t.start(); - slice sX = t.get<0>(); - slice sY = t.get<1>(); + auto sX = t.template get>(0); + auto sY = t.template get>(1); axpy<<<16, 128, 0, t.get_stream()>>>(sX.size(), alpha, sX.data_handle(), sY.data_handle()); t.end(); - ctx.sync(); + ctx.finalize(); } The ``ctx.task()`` call returns a task object. This object provides @@ -86,22 +85,22 @@ low-level interface. auto lY = ctx.logical_data(Y); for (int k = 0; k < 10; k++) { - graph_task t = ctx.task(); - t.add_deps(handle_X.rw()); + graph_task<> t = ctx.task(); + t.add_deps(lX.rw()); t.start(); cudaGraphNode_t n; cuda_safe_call(cudaGraphAddEmptyNode(&n, t.get_graph(), nullptr, 0)); t.end(); } - graph_task t2 = ctx.task(); + graph_task<> t2 = ctx.task(); t2.add_deps(lX.read(), lY.rw()); t2.start(); cudaGraphNode_t n2; cuda_safe_call(cudaGraphAddEmptyNode(&n2, t2.get_graph(), nullptr, 0)); t2.end(); - ctx.sync(); + ctx.finalize(); A task in the CUDA graph backend corresponds to a *child graph* automatically inserted into the CUDA graph associated to a ``graph_ctx``