.. _ogn_tutorial_cudaData:


Tutorial 8 - GPU Data Node
==========================

The GPU data node creates various attributes for use in a CUDA-based GPU compute. Several representative
types are used, though the list of potential attribute types is not exhaustive. See :ref:`ogn_attribute_types`
for the full list.

This node also introduces the notion of attribute typedefs; a useful concept when passing data around in functions.

OgnTutorialCudaData.ogn
-----------------------
The *ogn* file shows the implementation of a node named "omni.graph.tutorials.CudaData", which has inputs and outputs of
various types to use in various computations. Three different CUDA methods are created to show how each of the
types is passed through to the GPU and used by CUDA.

.. literalinclude:: ../../../../../source/extensions/omni.graph.tutorials/tutorials/tutorial8/OgnTutorialCudaData.ogn
    :linenos:
    :language: json

OgnTutorialCudaData.cpp
-----------------------
The *cpp* file contains the implementation of the compute method, which in turn calls the three CUDA algorithms.

.. literalinclude:: ../../../../../source/extensions/omni.graph.tutorials/tutorials/tutorial8/OgnTutorialCudaData.cpp
    :linenos:
    :language: c++

OgnTutorialCudaData_CUDA.cu
---------------------------
The *cu* file contains the implementation of the algorithms on the GPU using CUDA.

.. literalinclude:: ../../../../../source/extensions/omni.graph.tutorials/tutorials/tutorial8/OgnTutorialCudaData_CUDA.cu
    :linenos:
    :language: c++

GPU Attribute Access
--------------------
Here is the set of generated attributes from the database. The attributes living on the GPU return pointers to
memory as the CPU side cannot dereference it into its actual type (e.g. a `float` value, which would be returned as
a `float&` on the CPU side is returned instead as a `float*` on the GPU side.)

In addition, when calling into CUDA code the data changes type as it crosses the GPU boundary. On the CUDA side it
uses the CUDA native data types when it can, which are bytewise compatible with their CPU counterparts. Note that
in the case of the CPU attribute *multiplier* the data is passed to the CUDA code by value, since it has to be
copied from CPU to GPU.

+---------------------+---------+--------------------+-----------------+
| Database Function   | Is GPU? | CPU Type           | CUDA Type       |
+=====================+=========+====================+=================+
| inputs.a()          | Yes     | const float*       | const float*    |
+---------------------+---------+--------------------+-----------------+
| inputs.b()          | Yes     | const float*       | const float*    |
+---------------------+---------+--------------------+-----------------+
| outputs.sum()       | Yes     | float*             | float*          |
+---------------------+---------+--------------------+-----------------+
| inputs.half()       | Yes     | const pxr::GfHalf* | __half*         |
+---------------------+---------+--------------------+-----------------+
| outputs.half()      | Yes     | pxr::GfHalf*       | __half*         |
+---------------------+---------+--------------------+-----------------+
| inputs.color()      | Yes     | const GfVec3d*     | const double3*  |
+---------------------+---------+--------------------+-----------------+
| outputs.color()     | Yes     | GfVec3d*           | double3*        |
+---------------------+---------+--------------------+-----------------+
| inputs.matrix()     | Yes     | const GfMatrix4d*  | const Matrix4d* |
+---------------------+---------+--------------------+-----------------+
| outputs.matrix()    | Yes     | GfMatrix4d*        | Matrix4d*       |
+---------------------+---------+--------------------+-----------------+
| inputs.multiplier() | No      | const GfVec3f&     | const float3    |
+---------------------+---------+--------------------+-----------------+
| inputs.points()     | Yes     | const GfVec3f*     | const float3**  |
+---------------------+---------+--------------------+-----------------+
| outputs.points()    | Yes     | GfVec3f*           | float3**        |
+---------------------+---------+--------------------+-----------------+

The array attribute *points* does not have an array-like wrapper as the CUDA code would rather deal with
raw pointers. In order to provide the size information, when calling the CUDA code the value ``inputs.points.size()``
should also be passed in.

Notice the subtle difference in types on the CPU side for GPU-based data. Instead of references to data there are
pointers, necessary since the data lives in a different memory-space, and all pointers have an extra level of
indirection for the same reason.

There is also a section of this generated file dedicated to information relevant to the CUDA code. In this section
the CUDA attribute data types are defined.
It is protected with ``#ifdef __CUDACC__`` so that it is only processed when included through the CUDA
compiler (and vice versa, so none of the other setup code will be processed on the CUDA side).

.. code-block:: c++

    #include <cuda_fp16.h>
    #include <omni/graph/core/cuda/CUDAUtils.h>
    #include <omni/graph/core/cuda/Matrix4d.h>
    namespace OgnTutorialCudaDataCudaTypes
    {
    namespace inputs
    {
    using a_t = const float*;
    using b_t = const float*;
    using points_t = const float3**;
    using multiplier_t = const float3*;
    using half_t = const __half*;
    using color_t = const double3*;
    using matrix_t = const Matrix4d*;
    }
    namespace outputs
    {
    using sum_t = float*;
    using points_t = float3**;
    using half_t = __half*;
    using color_t = double3*;
    using matrix_t = Matrix4d*;
    }
    }
    using namespace OgnTutorialCudaDataCudaTypes;

Notice the inclusion of the file *cuda_fp16.h*, needed due to the use of the CUDA type *__half*, and the files
*omni/graph/core/cuda/CUDAUtils.h* and *omni/graph/core/cuda/Matrix4d.h*, which provide support functions for CUDA
math.

The data types used by CUDA are compatible with their equivalents on the C++ side, so you can specify passing arguments
into CUDA from C++ by using a declaration such as this on the C++ side:

.. code-block:: c++

    // In this code "inputs::points_t" is the type "GfVec3f*".
    // The size of that array must be passed in as well since it is not implicit in the data type.
    extern "C" void cudaCompute(
        inputs::points_t*,
        size_t pointSize,
        inputs::multiplier_t*,
        outputs::points_t*
    );

which corresponds to this function defined in the .cu file:

.. code-block:: c++

    // In this code "inputs::points_t" is the type "float3*"
    extern "C" void cudaCompute(
        inputs::points_t* inPoints,
        size_t pointSize,
        inputs::multiplier_t* multiplier,
        outputs::points_t* outPoints
    )
    {...}

Pointers are used in the calls rather than being part of the type definitions in order to emphasize the fact that the
values passed through are pointers to the real data in the fabric, which in this case is data in GPU memory.
