Add foundational CUDA support and basic kernel

CMakeLists.txt

@@ -597,6 +597,22 @@ install(
             COMPONENT library
             EXCLUDE_FROM_ALL)
 
+# Nothing links against nvrtc-builtins but if we don't include it we get a runtime crash.
+# Unfortunately this means we have to bypass CMake's dependency system and install it manually here.
+if(WIN32)
+    install(
+        FILES "${PROJECT_SOURCE_DIR}/extern/nvidia/_build/target-deps/cuda/cuda/bin/nvrtc-builtins64_118.dll"
+        DESTINATION "${KIT_EXTENSION_BIN_PATH}"
+        COMPONENT install)
+else()
+    install(
+        FILES "${PROJECT_SOURCE_DIR}/extern/nvidia/_build/target-deps/cuda/cuda/lib64/libnvrtc-builtins.so"
+              "${PROJECT_SOURCE_DIR}/extern/nvidia/_build/target-deps/cuda/cuda/lib64/libnvrtc-builtins.so.11.8"
+              "${PROJECT_SOURCE_DIR}/extern/nvidia/_build/target-deps/cuda/cuda/lib64/libnvrtc-builtins.so.11.8.89"
+        DESTINATION "${KIT_EXTENSION_BIN_PATH}"
+        COMPONENT install)
+endif()
+
 install(
     TARGETS CesiumOmniversePythonBindings
     ARCHIVE DESTINATION ${KIT_EXTENSION_BINDINGS_PATH} COMPONENT install

exts/cesium.omniverse/config/extension.toml

@@ -52,6 +52,12 @@ archiveDirs = ["vendor"]
 [[native.plugin]]
 path = "bin/cesium.omniverse.plugin"
 
+[[native.library]]
+"filter:platform"."windows-x86_64"."path" = "bin/${lib_prefix}nvrtc-builtins64_118${lib_ext}"
+
+[[native.library]]
+"filter:platform"."linux-x86_64"."path" = "bin/${lib_prefix}nvrtc-builtins${lib_ext}"
+
 [settings]
 exts."cesium.omniverse".defaultAccessToken = ""
 persistent.exts."cesium.omniverse".userAccessToken = ""

src/core/include/cesium/omniverse/Context.h

@@ -73,6 +73,7 @@ class Context {
 
     int64_t getContextId() const;
     int64_t getNextTilesetId() const;
+    int64_t getNextTileId() const;
 
     const CesiumGeospatial::Cartographic getGeoreferenceOrigin() const;
     void setGeoreferenceOrigin(const CesiumGeospatial::Cartographic& origin);
@@ -151,6 +152,7 @@ class Context {
     int64_t _contextId;
 
     mutable std::atomic<int64_t> _tilesetId{};
+    mutable std::atomic<int64_t> _tileId{};
 
     std::filesystem::path _cesiumExtensionLocation;
     std::filesystem::path _certificatePath;

src/core/include/cesium/omniverse/CudaKernels.h

@@ -0,0 +1,38 @@
+#pragma once
+
+namespace cesium::omniverse::cudaKernels {
+
+inline const char* printPointsKernel = R"(
+
+    extern "C" __global__ void printPoints(float3** points, int numPoints) {
+        const size_t i = blockIdx.x * blockDim.x + threadIdx.x;
+        if (i >= numPoints) return;
+
+        int pointIndex = static_cast<int>(i);
+
+        printf("point %d: %f, %f, %f\n", pointIndex, points[0][pointIndex].x, points[0][pointIndex].y, points[0][pointIndex].z);
+    }
+    )";
+
+inline const char* helloWorldKernel = R"(
+    extern "C" __global__
+    void helloWorld(double* values, size_t count)
+    {
+        size_t i = blockIdx.x * blockDim.x + threadIdx.x;
+        if (count <= i) return;
+
+        printf("Hello world, from index %llu\n", i);
+    }
+    )";
+
+inline const char* createVoxelsKernel = R"(
+    extern "C" __global__
+    void createVoxels(float3** points, size_t count)
+    {
+        size_t i = blockIdx.x * blockDim.x + threadIdx.x;
+        if (count <= i) return;
+
+        // printf("Placeholder: create voxel at index %llu\n", i);
+    }
+    )";
+} // namespace cesium::omniverse::cudaKernels

src/core/include/cesium/omniverse/CudaManager.h

