shader buffer nodes

zeno/include/zeno/extra/ShaderNode.h

@@ -21,15 +21,25 @@ struct ShaderNode : INode {
     ZENO_API ShaderNode();
     ZENO_API ~ShaderNode() override;
 };
-using ShaderDataTypeList = std::tuple<bool, int, unsigned int, float, vec2f, vec3f, vec4f>;
+using ShaderDataTypeList = std::tuple<bool, int32_t, uint32_t, int64_t, uint64_t, float, vec2f, vec3f, vec4f>;
 
-inline const auto ShaderDataTypeNames = std::array { "bool", "int", "uint", "float", "vec2", "vec3", "vec4" };
+inline const auto ShaderDataTypeNames = std::array { "bool", "int", "uint", "int64", "uint64", "float", "vec2", "vec3", "vec4" };
+
+static const inline std::string ShaderDataTypeNamesString = []() {
+    std::string result;
+    for (auto& name : ShaderDataTypeNames) {
+        result += name + std::string(" ");
+    }
+    return result;
+} ();
 
 static const inline std::map<std::string, int> TypeHint {
 
     {"bool", 0},
     {"int", 10},
     {"uint", 11},
+    {"int64", 12},
+    {"uint64", 13},
 
     {"float", 1},
     {"vec2", 2},

zeno/src/extra/ShaderNode.cpp

@@ -8,7 +8,8 @@
 
 namespace zeno {
 
-static std::string ftos(float x) {
+template<typename T>
+static std::string ftos(T x) {
     std::ostringstream ss;
     ss << x;
     return ss.str();
@@ -137,8 +138,42 @@ ZENO_API std::string EmissionPass::collectDefs() const {
     std::string res;
     int cnt = 0;
     for (auto const &var: constants) {
+
+        // auto type = std::visit([&] (auto const &value) -> std::string {
+        //     using T = std::decay_t<decltype(value)>;
+        //     std::string expression {};
+        //     std::string value_string = ftos(value);
+
+        //     zeno::static_for<0, std::tuple_size_v<ShaderDataTypeList>>([&] (auto i) {
+        //         using ThisType = std::tuple_element_t<i, ShaderDataTypeList>;
+
+        //         if (std::is_same_v<ThisType, T>) {
+        //             auto type_string = ShaderDataTypeNames[i];
+        //             //auto type_int = TypeHint.at(type_string);
+        //             expression = std::string(type_string) + "(" + value_string + ")";
+        //             return true;
+        //         }
+        //         return false;
+        //     });
+
+        //     return expression;
+        // }, var.value);
+
         auto expr = std::visit([&] (auto const &value) -> std::string {
             using T = std::decay_t<decltype(value)>;
+
+            if constexpr (std::is_same_v<bool, T>) 
+                return "bool(" + ftos(value) + ")";
+            if constexpr (std::is_same_v<int, T>) 
+                return "int(" + ftos(value) + ")";
+            if constexpr (std::is_same_v<unsigned int, T>) 
+                return "uint(" + ftos(value) + ")";
+
+            if constexpr (std::is_same_v<int64_t, T>) 
+                return "int64_t(" + ftos(value) + ")";
+            if constexpr (std::is_same_v<uint64_t, T>) 
+                return "uint64_t(" + ftos(value) + ")";
+
             if constexpr (std::is_same_v<float, T>) {
                 return typeNameOf(1) + "(" + ftos(value) + ")";
             } else if constexpr (std::is_same_v<vec2f, T>) {

zeno/src/nodes/mtl/MakeMaterial.cpp

@@ -4,7 +4,7 @@
 #include "zeno/types/ListObject.h"
 #include "zeno/types/StringObject.h"
 #include <zeno/types/UserData.h>
-
+#include <tinygltf/json.hpp>
 namespace zeno
 {
   /*struct MakeMaterial
@@ -235,4 +235,69 @@ struct ExtractMaterialShader : zeno::INode
             },
         });
 
+    struct PrimAttrAsShaderBuffer : zeno::INode {
+
+        static std::string aKey() {
+            return "attrNames";
+        }
+        static std::string bKey() { 
+            return "bindNames";
+        }
+
+        virtual void apply() override {
+
+            auto prim = get_input2<zeno::PrimitiveObject>("in");
+
+            auto a_value = get_input2<std::string>(aKey(), "");
+            auto b_value = get_input2<std::string>(bKey(), "");
+
+            auto task = [](const std::string& raw){
+
+                std::string segment;
+                std::stringstream test(raw);
+                std::vector<std::string> result;
+
+                while(std::getline(test, segment, ','))
+                {
+                    result.push_back(segment);
+                }
+                return result;
+            };
+
+            std::vector<std::string> a_list = task(a_value);
+            std::vector<std::string> b_list = task(b_value);
+
+            if (a_list.size() != b_list.size()) {
+                throw std::runtime_error("buffer count doesn't match attr count");
+            }
+
+            nlohmann::json json;
+
+            for (size_t i=0; i<a_list.size(); ++i) {
+                auto a = a_list[i];
+                auto b = b_list[i];
+                json[a] = b;
+            }
+
+            prim->userData().set2("ShaderAttributes", json.dump());
+            set_output("out", std::move(prim));
+        }
+    };
+
+    ZENDEFNODE( PrimAttrAsShaderBuffer,
+    {
+        {
+            {"in"},
+            {"string", PrimAttrAsShaderBuffer::aKey(), ""},
+            {"string", PrimAttrAsShaderBuffer::bKey(), ""}
+        },
+        {
+            {"out"},
+        },
+        {},
+        {
+            "shader",
+        },
+    });
+
 } // namespace zeno

zeno/src/nodes/mtl/ShaderAttrs.cpp

@@ -27,12 +27,12 @@ enum struct RayAttr {
     rayLength, isBackFace, isShadowRay,
 };
 
-static std::string dataTypeDefaultString() {
+static std::string shaderAttrDefaultString() {
     auto name = magic_enum::enum_name(SurfaceAttr::pos);
     return std::string(name);
 }
 
-static std::string dataTypeListString() {
+static std::string shaderAttrListString() {
     auto list0 = magic_enum::enum_names<SurfaceAttr>();
     auto list1 = magic_enum::enum_names<InstAttr>();
     auto list2 = magic_enum::enum_names<RayAttr>();
@@ -72,8 +72,8 @@ struct ShaderInputAttr : ShaderNodeClone<ShaderInputAttr> {
 
 ZENDEFNODE(ShaderInputAttr, {
     {
-        {"enum" + dataTypeListString(), "attr", dataTypeDefaultString()},
-        {"enum float vec2 vec3 vec4 bool int uint", "type", "vec3"},
+        {"enum" + shaderAttrListString(), "attr", shaderAttrDefaultString()},
+        {"enum " + ShaderDataTypeNamesString, "type", "float"},
     },
     {
         {"shader", "out"},

zeno/src/nodes/mtl/ShaderBuffer.cpp

@@ -0,0 +1,66 @@
+#include <zeno/zeno.h>
+#include <zeno/extra/ShaderNode.h>
+#include <zeno/types/ShaderObject.h>
+
+namespace zeno {
+
+struct ShaderBuffer : ShaderNodeClone<ShaderBuffer> {
+    virtual int determineType(EmissionPass *em) override {
+        return TypeHint.at("uint64");
+    }
+
+    virtual void emitCode(EmissionPass *em) override {
+
+        auto name = get_input2<std::string>("name");
+        return em->emitCode("reinterpret_cast<uint64_t>("+ name + "_buffer" +")");
+    }
+};
+
+ZENDEFNODE(ShaderBuffer, {
+                {
+                    {"string", "name", ""},
+                },
+                {
+                    {"shader", "out"},
+                    {"int",   "size"},
+                },
+                {},
+                {"shader"},
+            });
+
+struct ShaderBufferRead : ShaderNodeClone<ShaderBufferRead> {
+
+    virtual int determineType(EmissionPass *em) override {
+
+        em->determineType(get_input("buffer").get());
+
+        auto type = get_input2<std::string>("type");
+        return TypeHint.at(type);
+    }
+
+    virtual void emitCode(EmissionPass *em) override {
+
+        auto buffer = get_input("buffer").get();
+
+        auto in = em->determineExpr(buffer);
+        auto type = get_input2<std::string>("type");
+        auto offset = get_input2<int>("offset");
+
+        em->emitCode("buffer_read<" + type + ">("+ in + "," + std::to_string(offset) + ")" );
+    }
+};
+
+ZENDEFNODE(ShaderBufferRead, {
+                {
+                    {"buffer"},
+                    {"int", "offset", "0"},
+                    {"enum " + ShaderDataTypeNamesString, "type", "float"},
+                },
+                {
+                    {"out"},
+                },
+                {},
+                {"shader"},
+            });
+
+}

zeno/src/nodes/mtl/ShaderTypeCast.cpp

@@ -0,0 +1,51 @@
+#include <zeno/zeno.h>
+#include <zeno/extra/ShaderNode.h>
+#include <zeno/types/ShaderObject.h>
+
+namespace zeno {
+
+static std::string dataTypeDefaultString() {
+    return ShaderDataTypeNames.front();
+}
+
+struct ShaderTypeCast : ShaderNodeClone<ShaderTypeCast> {
+    virtual int determineType(EmissionPass *em) override {
+
+        auto obj = get_input("in").get();
+        em->determineType(obj);
+
+        auto type = get_input2<std::string>("type:");
+        return TypeHint.at(type);
+    }
+
+    virtual void emitCode(EmissionPass *em) override {
+
+        auto op = get_input2<std::string>("op:");
+        auto type = get_input2<std::string>("type:");
+
+        auto obj = get_input("in").get();
+        auto in = em->determineExpr(obj);
+
+        if (op == "bit_cast") {
+            em->emitCode("reinterpret_cast<"+type+"&>("+in+")");
+        } else {
+            em->emitCode(type + "(" + in + ")" );
+        }
+    }
+};
+
+ZENDEFNODE(ShaderTypeCast, {
+    {
+        {"in"}
+    },
+    {
+        {"shader", "out"},
+    },
+    {
+        {"enum bit_cast data_cast ", "op", "bit_cast"},
+        {"enum" + ShaderDataTypeNamesString, "type", "bool"},
+    },
+    {"shader"},
+});
+
+}

zenovis/src/optx/RenderEngineOptx.cpp

@@ -32,6 +32,7 @@
 #include "../../xinxinoptix/OptiXStuff.h"
 #include <zeno/types/PrimitiveTools.h>
 #include <zeno/types/StringObject.h>
+#include <zeno/types/AttrVector.h>
 #include <tinygltf/json.hpp>
 
 #include <map>
@@ -42,6 +43,9 @@
 #include <hair/Hair.h>
 #include <hair/optixHair.h>
 
+#include "ShaderBuffer.h"
+#include <zeno/extra/ShaderNode.h>
+
 namespace zenovis::optx {
 
 struct CppTimer {
@@ -335,6 +339,49 @@ struct GraphicsManager {
                 // ^^^ Don't wuhui, I mean: Literial Synthetic Lazy internal static Local Shared Pointer
                 auto prim_in = prim_in_lslislSp.get();
 
+                if ( prim_in->userData().has("ShaderAttributes") ) {
+                    auto attritbutes  = prim_in->userData().get2<std::string>("ShaderAttributes");
+
+                    using VarType = zeno::AttrVector<zeno::vec3f>::AttrVectorVariant;
+
+                    auto json = nlohmann::json::parse(attritbutes);
+
+                    for (auto& [attrName, bufferName] : json.items()) {
+                        //for (auto& kname : keys) {
+                        auto& val = prim_in->verts.attrs[attrName];
+
+                        std::visit([&, &bufferName=bufferName](auto&& arg) {
+                            using T = std::decay_t<decltype(arg)>;
+
+                            constexpr auto vsize = std::variant_size_v<VarType>;
+
+                            zeno::static_for<0, vsize>([&, &bufferName=bufferName] (auto i) {
+                                using ThisType = std::variant_alternative_t<i, VarType>;
+                                using EleType = typename ThisType::value_type;
+
+                                if constexpr (std::is_same_v<T, ThisType>) {
+
+                                    auto& obj = reinterpret_cast<ThisType&>(val);
+
+                                    if (obj.size() > 0) {
+
+                                        size_t byte_size = obj.size() * sizeof(EleType);
+                                        auto tmp_ptr = std::make_shared<xinxinoptix::raii<CUdeviceptr>>();
+                                        tmp_ptr->resize(byte_size);
+                                        cudaMemcpy((void*)tmp_ptr->handle, obj.data(), byte_size, cudaMemcpyHostToDevice);
+
+                                        load_buffer_group(bufferName, tmp_ptr);
+                                    }
+                                    return true;
+                                }
+                                return false;
+                            });
+
+                        }, val);
+                    }
+                }
+
+
                 if (prim_in->userData().has("curve") && prim_in->verts->size() && prim_in->verts.has_attr("width")) {
 
                     auto& ud = prim_in->userData();

zenovis/xinxinoptix/CMakeLists.txt

@@ -14,7 +14,6 @@ target_sources(zenovis PRIVATE
   Portal.h
   Shape.h
   XAS.h
-
   BCX.h
 
   LightTree.cpp
@@ -26,6 +25,7 @@ target_sources(zenovis PRIVATE
   TypeCaster.cpp
   TypeCaster.h
 
+  ShaderBuffer.h
   GeometryAux.h
 
   SDK/sutil/Aabb.h              SDK/sutil/Quaternion.h

zenovis/xinxinoptix/CallableDefault.cu

@@ -9,7 +9,7 @@
 
 //COMMON_CODE
 
-extern "C" __device__ MatOutput __direct_callable__evalmat(cudaTextureObject_t zenotex[], float4* uniforms, const MatInput& attrs) {
+extern "C" __device__ MatOutput __direct_callable__evalmat(cudaTextureObject_t zenotex[], const float4* uniforms, const void** buffers, const MatInput& attrs) {
 
     /* MODMA */
     auto att_pos = attrs.pos;