From edab05b90ac9ffc24096daaeb48f00f5d32fcff9 Mon Sep 17 00:00:00 2001
From: iaomw <iaomw@live.com>
Date: Mon, 6 Jan 2025 20:04:19 +0800
Subject: [PATCH] improve shader nodes

---
 zeno/src/nodes/mtl/ShaderBuffer.cpp    | 28 +++++++++++++-
 zeno/src/nodes/mtl/ShaderTypeCast.cpp  | 53 +++++++++++++++++++++++++-
 zenovis/xinxinoptix/CallableDefault.cu |  4 +-
 zenovis/xinxinoptix/CallableVolume.cu  |  5 ++-
 4 files changed, 83 insertions(+), 7 deletions(-)

diff --git a/zeno/src/nodes/mtl/ShaderBuffer.cpp b/zeno/src/nodes/mtl/ShaderBuffer.cpp
index 957b5143f..c4aed0e85 100644
--- a/zeno/src/nodes/mtl/ShaderBuffer.cpp
+++ b/zeno/src/nodes/mtl/ShaderBuffer.cpp
@@ -1,10 +1,13 @@
+#include <array>
 #include <zeno/zeno.h>
 #include <zeno/extra/ShaderNode.h>
 #include <zeno/types/ShaderObject.h>
 
 namespace zeno {
+    
+struct ImplShaderBuffer : ShaderNodeClone<ImplShaderBuffer> {
+    int out;
 
-struct ShaderBuffer : ShaderNodeClone<ShaderBuffer> {
     virtual int determineType(EmissionPass *em) override {
         return TypeHint.at("uint64");
     }
@@ -12,7 +15,28 @@ struct ShaderBuffer : ShaderNodeClone<ShaderBuffer> {
     virtual void emitCode(EmissionPass *em) override {
 
         auto name = get_input2<std::string>("name");
-        return em->emitCode("reinterpret_cast<uint64_t>("+ name + "_buffer" +")");
+
+        if ( out > 0 ) {
+            return em->emitCode(name + "_bfsize");
+        }
+        return em->emitCode("reinterpret_cast<uint64_t>("+ name + "_buffer)");
+    }
+};
+
+struct ShaderBuffer : INode {
+
+    virtual void apply() override {
+
+        static const auto list = std::array {"out", "size"}; 
+
+        for(int i=0; i<list.size(); ++i) {
+
+            auto node = std::make_shared<ImplShaderBuffer>();
+            node->inputs["name"] = get_input("name");
+            node->out = i;
+            auto shader = std::make_shared<ShaderObject>(node.get());
+            set_output(list[i], std::move(shader));
+        }
     }
 };
 
diff --git a/zeno/src/nodes/mtl/ShaderTypeCast.cpp b/zeno/src/nodes/mtl/ShaderTypeCast.cpp
index 750598ca6..1ecbeeb09 100644
--- a/zeno/src/nodes/mtl/ShaderTypeCast.cpp
+++ b/zeno/src/nodes/mtl/ShaderTypeCast.cpp
@@ -2,6 +2,10 @@
 #include <zeno/extra/ShaderNode.h>
 #include <zeno/types/ShaderObject.h>
 
+#include <iomanip>
+#include <sstream>
+#include <iostream>
+
 namespace zeno {
 
 static std::string dataTypeDefaultString() {
@@ -43,8 +47,55 @@ ZENDEFNODE(ShaderTypeCast, {
     },
     {
         {"enum bit_cast data_cast ", "op", "bit_cast"},
-        {"enum" + ShaderDataTypeNamesString, "type", "bool"},
+        {"enum " + ShaderDataTypeNamesString, "type", "bool"},
+    },
+    {"shader"},
+});
+
+struct ShaderPrint : ShaderNodeClone<ShaderPrint> {
+    virtual int determineType(EmissionPass *em) override {
+        auto in = get_input("in").get();
+        auto in_type = em->determineType(in);
+        return in_type;
+    }
+
+    virtual void emitCode(EmissionPass *em) override {
+        
+        auto in = get_input("in").get();
+        auto in_type = em->determineType(in);
+        auto in_expr = em->determineExpr(in);
+
+        auto str = get_input2<std::string>("str");
+
+        static const std::map<std::string, std::string> typeTable {
+            {"float", "%f"}, {"bool", "%d"},
+            {"int", "%d"}, {"uint", "%u"},
+            {"int64", "%ll"}, {"uint64", "%llu"}
+        };
+
+        auto typeStr = TypeHintReverse.at(in_type);
+        auto typePri = typeTable.at(typeStr);
+        auto content = str + ": " + typePri + "\\n";
+
+        std::stringstream ss;
+        ss << "printf(";
+        ss << std::quoted(content, '"', '\n');
+        ss << "," << in_expr << ");";
+
+        em->lines.back() += ss.str();
+        em->emitCode(in_expr);
+    }
+};
+
+ZENDEFNODE(ShaderPrint, {
+    {
+        {"in"},
+        {"string", "str", ""}
+    },
+    {
+        {"shader", "out"},
     },
+    {},
     {"shader"},
 });
 
diff --git a/zenovis/xinxinoptix/CallableDefault.cu b/zenovis/xinxinoptix/CallableDefault.cu
index a1eedc2f4..9a15c60cd 100644
--- a/zenovis/xinxinoptix/CallableDefault.cu
+++ b/zenovis/xinxinoptix/CallableDefault.cu
@@ -26,8 +26,8 @@ extern "C" __device__ MatOutput __direct_callable__evalmat(cudaTextureObject_t z
     auto att_instTang = attrs.instTang;
     auto att_rayLength = attrs.rayLength;
 
-    auto att_isBackFace = attrs.isBackFace ? 1.0f:0.0f;
-    auto att_isShadowRay = attrs.isShadowRay ? 1.0f:0.0f;
+    auto att_isBackFace = attrs.isBackFace;
+    auto att_isShadowRay = attrs.isShadowRay;
 
     vec3 b = normalize(cross(attrs.T, attrs.N));
     vec3 t = normalize(cross(attrs.N, b));
diff --git a/zenovis/xinxinoptix/CallableVolume.cu b/zenovis/xinxinoptix/CallableVolume.cu
index b60f51641..a3e71c293 100644
--- a/zenovis/xinxinoptix/CallableVolume.cu
+++ b/zenovis/xinxinoptix/CallableVolume.cu
@@ -242,8 +242,9 @@ extern "C" __device__ void __direct_callable__evalmat(const float4* uniforms, vo
     auto vdb_grids = sbt_data->vdb_grids;
     auto vdb_max_v = sbt_data->vdb_max_v;
 
-    auto att_isShadowRay = attrs.isShadowRay ? 1.0f:0.0f;
-
+    auto att_isBackFace = false;
+    auto att_isShadowRay = attrs.isShadowRay;
+    
 #ifndef _FALLBACK_
 
     //GENERATED_BEGIN_MARK