@@ -0,0 +1,97 @@
+#pragma once
+
+#include "cesium/omniverse/CudaKernels.h"
+#include "cesium/omniverse/CudaManager.h"
+
+#include <cuda/include/cuda.h>
+#include <cuda/include/cuda_runtime.h>
+#include <cuda/include/nvrtc.h>
+#include <omni/fabric/FabricUSD.h>
+#include <omni/fabric/IFabric.h>
+
+#include <any>
+#include <stdexcept>
+#include <string>
+#include <unordered_map>
+#include <utility>
+#include <vector>
+
+namespace cesium::omniverse {
+class CudaManager;
+class CudaRunner;
+struct CudaKernel;
+struct CudaKernelArgs;
+
+enum CudaKernelType { HELLO_WORLD, CREATE_VOXELS, PRINT_POINTS };
+enum CudaUpdateType { ONCE, ON_UPDATE_FRAME };
+
+struct CudaKernelArgs {
+    std::unordered_map<std::string, std::any> args;
+};
+
+struct CudaKernel {
+    nvrtcProgram program;
+    char* ptx;
+    CUmodule module;
+    CUfunction function;
+};
+
+class CudaRunner {
+  public:
+    CudaKernelType kernelType;
+
+    CudaRunner() = delete;
+    CudaRunner(
+        CudaKernelType cudaKernelType,
+        CudaUpdateType updateType,
+        int64_t tileId,
+        CudaKernelArgs args,
+        int elementCountArg)
+        : kernelType(cudaKernelType)
+        , kernelArgs(std::move(args))
+        , elementCount(elementCountArg)
+        , _tileId(tileId)
+        , _updateType(updateType){};
+    [[nodiscard]] int64_t getTileId() const {
+        return _tileId;
+    }
+    CudaKernelArgs kernelArgs;
+    [[nodiscard]] const CudaUpdateType& getUpdateType() const {
+        return _updateType;
+    }
+    int elementCount;
+
+  private:
+    // omni::fabric::PrimBucketList _bucketList;
+    int64_t _tileId;
+    CudaUpdateType _updateType;
+};
+
+class CudaManager {
+  public:
+    static CudaManager& getInstance() {
+        static CudaManager instance;
+        return instance;
+    }
+
+    void onUpdateFrame();
+    void addRunner(CudaRunner& cudaRunner);
+    void removeRunner(int64_t tileId);
+    [[nodiscard]] const char* getKernelCode(CudaKernelType kernelType) const;
+    [[nodiscard]] const char* getFunctionName(CudaKernelType kernelType) const;
+
+  private:
+    CUdevice _device;
+    CUcontext _context;
+    bool _initialized = false;
+    std::unordered_map<CudaUpdateType, std::unordered_map<int64_t, CudaRunner>> _runnersByUpdateType;
+    std::unordered_map<CudaKernelType, CudaKernel> _kernels;
+    int _blockSize, _numBlocks;
+
+    void compileKernel(CudaKernelType kernelType);
+    void runAllRunners();
+    void initialize();
+    void runRunner(CudaRunner& runner);
+    void** packArgs(CudaKernelArgs cudaKernelArgs, CudaKernelType cudaKernelType);
+};
+} // namespace cesium::omniverse

src/core/include/cesium/omniverse/FabricGeometry.h

@@ -1,5 +1,6 @@
 #pragma once
 
+#include "cesium/omniverse/CudaManager.h"
 #include "cesium/omniverse/FabricGeometryDefinition.h"
 
 #include <glm/glm.hpp>
@@ -24,13 +25,15 @@ class FabricGeometry {
 
     void setGeometry(
         int64_t tilesetId,
+        int64_t tileId,
         const glm::dmat4& ecefToUsdTransform,
         const glm::dmat4& gltfToEcefTransform,
         const glm::dmat4& nodeTransform,
         const CesiumGltf::Model& model,
         const CesiumGltf::MeshPrimitive& primitive,
         bool smoothNormals,
-        bool hasImagery);
+        bool hasImagery,
+        float geometricError);
 
     void setActive(bool active);
     void setVisibility(bool visible);

src/core/include/cesium/omniverse/FabricMaterial.h

@@ -23,7 +23,7 @@ class FabricMaterial {
         long stageId);
     ~FabricMaterial();
 
-    void setMaterial(int64_t tilesetId, const MaterialInfo& materialInfo);
+    void setMaterial(int64_t tilesetId, int64_t tileId, const MaterialInfo& materialInfo);
     void setBaseColorTexture(const pxr::TfToken& textureAssetPathToken, const TextureInfo& textureInfo);
 
     void clearMaterial();
@@ -50,7 +50,7 @@ class FabricMaterial {
         const omni::fabric::Path& texturePath,
         const pxr::TfToken& textureAssetPathToken,
         const TextureInfo& textureInfo);
