improve shader nodes

zeno/src/nodes/mtl/ShaderBuffer.cpp

@@ -1,18 +1,42 @@
+#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");
     }
 
     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));
+        }
     }
 };
 

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"},
 });
 

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));

zenovis/xinxinoptix/CallableVolume.cu

@@ -226,8 +226,9 @@ static __inline__ __device__ vec2 samplingVDB(const unsigned long long grid_ptr,
     return vec2 { nanoSampling<decltype(_acc), DataTypeNVDB, Order>(_acc, pos_indexed, volin), _grid->tree().root().maximum() };
 }
 
-extern "C" __device__ VolumeOut __direct_callable__evalmat(const float4* uniforms, VolumeIn2& attrs) {
+extern "C" __device__ void __direct_callable__evalmat(const float4* uniforms, void** buffers, void* attrs_ptr, VolumeOut& output) {
 
+    auto& attrs = *reinterpret_cast<VolumeIn2*>(attrs_ptr);
     auto& prd = attrs;
 
     vec3& att_pos = reinterpret_cast<vec3&>(attrs.pos_world);
@@ -241,8 +242,9 @@ extern "C" __device__ VolumeOut __direct_callable__evalmat(const float4* uniform
     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 
@@ -264,8 +266,6 @@ extern "C" __device__ VolumeOut __direct_callable__evalmat(const float4* uniform
     auto extinction = vec3(1.0f);
 #endif // _FALLBACK_
 
-VolumeOut output;
-
 #if _USING_NANOVDB_
 
     output.albedo = clamp(albedo, 0.0f, 1.0f);
@@ -294,5 +294,4 @@ VolumeOut output;
     //USING 3D ARRAY
     //USING 3D Noise 
 #endif
-	return output;
 }

zenovis/xinxinoptix/volume.cu

@@ -123,7 +123,8 @@ extern "C" __global__ void __intersection__volume()
     }
 }
 
-__forceinline__ __device__ auto EvalVolume(uint32_t* seed, float* m16, float sigma_t, float3& pos, bool isShadowRay=false) {
+
+__device__ __inline__ auto EvalVolume(uint32_t& seed, float* m16, float sigma_t, float3& pos, VolumeOut& out, bool isShadowRay=false) {
 
     const HitGroupData* sbt_data = reinterpret_cast<HitGroupData*>( optixGetSbtDataPointer() );
 
@@ -133,14 +134,14 @@ __forceinline__ __device__ auto EvalVolume(uint32_t* seed, float* m16, float sig
 
     vin.isShadowRay = isShadowRay;
 
-    vin.seed = seed;
+    vin.seed = &seed;
     vin.sigma_t = sigma_t;
     vin.sbt_ptr = (void*)sbt_data;
 
     vin.world2object = m16;
 
-
-    return optixDirectCall<VolumeOut, const float4*, const VolumeIn&>( sbt_data->dc_index, sbt_data->uniforms, vin );
+    auto buffers = (void**)params.global_buffers;
+    optixDirectCall<void, const float4*, void**, void*, VolumeOut&>( sbt_data->dc_index, sbt_data->uniforms, buffers, (void*)&vin, out);
 }
 
 extern "C" __global__ void __closesthit__radiance_volume()
@@ -206,7 +207,8 @@ extern "C" __global__ void __closesthit__radiance_volume()
     if (0 == sbt_data->vol_depth) { // Homogeneous
 
         new_orig = ray_orig + 0.5f * (t0 + t1) * ray_dir;
-        VolumeOut homo_out = EvalVolume(&prd->seed, m16, 0.0f, new_orig);
+        VolumeOut homo_out; 
+        EvalVolume(prd->seed, m16, 0.0f, new_orig, homo_out);
         //auto hg = pbrt::HenyeyGreenstein(vol_out.anisotropy);
 
         float3 transmittance = vec3(1.0f);
@@ -298,7 +300,7 @@ extern "C" __global__ void __closesthit__radiance_volume()
     float v_density = 0.0;
     float sigma_t = sbt_data->vol_extinction;
 
-    VolumeOut vol_out;
+    VolumeOut vol_out {};
     auto level = sbt_data->vol_depth;
     auto step_scale = 1.0f / sigma_t;
 
@@ -321,7 +323,7 @@ extern "C" __global__ void __closesthit__radiance_volume()
         } // over shoot, outside of volume
 
         new_orig = ray_orig + (t0+t_ele) * ray_dir;
-        vol_out = EvalVolume(&prd->seed, m16, sigma_t, new_orig);
+        EvalVolume(prd->seed, m16, sigma_t, new_orig, vol_out);
 
         v_density = clamp(vol_out.density / sigma_t, 0.0f, 1.0f);
         emitting += vol_out.emission;
@@ -416,7 +418,8 @@ extern "C" __global__ void __anyhit__occlusion_volume()
 
         test_point += ray_dir * 0.5f * (t0 + t1);
 
-        VolumeOut homo_out = EvalVolume(&prd->seed, m16, sigma_t, test_point);
+        VolumeOut homo_out;
+        EvalVolume(prd->seed, m16, sigma_t, test_point, homo_out);
         hg = pbrt::HenyeyGreenstein(homo_out.anisotropy);
 
         transmittance = expf(-homo_out.extinction * t_max);
@@ -439,7 +442,8 @@ extern "C" __global__ void __anyhit__occlusion_volume()
             break;
         } // over shoot, outside of volume
 
-        VolumeOut vol_out = EvalVolume(&prd->seed, m16, sigma_t, test_point, true);
+        VolumeOut vol_out;
+        EvalVolume(prd->seed, m16, sigma_t, test_point, vol_out, true);
 
         const auto v_density = vol_out.density / sigma_t;
 

zenovis/xinxinoptix/volume.h

@@ -95,53 +95,4 @@ inline float HenyeyGreenstein::sample(const float3 &wo, float3 &wi, const float2
     return PhaseHG(cosTheta, g, gg);
 }
 
-// Spectrum Function Declarations
-__device__ inline float Blackbody(float lambda, float T) {
-    if (T <= 0)
-        return 0;
-    const float c = 299792458.f;
-    const float h = 6.62606957e-34f;
-    const float kb = 1.3806488e-23f;
-    // Return emitted radiance for blackbody at wavelength _lambda_ 
-    float l = lambda * 1e-9f; 
-    float Le = (2 * h * c * c) / (powf(l, 5) * (expf((h * c) / (l * kb * T)) - 1));
-    //CHECK(!IsNaN(Le));
-    return Le;
-}
-
-class BlackbodySpectrum {
-  public:
-    // BlackbodySpectrum Public Methods
-    __device__ 
-    BlackbodySpectrum(float T) : T(T) {
-        // Compute blackbody normalization constant for given temperature
-        float lambdaMax = 2.8977721e-3f / T;
-        normalizationFactor = 1 / Blackbody(lambdaMax * 1e9f, T);
-    }
-
-    float operator()(float lambda) const {
-        return Blackbody(lambda, T) * normalizationFactor;
-    }
-
-    const static int NSpectrumSamples = 3;
-
-    float3 Sample(const float3 &lambda) const {
-        float3 s;
-
-        s.x = Blackbody(lambda.x, T) * normalizationFactor;
-        s.y = Blackbody(lambda.y, T) * normalizationFactor;
-        s.z = Blackbody(lambda.z, T) * normalizationFactor;
-
-        return s;
-    }
-
-    float MaxValue() const { return 1.f; }
-    //std::string ToString() const;
-
-  private:
-    // BlackbodySpectrum Private Members
-    float T;
-    float normalizationFactor;
-};
-
 } // namespace pbrt