-    void setTilesetId(int64_t tilesetId);
+    void setTilesetIdAndTileId(int64_t tilesetId, int64_t tileId);
     bool stageDestroyed();
 
     omni::fabric::Path _materialPath;

src/core/include/cesium/omniverse/FabricPrepareRenderResources.h

@@ -27,6 +27,7 @@ struct FabricMesh {
 struct TileRenderResources {
     glm::dmat4 tileTransform;
     std::vector<FabricMesh> fabricMeshes;
+    int64_t tileId;
 };
 
 class FabricPrepareRenderResources final : public Cesium3DTilesSelection::IPrepareRendererResources {

src/core/include/cesium/omniverse/FabricUtil.h

@@ -22,6 +22,7 @@ struct FabricStatistics {
 
 // -1 means the prim is not yet associated with a tileset
 const auto NO_TILESET_ID = int64_t(-1);
+const auto NO_TILE_ID = int64_t(-1);
 
 } // namespace cesium::omniverse
 
@@ -31,7 +32,7 @@ std::string printFabricStage();
 FabricStatistics getStatistics();
 void destroyPrim(const omni::fabric::Path& path);
 void setTilesetTransform(int64_t tilesetId, const glm::dmat4& ecefToUsdTransform);
-void setTilesetId(const omni::fabric::Path& path, int64_t tilesetId);
+void setTilesetIdAndTileId(const omni::fabric::Path& pathFabric, int64_t tilesetId, int64_t tileId);
 omni::fabric::Path toFabricPath(const pxr::SdfPath& path);
 omni::fabric::Path joinPaths(const omni::fabric::Path& absolutePath, const omni::fabric::Token& relativePath);
 

src/core/include/cesium/omniverse/Tokens.h

@@ -34,6 +34,7 @@ __pragma(warning(push)) __pragma(warning(disable : 4003))
     (vertexColor) \
     (_cesium_localToEcefTransform) \
     (_cesium_tilesetId) \
+    (_cesium_tileId) \
     (_deletedPrims) \
     (_paramColorSpace) \
     (_sdrMetadata) \
@@ -144,6 +145,7 @@ const omni::fabric::Type primvars_vertexColor(omni::fabric::BaseDataType::eFloat
 const omni::fabric::Type Shader(omni::fabric::BaseDataType::eTag, 1, 0, omni::fabric::AttributeRole::ePrimTypeName);
 const omni::fabric::Type subdivisionScheme(omni::fabric::BaseDataType::eToken, 1, 0, omni::fabric::AttributeRole::eNone);
 const omni::fabric::Type _cesium_localToEcefTransform(omni::fabric::BaseDataType::eDouble, 16, 0, omni::fabric::AttributeRole::eMatrix);
+const omni::fabric::Type _cesium_tileId(omni::fabric::BaseDataType::eInt64, 1, 0, omni::fabric::AttributeRole::eNone);
 const omni::fabric::Type _cesium_tilesetId(omni::fabric::BaseDataType::eInt64, 1, 0, omni::fabric::AttributeRole::eNone);
 const omni::fabric::Type _paramColorSpace(omni::fabric::BaseDataType::eToken, 1, 1, omni::fabric::AttributeRole::eNone);
 const omni::fabric::Type _sdrMetadata(omni::fabric::BaseDataType::eToken, 1, 1, omni::fabric::AttributeRole::eNone);

src/core/src/Context.cpp

@@ -3,6 +3,7 @@
 #include "cesium/omniverse/AssetRegistry.h"
 #include "cesium/omniverse/Broadcast.h"
 #include "cesium/omniverse/CesiumIonSession.h"
+#include "cesium/omniverse/CudaManager.h"
 #include "cesium/omniverse/FabricResourceManager.h"
 #include "cesium/omniverse/FabricUtil.h"
 #include "cesium/omniverse/GeospatialUtil.h"
@@ -205,6 +206,8 @@ void Context::onUpdateFrame(const std::vector<Viewport>& viewports) {
     for (const auto& tileset : tilesets) {
         tileset->onUpdateFrame(viewports);
     }
+
+    CudaManager::getInstance().onUpdateFrame();
 }
 
 void Context::processPropertyChanged(const ChangedPrim& changedPrim) {
@@ -420,6 +423,10 @@ int64_t Context::getNextTilesetId() const {
     return _tilesetId++;
 }
 
+int64_t Context::getNextTileId() const {
+    return _tileId++;
+}
+
 const CesiumGeospatial::Cartographic Context::getGeoreferenceOrigin() const {
     const auto georeference = UsdUtil::getOrCreateCesiumGeoreference();