diff --git a/ui/zenoedit/launch/serialize.cpp b/ui/zenoedit/launch/serialize.cpp
index 27a8877b5c..5493f8e91c 100644
--- a/ui/zenoedit/launch/serialize.cpp
+++ b/ui/zenoedit/launch/serialize.cpp
@@ -13,7 +13,7 @@ using namespace JsonHelper;
 
 QSet<QString> lightCameraNodes({
     "CameraEval", "CameraNode", "CihouMayaCameraFov", "ExtractCameraData", "GetAlembicCamera","MakeCamera",
-    "LightNode", "BindLight", "ProceduralSky", "HDRSky",
+    "LightNode", "BindLight", "ProceduralSky", "HDRSky", "SkyComposer"
     });
 
 std::set<std::string> matNodeNames = {"ShaderFinalize", "ShaderVolume", "ShaderVolumeHomogeneous"};
diff --git a/zeno/include/zeno/types/LightObject.h b/zeno/include/zeno/types/LightObject.h
index c6c7f7321e..374ceece56 100644
--- a/zeno/include/zeno/types/LightObject.h
+++ b/zeno/include/zeno/types/LightObject.h
@@ -5,6 +5,12 @@
 #include <zeno/core/IObject.h>
 #include <zeno/utils/vec.h>
 
+#else
+
+#ifndef vec3f
+#define vec3f vec3
+#endif
+
 #endif
 
 namespace zeno {
@@ -23,6 +29,15 @@ namespace zeno {
         LightConfigDoubleside = 2u
     };
 
+    struct DistantLightData {
+        vec3f direction;
+        float angle;
+        vec3f color;
+        float intensity;
+
+        DistantLightData() = default;
+    };
+
 #ifndef __CUDACC_RTC__ 
 
 struct LightData {
diff --git a/zeno/include/zeno/utils/pfm.h b/zeno/include/zeno/utils/pfm.h
new file mode 100644
index 0000000000..7efe11a952
--- /dev/null
+++ b/zeno/include/zeno/utils/pfm.h
@@ -0,0 +1,27 @@
+#pragma once
+
+#include <vector>
+#include <string>
+#include <cstring>
+
+#include <zeno/utils/format.h>
+#include <zeno/utils/fileio.h>
+
+namespace zeno {
+
+static void write_pfm(const char* path, int w, int h, const float *pixel, bool mono=false) {
+    std::string header = zeno::format("PF\n{} {}\n-1.0\n", w, h);
+    char channel = 3;
+
+    if (mono) {
+        header = zeno::format("Pf\n{} {}\n-1.0\n", w, h);
+        channel = 1;
+    }
+
+    std::vector<char> data(header.size() + w * h * sizeof(float) * channel);
+    memcpy(data.data(), header.data(), header.size());
+    memcpy(data.data() + header.size(), pixel, w * h * sizeof(float) * channel);
+    zeno::file_put_binary(data, path);
+}
+
+}
\ No newline at end of file
diff --git a/zeno/src/extra/GlobalComm.cpp b/zeno/src/extra/GlobalComm.cpp
index ec3b8b6eb6..b09c34f9be 100644
--- a/zeno/src/extra/GlobalComm.cpp
+++ b/zeno/src/extra/GlobalComm.cpp
@@ -21,7 +21,7 @@ namespace zeno {
 std::vector<std::filesystem::path> cachepath(3);
 std::unordered_set<std::string> lightCameraNodes({
     "CameraEval", "CameraNode", "CihouMayaCameraFov", "ExtractCameraData", "GetAlembicCamera","MakeCamera",
-    "LightNode", "BindLight", "ProceduralSky", "HDRSky",
+    "LightNode", "BindLight", "ProceduralSky", "HDRSky", "SkyComposer"
     });
 std::set<std::string> matNodeNames = {"ShaderFinalize", "ShaderVolume", "ShaderVolumeHomogeneous"};
 
diff --git a/zeno/src/nodes/ProcedrualSkyNode.cpp b/zeno/src/nodes/ProcedrualSkyNode.cpp
index e6ad587fd4..c96a0927a9 100644
--- a/zeno/src/nodes/ProcedrualSkyNode.cpp
+++ b/zeno/src/nodes/ProcedrualSkyNode.cpp
@@ -2,6 +2,12 @@
 #include <zeno/types/PrimitiveObject.h>
 #include <zeno/types/UserData.h>
 #include <zeno/utils/fileio.h>
+#include <zeno/ListObject.h>
+#include <tinygltf/json.hpp>
+
+#include <zeno/utils/eulerangle.h>
+#include <zeno/types/LightObject.h>
+
 
 namespace zeno {
 struct ProceduralSky : INode {
@@ -78,6 +84,234 @@ ZENDEFNODE(HDRSky, {
     {"shader"},
 });
 
+struct DistantLightWrapper : IObject{
+    DistantLightData data;
+};
+
+struct DistantLight : INode {
+
+    virtual void apply() override {
+        auto dir2 = get_input2<zeno::vec2f>("Lat-Lon");
+        // dir2[0] = fmod(dir2[0], 180.f);
+        // dir2[1] = fmod(dir2[1], 180.f);
+        dir2 *= M_PIf / 180.f;
+
+        zeno::vec3f dir3;
+        dir3[1] = std::sin(dir2[0]);
+        
+        dir3[2] = std::cos(dir2[0]) * std::cos(dir2[1]);
+        dir3[0] = std::cos(dir2[0]) * std::sin(dir2[1]);
+        
+        //dir3 = zeno::normalize(dir3);
+    
+        auto angleExtent = get_input2<float>("angleExtent");
+        angleExtent = zeno::clamp(angleExtent, 0.0f, 60.0f);
+
+        auto color = get_input2<zeno::vec3f>("color");
+        auto intensity = get_input2<float>("intensity");
+        intensity = fmaxf(0.0, intensity);
+
+        auto result = std::make_shared<DistantLightWrapper>();
+        result->data.direction = dir3;
+        result->data.angle = angleExtent;
+        result->data.color = color;
+        result->data.intensity = intensity;
+        set_output2("out", std::move(result) );
+    }
+};
+
+ZENDEFNODE(DistantLight, {
+    {
+        {"vec2f", "Lat-Lon", "45, 90"},
+        {"float", "angleExtent", "0.5"},
+        {"colorvec3f", "color", "1,1,1"},
+        {"float", "intensity", "1"}
+    },
+    {
+        {"out"},
+    },
+    {
+    },
+    {"shader"},
+});
+
+struct PortalLight : INode {
+    virtual void apply() override {
+
+        auto pos = get_input2<zeno::vec3f>("pos");
+        auto scale = get_input2<zeno::vec2f>("scale");
+        auto rotate = get_input2<zeno::vec3f>("rotate");
+        auto size = get_input2<int>("size");
+        size = std::max(size, 180);
+
+        scale = 0.5f * abs(scale);
+
+        auto order = get_input2<std::string>("EulerRotationOrder:");
+        auto orderTyped = magic_enum::enum_cast<EulerAngle::RotationOrder>(order).value_or(EulerAngle::RotationOrder::XYZ);
+
+        auto measure = get_input2<std::string>("EulerAngleMeasure:");
+        auto measureTyped = magic_enum::enum_cast<EulerAngle::Measure>(measure).value_or(EulerAngle::Measure::Radians);
+
+        glm::vec3 eularAngleXYZ = glm::vec3(rotate[0], rotate[1], rotate[2]);
+        glm::mat4 rotation = EulerAngle::rotate(orderTyped, measureTyped, eularAngleXYZ);
+
+        glm::mat4 transform(1.0f);
+
+        transform = glm::translate(transform, glm::vec3(pos[0], pos[1], pos[2]));
+        transform = transform * rotation;
+        transform = glm::scale(transform, glm::vec3(scale[0], 0.5 * (scale[0] + scale[1]), scale[1]));
+
+        auto prim = std::make_shared<zeno::PrimitiveObject>();
+        prim->verts->resize(8);
+
+        prim->verts[0] = zeno::vec3f(-1, 0, -1);
+        prim->verts[1] = zeno::vec3f(+1, 0, -1);
+        prim->verts[2] = zeno::vec3f(+1, 0, +1);
+        prim->verts[3] = zeno::vec3f(-1, 0, +1);
+
+        prim->verts[4] = zeno::vec3f(0, 0, 0);
+        prim->verts[5] = zeno::vec3f(0.5, 0, 0);
+        prim->verts[6] = zeno::vec3f(0, 0.5, 0);
+        prim->verts[7] = zeno::vec3f(0, 0, 0.5);
+
+        for (size_t i=0; i<prim->verts->size(); ++i) {
+            auto& ele = prim->verts[i];
+            auto ttt = transform * glm::vec4(ele[0], ele[1], ele[2], 1.0f);
+            prim->verts[i] = zeno::vec3f(ttt.x, ttt.y, ttt.z);
+        }
+
+        //prim->lines.attrs.clear();
+        prim->lines->resize(8);
+        prim->lines[0] = {0, 1};
+        prim->lines[1] = {1, 2};
+        prim->lines[2] = {2, 3};
+        prim->lines[3] = {3, 0};
+
+        prim->lines[4] = {4, 5};
+        prim->lines[5] = {4, 6};
+        prim->lines[6] = {4, 7};
+
+        auto& color = prim->verts.add_attr<zeno::vec3f>("clr");
+        color.resize(8);
+        color[0] = {1,1,1};
+        color[1] = {1,1,1};
+        color[2] = {1,1,1};
+        color[3] = {1,1,1};
+        
+        color[4] = {1, 1, 1};
+        color[5] = {1, 0, 0};
+        color[6] = {0, 1, 0};
+        color[7] = {0, 0, 1};
+        //prim->lines.update();
+        prim->userData().set2("size", size);
+        set_output2("out", std::move(prim));
+    }
+};
+
+ZENDEFNODE(PortalLight, {
+    {
+        {"vec3f", "pos", "0,0,0"},
+        {"vec2f", "scale", "1, 1"},
+        {"vec3f", "rotate", "0,0,0"},
+        {"int", "size", "180"}
+    },
+    {
+        {"out"},
+    },
+    {
+        {"enum " + EulerAngle::RotationOrderListString(), "EulerRotationOrder", "XYZ"},
+        {"enum " + EulerAngle::MeasureListString(), "EulerAngleMeasure", "Degree"}
+    },
+    {"shader"},
+});
+
+struct SkyComposer : INode {
+    virtual void apply() override {
+
+        auto prim = std::make_shared<zeno::PrimitiveObject>();
+
+        if (has_input("dlights")) {
+            auto dlights = get_input<ListObject>("dlights")->get<DistantLightWrapper>();
+            if (dlights.empty()) {
+                throw zeno::makeError("Bad input for dlights");
+            }
+
+            prim->verts->resize(dlights.size());
+            auto& attr_rad = prim->verts.add_attr<float>("rad");
+            auto& attr_angle = prim->verts.add_attr<float>("angle");
+            auto& attr_color = prim->verts.add_attr<zeno::vec3f>("color");
+            auto& attr_inten = prim->verts.add_attr<float>("inten");
+
+            unsigned i = 0;
+            for (const auto& dlight : dlights) {
+                
+                prim->verts[i] = dlight->data.direction;
+                attr_rad[i] = 0.0f;
+                attr_angle[i] = dlight->data.angle;
+                attr_color[i] = dlight->data.color;
+                attr_inten[i] = dlight->data.intensity;
+
+                ++i;
+            }
+        }
+
+        if (has_input("portals")) {
+            auto portals = get_input<ListObject>("portals")->get<zeno::PrimitiveObject>();
+            if (portals.empty()) {
+                throw zeno::makeError("Bad input for portals");
+            }
+
+            using json = nlohmann::json;
+            std::vector<zeno::vec3f> raw(4 * portals.size());
+            std::vector<int> psizes(portals.size());
+
+            for (size_t i=0; i<portals.size(); ++i) {
+                auto &rect = portals[i];
+
+                auto p0 = rect->verts[0];
+                auto p1 = rect->verts[1];
+                auto p2 = rect->verts[2];
+                auto p3 = rect->verts[3];
+
+                /* p0 --- p1 */
+                /* --------- */
+                /* p3 --- p2 */
+
+                raw[4 * i + 0] = p0;
+                raw[4 * i + 1] = p1;
+                raw[4 * i + 2] = p2;
+                raw[4 * i + 3] = p3;
+
+                auto psize = rect->userData().get2<int>("size");
+                psizes[i] = psize;
+            }
+
+            json aux(raw);
+            prim->userData().set2("portals", std::move(aux.dump()));
+            prim->userData().set2("psizes", json(psizes).dump());
+        }
+
+        prim->userData().set2("SkyComposer", std::move(1));
+        prim->userData().set2("isRealTimeObject", std::move(1));
+        set_output2("out", std::move(prim));
+    }
+};
+
+ZENDEFNODE(SkyComposer, {
+    {
+
+        {"list", "dlights"},
+        {"list", "portals"}
+    },
+    {
+        {"out"},
+    },
+    {
+        {"enum SphereUnbounded", "proxy", "SphereUnbounded"},
+    },
+    {"shader"},
+});
+
 vec3f colorTemperatureToRGB(float temperatureInKelvins)
 {
     vec3f retColor;
diff --git a/zenovis/CMakeLists.txt b/zenovis/CMakeLists.txt
index 7389441083..f7fd3cd746 100644
--- a/zenovis/CMakeLists.txt
+++ b/zenovis/CMakeLists.txt
@@ -2,8 +2,6 @@ file(GLOB_RECURSE source CONFIGURE_DEPENDS include/*.h src/*.cpp)
 file(GLOB_RECURSE glad_source CONFIGURE_DEPENDS glad/include/*.h glad/src/*.c)
 file(GLOB_RECURSE stbi_source CONFIGURE_DEPENDS stbi/include/*.h stbi/src/*.cpp stbi/src/*.c)
 
-#OPTION(OPTIX_USE_20XX "turn on if on a 20xx gpu" OFF)
-
 add_library(zenovis OBJECT ${source} ${glad_source} ${stbi_source})
 target_link_libraries(zenovis PRIVATE ${CMAKE_DL_LIBS})
 target_link_libraries(zenovis PUBLIC zeno)
diff --git a/zenovis/include/zenovis/RenderEngine.h b/zenovis/include/zenovis/RenderEngine.h
index a5c0ad58b9..4b275d16b4 100644
--- a/zenovis/include/zenovis/RenderEngine.h
+++ b/zenovis/include/zenovis/RenderEngine.h
@@ -15,7 +15,7 @@ struct Scene;
 struct RenderEngine {
     virtual void draw(bool record) = 0;
     virtual void update() = 0;
-    virtual void cleanupOptix() = 0;
+    virtual void cleanupAssets() = 0;
 
     virtual ~RenderEngine() = default;
 };
diff --git a/zenovis/src/Scene.cpp b/zenovis/src/Scene.cpp
index 347814a3ae..e67af2ea14 100644
--- a/zenovis/src/Scene.cpp
+++ b/zenovis/src/Scene.cpp
@@ -59,7 +59,7 @@ void Scene::cleanUpScene()
     RenderEngine* pEngine = renderMan->getEngine();
     if (pEngine) {
         pEngine->update();
-        pEngine->cleanupOptix();
+        pEngine->cleanupAssets();
     }
 }
 
diff --git a/zenovis/src/bate/RenderEngineBate.cpp b/zenovis/src/bate/RenderEngineBate.cpp
index 2a67d741ac..c4fc891593 100644
--- a/zenovis/src/bate/RenderEngineBate.cpp
+++ b/zenovis/src/bate/RenderEngineBate.cpp
@@ -92,7 +92,7 @@ struct RenderEngineBate : RenderEngine {
         }
     }
 
-    void cleanupOptix() override {
+    void cleanupAssets() override {
 
     }
 };
diff --git a/zenovis/src/optx/RenderEngineOptx.cpp b/zenovis/src/optx/RenderEngineOptx.cpp
index 4e8542328c..049e550789 100644
--- a/zenovis/src/optx/RenderEngineOptx.cpp
+++ b/zenovis/src/optx/RenderEngineOptx.cpp
@@ -32,6 +32,7 @@
 #include "../../xinxinoptix/OptiXStuff.h"
 #include <zeno/types/PrimitiveTools.h>
 #include <zeno/types/StringObject.h>
+#include <tinygltf/json.hpp>
 
 #include <map>
 #include <string>
@@ -694,14 +695,88 @@ struct GraphicsManager {
                 float evnTexStrength = prim_in->userData().get2<float>("evnTexStrength");
                 bool enableHdr = prim_in->userData().get2<bool>("enable");
                 if (!path.empty()) {
-                    if (OptixUtil::sky_tex.has_value() && OptixUtil::sky_tex.value() != path) {
+                    if (OptixUtil::sky_tex.has_value() && OptixUtil::sky_tex.value() != path
+                        && OptixUtil::sky_tex.value() != OptixUtil::default_sky_tex ) {
                         OptixUtil::removeTexture(OptixUtil::sky_tex.value());
                     }
+
                     OptixUtil::sky_tex = path;
                     OptixUtil::addTexture(path);
+                } else {
+                    OptixUtil::sky_tex = OptixUtil::default_sky_tex;
                 }
+
                 xinxinoptix::update_hdr_sky(evnTexRotation, evnTex3DRotation, evnTexStrength);
                 xinxinoptix::using_hdr_sky(enableHdr);
+
+                if (OptixUtil::portal_delayed.has_value()) {
+                    OptixUtil::portal_delayed.value()();
+                    //OptixUtil::portal_delayed.reset();
+                }
+            }
+            else if (prim_in->userData().has<int>("SkyComposer")) {
+
+                auto& attr_dir = prim_in->verts;
+
+                std::vector<zeno::DistantLightData> dlights;
+                dlights.reserve(attr_dir->size());
+                
+                if (attr_dir->size()) {
+
+                    auto& attr_angle = attr_dir.attr<float>("angle");
+                    auto& attr_color = attr_dir.attr<zeno::vec3f>("color");
+                    auto& attr_inten = attr_dir.attr<float>("inten");
+
+                    for (size_t i=0; i<attr_dir->size(); ++i) {
+
+                        auto& dld = dlights.emplace_back();
+                        dld.direction = attr_dir[i];
+                        dld.angle = attr_angle[i];
+                        dld.color = attr_color[i];
+                        dld.intensity = attr_inten[i];
+                    }
+                }
+                xinxinoptix::updateDistantLights(dlights);
+
+                if(prim_in->userData().has<std::string>("portals")) {
+
+                    auto portals_string = prim_in->userData().get2<std::string>("portals");
+                    auto portals_json = nlohmann::json::parse(portals_string);
+
+                    auto ps_string = prim_in->userData().get2<std::string>("psizes");
+                    auto ps_json = nlohmann::json::parse(ps_string);
+
+                    std::vector<Portal> portals {};
+
+                    if (portals_json.is_array() && portals_json.size()%4 == 0) {
+                        
+                        portals.reserve(portals_json.size()/4);
+
+                        auto pack = [&portals_json](size_t i) {
+                            auto x = portals_json[i][0].template get<float>();
+                            auto y = portals_json[i][1].template get<float>();
+                            auto z = portals_json[i][2].template get<float>();
+                            return zeno::vec3f(x, y, z);
+                        };
+                        
+                        for (size_t i=0; i<portals_json.size(); i+=4) {
+                            auto v0 = pack(i+0);
+                            auto v1 = pack(i+1);
+                            auto v2 = pack(i+2);
+                            auto v3 = pack(i+3);
+
+                            uint32_t psize = ps_json[i/4].template get<int>(); 
+                            portals.push_back({v0, v1, v2, v3, psize});
+                        }
+                    } 
+
+                    if (OptixUtil::sky_tex.has_value()) {
+                        xinxinoptix::updatePortalLights(portals);
+                    }
+                    OptixUtil::portal_delayed = [=]() {
+                        xinxinoptix::updatePortalLights(portals);
+                    }; 
+                } //portals
             }
         }
         return sky_found;
@@ -1160,6 +1235,9 @@ struct RenderEngineOptx : RenderEngine, zeno::disable_copy {
                     if (OptixUtil::sky_tex.has_value() && tex == OptixUtil::sky_tex.value()) {
                         continue;
                     }
+                    if (tex == OptixUtil::default_sky_tex) {
+                        continue;
+                    }
                     needToRemoveTexPaths.emplace_back(tex);
                 }
                 for (const auto& need_remove_tex: needToRemoveTexPaths) {
@@ -1389,11 +1467,11 @@ struct RenderEngineOptx : RenderEngine, zeno::disable_copy {
     }
 
     ~RenderEngineOptx() override {
-        xinxinoptix::optixcleanup();
+        xinxinoptix::optixDestroy();
     }
 
-    void cleanupOptix() override {
-
+    void cleanupAssets() override {
+        xinxinoptix::optixCleanup();
     }
 };
 
diff --git a/zenovis/src/zhxx/RenderEngineZhxx.cpp b/zenovis/src/zhxx/RenderEngineZhxx.cpp
index 8ee8ed1621..1dfe495fe7 100644
--- a/zenovis/src/zhxx/RenderEngineZhxx.cpp
+++ b/zenovis/src/zhxx/RenderEngineZhxx.cpp
@@ -123,7 +123,7 @@ struct RenderEngineZhxx : RenderEngine, zeno::disable_copy {
         zenvis::finalize();
     }
 
-    void cleanupOptix() override {
+    void cleanupAssets() override {
 
     }
 };
diff --git a/zenovis/xinxinoptix/CMakeLists.txt b/zenovis/xinxinoptix/CMakeLists.txt
index 3f558d2378..43e552e726 100644
--- a/zenovis/xinxinoptix/CMakeLists.txt
+++ b/zenovis/xinxinoptix/CMakeLists.txt
@@ -11,6 +11,7 @@ target_sources(zenovis PRIVATE
   volume/optixVolume.h
   optixSphere.cpp
   optixSphere.h
+  Portal.h
   Shape.h
   XAS.h
 
@@ -36,10 +37,6 @@ target_sources(zenovis PRIVATE
   SDK/sutil/Preprocessor.h      SDK/sutil/WorkDistribution.h
   )
 
-#if (OPTIX_USE_20XX)
-    #target_compile_definitions(zenovis PRIVATE -DUSING_20XX)
-#endif()
-
 find_package(CUDAToolkit REQUIRED COMPONENTS cudart nvrtc REQUIRED)
 target_link_libraries(zenovis PRIVATE CUDA::cudart CUDA::nvrtc)
 
@@ -133,6 +130,8 @@ set(FILE_LIST
     ${CMAKE_CURRENT_SOURCE_DIR}/@LightBounds.h
     ${CMAKE_CURRENT_SOURCE_DIR}/@LightTree.h
 
+    ${CMAKE_CURRENT_SOURCE_DIR}/@Portal.h
+
     ${CMAKE_CURRENT_SOURCE_DIR}/@Sampling.h    
     ${CMAKE_CURRENT_SOURCE_DIR}/@Shape.h
     ${CMAKE_CURRENT_SOURCE_DIR}/@Light.cu
diff --git a/zenovis/xinxinoptix/Light.h b/zenovis/xinxinoptix/Light.h
index 6554fe16de..97bdd8697d 100644
--- a/zenovis/xinxinoptix/Light.h
+++ b/zenovis/xinxinoptix/Light.h
@@ -7,6 +7,8 @@
 // #include "DisneyBSDF.h"
 #include "proceduralSky.h"
 
+#include "Portal.h"
+
 static __inline__ __device__
 vec3 ImportanceSampleEnv(float* env_cdf, int* env_start, int nx, int ny, float p, float &pdf)
 {
@@ -210,8 +212,7 @@ static __forceinline__ __device__
 void DirectLighting(RadiancePRD *prd, ShadowPRD& shadowPRD, const float3& shadingP, const float3& ray_dir, 
                     TypeEvalBxDF& evalBxDF, TypeAux* taskAux=nullptr, float3* RadianceWithoutShadow=nullptr) {
 
-    const float3 wo = normalize(-ray_dir); 
-    float3 light_attenuation = vec3(1.0f);
+    const float3 wo = normalize(-ray_dir);
 
     const float _SKY_PROB_ = params.skyLightProbablity();
 
@@ -415,7 +416,7 @@ void DirectLighting(RadiancePRD *prd, ShadowPRD& shadowPRD, const float3& shadin
             shadowPRD.maxDistance = lsr.dist;
             
             traceOcclusion(params.handle, shadowPRD.origin, lsr.dir, 0, lsr.dist, &shadowPRD);
-            light_attenuation = shadowPRD.attanuation;
+            auto light_attenuation = shadowPRD.attanuation;
 
             if (nullptr==RadianceWithoutShadow && lengthSquared(light_attenuation) == 0.0f) return;
 
@@ -446,69 +447,125 @@ void DirectLighting(RadiancePRD *prd, ShadowPRD& shadowPRD, const float3& shadin
     
     } else {
 
-        float env_weight_sum = 1e-8f;
-        int NSamples = prd->depth<=2?1:1;//16 / pow(4.0f, (float)prd->depth-1);
-        for(int samples=0;samples<NSamples;samples++) {
-
-            bool hasenv = params.skynx | params.skyny;
-            hasenv = params.usingHdrSky && hasenv;
-            float envpdf = 1.0f;
-
-            vec3 sunLightDir = hasenv? ImportanceSampleEnv(params.skycdf, params.sky_start,
-                                                            params.skynx, params.skyny, rnd(prd->seed), envpdf)
-                                    : vec3(params.sunLightDirX, params.sunLightDirY, params.sunLightDirZ);
-            auto sun_dir = BRDFBasics::halfPlaneSample(prd->seed, sunLightDir,
-                                                    params.sunSoftness * 0.0f); //perturb the sun to have some softness
-            sun_dir = hasenv ? normalize(sunLightDir):normalize(sun_dir);
-
-            float tmpPdf;
-            auto illum = float3(envSky(sun_dir, sunLightDir, make_float3(0., 0., 1.),
-                                        40, // be careful
-                                        .45, 15., 1.030725f * 0.3f, params.elapsedTime, tmpPdf));
-            if(tmpPdf <= 0.0f) { return; }
+        auto shadeTask = [&](float3 sampleDir, float samplePDF, float3 illum, const bool mis) {\
 
-            auto Ldir = sun_dir;
 
-            if (envpdf < __FLT_DENORM_MIN__) {
-                return;
-            }
-
-            shadowPRD.maxDistance = 1e16f;
-            traceOcclusion(params.handle, shadowPRD.origin, sun_dir,
-                        1e-5f, // tmin
-                        1e16f, // tmax,
+            shadowPRD.attanuation = vec3(1.0);
+            shadowPRD.maxDistance = FLT_MAX;
+            traceOcclusion(params.handle, shadowPRD.origin, sampleDir,
+                        0, // tmin
+                        FLT_MAX, // tmax,
                         &shadowPRD);
-            light_attenuation = shadowPRD.attanuation;
 
-            if (nullptr==RadianceWithoutShadow && lengthSquared(light_attenuation) == 0.0f) return;
+            if (nullptr==RadianceWithoutShadow && lengthSquared(shadowPRD.attanuation) == 0.0f) return;
 
-            auto inverseProb = 1.0f/_SKY_PROB_;
-            auto bxdf_value = evalBxDF(sun_dir, wo, scatterPDF);
+            auto bxdf_value = evalBxDF(sampleDir, wo, scatterPDF);
 
-            float tmp = 1.0f;
+            float tmp = 1.0f / samplePDF;
 
-            if constexpr(_MIS_) {
-                float misWeight = BRDFBasics::PowerHeuristic(tmpPdf, scatterPDF);
+            if (mis) {
+                float misWeight = BRDFBasics::PowerHeuristic(samplePDF, scatterPDF);
                 misWeight = misWeight>0.0f?misWeight:1.0f;
                 misWeight = scatterPDF>1e-5f?misWeight:0.0f;
-                misWeight = tmpPdf>1e-5f?misWeight:0.0f;
+                misWeight = samplePDF>1e-5f?misWeight:0.0f;
 
-                tmp = (1.0f / NSamples) * misWeight * inverseProb  / tmpPdf;
-            } else {
-                tmp = (1.0f / NSamples) * inverseProb / tmpPdf;
-            }
+                tmp *= misWeight;
+            } 
 
             float3 radianceNoShadow = illum * tmp * bxdf_value; 
 
             if (nullptr != RadianceWithoutShadow) {
-                *RadianceWithoutShadow = radianceNoShadow;
+                *RadianceWithoutShadow += radianceNoShadow;
             }
 
             if constexpr (!detail::is_void<TypeAux>::value) {
-                (*taskAux)(illum * tmp * light_attenuation);
+                (*taskAux)(illum * tmp * shadowPRD.attanuation);
             }// TypeAux
 
-            prd->radiance += radianceNoShadow * light_attenuation; // with shadow
+            prd->radiance += radianceNoShadow * shadowPRD.attanuation; // with shadow
+        }; // shadeTask
+
+        UF = UF / _SKY_PROB_;
+        UF = clamp(UF, 0.0f, 1.0f);
+
+        auto binsearch = [&](float* cdf, uint min, uint max) {
+            //auto idx = min;
+            while(min < max) {
+                auto _idx_ = (min + max) / 2;
+                auto _cdf_ = cdf[_idx_];
+
+                if (_cdf_ > UF) {
+                    max = _idx_; continue; //include
+                }
+                if (_cdf_ < UF) {
+                    min = _idx_+1; continue;
+                }
+                min = _idx_; break;
+            }
+            return min;
+        };
+
+        auto dlights = reinterpret_cast<DistantLightList*>(params.dlights_ptr);
+        
+        if (nullptr != dlights && dlights->COUNT()) {
+
+            auto idx = binsearch(dlights->cdf, 0, dlights->COUNT());
+            auto& dlight = dlights->list[idx];
+            auto dlight_dir = reinterpret_cast<vec3&>(dlight.direction);
+
+            auto sample_dir = BRDFBasics::halfPlaneSample(prd->seed, dlight_dir, dlight.angle/180.0f);
+            auto sample_prob = _SKY_PROB_ / dlights->COUNT();
+
+            if (dlight.intensity > 0) {
+                auto ccc = dlight.color * dlight.intensity;
+                auto illum = reinterpret_cast<float3&>(ccc);
+                shadeTask(sample_dir, sample_prob, illum, false);
+            }
+        }
+        
+        auto plights = reinterpret_cast<PortalLightList*>(params.plights_ptr);
+
+        if (plights != nullptr && plights->COUNT()) {
+
+            uint idx = binsearch(plights->cdf, 0, plights->COUNT());
+            auto plight = &plights->list[idx];
+
+            LightSampleRecord lsr; lsr.PDF = 0.0f;
+            float2 uu = { prd->rndf(), prd->rndf() };
+            float3 color {};
+            
+            plight->sample(lsr, reinterpret_cast<const Vector3f&>(shadingP), uu, color);
+            
+            lsr.PDF *= plights->pdf[idx] * _SKY_PROB_;
+            if (lsr.PDF > 0) {
+                //auto suv = sphereUV(lsr.dir, true);
+                //color = (vec3)texture2D(params.sky_texture, vec2(suv.x, suv.y));
+                shadeTask(lsr.dir, lsr.PDF, color * params.sky_strength, false);
+            }
+            return;
+        }
+
+        { // SKY
+            bool hasenv = params.skynx | params.skyny;
+            hasenv = params.usingHdrSky && hasenv;
+            float envpdf = 1.0f;
+
+            vec3 sunLightDir = vec3(params.sunLightDirX, params.sunLightDirY, params.sunLightDirZ);
+
+            vec3 sample_dir = hasenv? ImportanceSampleEnv(params.skycdf, params.sky_start,
+                                                            params.skynx, params.skyny, rnd(prd->seed), envpdf)
+                                    : BRDFBasics::halfPlaneSample(prd->seed, sunLightDir,
+                                                    params.sunSoftness * 0.0f);
+            sample_dir = normalize(sample_dir);
+
+            float samplePDF;
+            float3 illum = envSky(sample_dir, sunLightDir, make_float3(0., 0., 1.),
+                                        40, // be careful
+                                        .45, 15., 1.030725f * 0.3f, params.elapsedTime, samplePDF);\
+            samplePDF *= _SKY_PROB_;
+            if(samplePDF <= 0.0f) { return; }
+
+            shadeTask(sample_dir, samplePDF, illum, true);
         }
     }
 };
\ No newline at end of file
diff --git a/zenovis/xinxinoptix/OptiXStuff.h b/zenovis/xinxinoptix/OptiXStuff.h
index c1cbaab36b..8c7d168f2f 100644
--- a/zenovis/xinxinoptix/OptiXStuff.h
+++ b/zenovis/xinxinoptix/OptiXStuff.h
@@ -336,9 +336,21 @@ inline void createRTProgramGroups(OptixDeviceContext &context, OptixModule &_mod
 }
 struct cuTexture{
     std::string md5;
-    cudaArray_t gpuImageArray;
-    cudaTextureObject_t texture;
-    cuTexture(){gpuImageArray = nullptr;texture=0;}
+    
+    cudaArray_t gpuImageArray = nullptr;
+    cudaTextureObject_t texture = 0llu;
+
+    uint32_t width, height;
+    float average = 0.0f;
+
+    std::vector<float> cdf;
+    std::vector<float> pdf; 
+    std::vector<int> start;
+
+    std::vector<float> rawData;
+
+    cuTexture() {}
+    cuTexture(uint32_t w, uint32_t h) : width(w), height(h) {}
     ~cuTexture()
     {
         if(gpuImageArray!=nullptr)
@@ -355,7 +367,7 @@ inline sutil::Texture loadCubeMap(const std::string& ppm_filename)
 }
 inline std::shared_ptr<cuTexture> makeCudaTexture(unsigned char* img, int nx, int ny, int nc)
 {
-    auto texture = std::make_shared<cuTexture>();
+    auto texture = std::make_shared<cuTexture>(nx, ny);
     std::vector<uchar4> data;
     data.resize(nx*ny);
     for(int j=0;j<ny;j++)
@@ -413,7 +425,7 @@ inline std::shared_ptr<cuTexture> makeCudaTexture(unsigned char* img, int nx, in
 }
 inline std::shared_ptr<cuTexture> makeCudaTexture(float* img, int nx, int ny, int nc)
 {
-    auto texture = std::make_shared<cuTexture>();
+    auto texture = std::make_shared<cuTexture>(nx, ny);
     std::vector<float4> data;
     data.resize(nx*ny);
     for(int j=0;j<ny;j++)
@@ -612,14 +624,14 @@ inline std::vector<float> loadIES(const std::string& path, float& coneAngle)
 
     return iesData;
 }
+
 inline std::map<std::string, std::shared_ptr<cuTexture>> g_tex;
 inline std::map<std::string, std::filesystem::file_time_type> g_tex_last_write_time;
 inline std::map<std::string, std::string> md5_path_mapping;
 inline std::optional<std::string> sky_tex;
-inline std::map<std::string, int> sky_nx_map;
-inline std::map<std::string, int> sky_ny_map;
-inline std::map<std::string, float> sky_avg_map;
+inline std::string default_sky_tex;
 
+inline std::optional<std::function<void(void)>> portal_delayed;
 
 struct WrapperIES {
     raii<CUdeviceptr> ptr;
@@ -628,18 +640,19 @@ struct WrapperIES {
 
 inline std::map<std::string, WrapperIES> g_ies;
 
-inline std::map<std::string, std::vector<float>> sky_cdf_map;
-inline std::map<std::string, std::vector<float>> sky_pdf_map;
-inline std::map<std::string, std::vector<int>>   sky_start_map;
-
-template<typename T>
-inline void calc_sky_cdf_map(int nx, int ny, int nc, T *img) {
-    auto &sky_nx = sky_nx_map[sky_tex.value()];
-    auto &sky_ny = sky_ny_map[sky_tex.value()];
-    auto &sky_cdf = sky_cdf_map[sky_tex.value()];
-    auto &sky_pdf = sky_pdf_map[sky_tex.value()];
-    auto &sky_start = sky_start_map[sky_tex.value()];
-    auto &sky_avg = sky_avg_map[sky_tex.value()];
+inline void calc_sky_cdf_map(int nx, int ny, int nc, std::function<float(uint32_t)>& look) {
+
+    auto& tex = g_tex[sky_tex.value()];
+
+    auto &sky_nx = tex->width;
+    auto &sky_ny = tex->height;
+
+    auto &sky_avg = tex->average;
+
+    auto &sky_cdf = tex->cdf;
+    auto &sky_pdf = tex->pdf;
+    auto &sky_start = tex->start;
+
     sky_nx = nx;
     sky_ny = ny;
     //we need to recompute cdf
@@ -658,7 +671,7 @@ inline void calc_sky_cdf_map(int nx, int ny, int nc, T *img) {
             size_t idx2 = jj*nx*nc + ii*nc;
             size_t idx = jj*nx + ii;
             float illum = 0.0f;
-            auto color = zeno::vec3f(img[idx2+0], img[idx2+1], img[idx2+2]);
+            auto color = zeno::vec3f(look(idx2+0), look(idx2+1), look(idx2+2));
             illum = zeno::dot(color, zeno::vec3f(0.33333333f,0.33333333f, 0.33333333f));
             //illum = illum > 0.5? illum : 0.0f;
             illum = abs(illum) * sinf(3.1415926f*((float)jj + 0.5f)/(float)ny);
@@ -687,6 +700,7 @@ inline void calc_sky_cdf_map(int nx, int ny, int nc, T *img) {
         }
     }
 }
+
 static std::string calculateMD5(const std::vector<char>& input) {
     CryptoPP::byte digest[CryptoPP::Weak::MD5::DIGESTSIZE];
     CryptoPP::Weak::MD5().CalculateDigest(digest, (const CryptoPP::byte*)input.data(), input.size());
@@ -724,9 +738,13 @@ inline void addTexture(std::string path)
     else {
         md5_path_mapping[md5Hash] = path;
     }
+
     int nx, ny, nc;
     stbi_set_flip_vertically_on_load(true);
 
+    std::function<float(uint32_t)> lookupTexture = [](uint32_t x) {return 0.0f;};
+    std::function<void(void)>     cleanupTexture = [](){};
+
     if (zeno::ends_with(path, ".exr", false)) {
         float* rgba;
         const char* err;
@@ -748,12 +766,15 @@ inline void addTexture(std::string path)
             }
         }
         assert(rgba);
-        if(sky_tex.value() == path)//if this is a loading of a sky texture
-        {
-            calc_sky_cdf_map(nx, ny, nc, rgba);
-        }
+
         g_tex[path] = makeCudaTexture(rgba, nx, ny, nc);
-        free(rgba);
+
+        lookupTexture = [&](uint32_t idx) {
+            return rgba[idx];
+        };
+        cleanupTexture = [&]() {
+            free(rgba);
+        };
     }
     else if (zeno::ends_with(path, ".ies", false)) {
         float coneAngle;
@@ -783,13 +804,11 @@ inline void addTexture(std::string path)
             g_tex[path] = std::make_shared<cuTexture>();
             return;
         }
-        int nx = std::max(img->userData().get2<int>("w"), 1);
-        int ny = std::max(img->userData().get2<int>("h"), 1);
+        nx = std::max(img->userData().get2<int>("w"), 1);
+        ny = std::max(img->userData().get2<int>("h"), 1);
         int channels = std::max(img->userData().get2<int>("channels"), 3);
-        if(sky_tex.value() == path)//if this is a loading of a sky texture
-        {
-            calc_sky_cdf_map(nx, ny, 3, (float *)img->verts.data());
-        }
+        nc = 3;
+
         if (channels == 3) {
             std::vector<unsigned char> ucdata;
             ucdata.resize(img->verts.size()*3);
@@ -811,6 +830,11 @@ inline void addTexture(std::string path)
             }
             g_tex[path] = makeCudaTexture((unsigned char *)data.data(), nx, ny, 4);
         }
+        
+        lookupTexture = [img](uint32_t idx) {
+            auto ptr = (float*)img->verts->data();
+            return ptr[idx];
+        };
     }
     else if (stbi_is_hdr(native_path.c_str())) {
         float *img = stbi_loadf(native_path.c_str(), &nx, &ny, &nc, 0);
@@ -822,12 +846,15 @@ inline void addTexture(std::string path)
         nx = std::max(nx, 1);
         ny = std::max(ny, 1);
         assert(img);
-        if(sky_tex.value() == path)//if this is a loading of a sky texture
-        {
-            calc_sky_cdf_map(nx, ny, nc, img);
-        }
+        
         g_tex[path] = makeCudaTexture(img, nx, ny, nc);
-        stbi_image_free(img);
+
+        lookupTexture = [&](uint32_t idx) {
+            return img[idx];
+        };
+        cleanupTexture = [&]() {
+            stbi_image_free(img);
+        };
     }
     else {
         unsigned char *img = stbi_load(native_path.c_str(), &nx, &ny, &nc, 0);
@@ -839,15 +866,31 @@ inline void addTexture(std::string path)
         nx = std::max(nx, 1);
         ny = std::max(ny, 1);
         assert(img);
-        if(sky_tex.value() == path)//if this is a loading of a sky texture
-        {
-            calc_sky_cdf_map(nx, ny, nc, img);
-        }
+        
         g_tex[path] = makeCudaTexture(img, nx, ny, nc);
-        stbi_image_free(img);
+
+        lookupTexture = [&](uint32_t idx) {
+            return (float)img[idx] / 255;
+        };
+        cleanupTexture = [&]() {
+            stbi_image_free(img);
+        };
     }
     g_tex[path]->md5 = md5Hash;
 
+    if(sky_tex.value() == path)
+    {
+        calc_sky_cdf_map(nx, ny, nc, lookupTexture);
+        auto& tex = g_tex[sky_tex.value()];
+        auto float_count = nx * ny * nc;
+        tex->rawData.resize(float_count);
+
+        for (size_t i=0; i<float_count; ++i) {
+            tex->rawData.at(i) = lookupTexture(i);
+        }
+    }
+    cleanupTexture();
+
     for (auto i = g_tex.begin(); i != g_tex.end(); i++) {
         zeno::log_info("-{}", i->first);
     }
@@ -863,12 +906,6 @@ inline void removeTexture(std::string path) {
             zeno::log_error("removeTexture: {} not exists!", path);
         }
         g_tex.erase(path);
-        sky_nx_map.erase(path);
-        sky_ny_map.erase(path);
-        sky_cdf_map.erase(path);
-        sky_pdf_map.erase(path);
-        sky_start_map.erase(path);
-        sky_avg_map.erase(path);
         g_tex_last_write_time.erase(path);
     }
 }
diff --git a/zenovis/xinxinoptix/PTKernel.cu b/zenovis/xinxinoptix/PTKernel.cu
index 75dfd235fb..e5ee736af3 100644
--- a/zenovis/xinxinoptix/PTKernel.cu
+++ b/zenovis/xinxinoptix/PTKernel.cu
@@ -306,7 +306,7 @@ extern "C" __global__ void __raygen__rg()
 
             //if(prd.depth>prd.max_depth) {
             float RRprob = max(max(prd.attenuation.x, prd.attenuation.y), prd.attenuation.z);
-            if(rnd(prd.seed) > RRprob || prd.depth > prd.max_depth*2) {
+            if(rnd(prd.seed) > RRprob || prd.depth > prd.max_depth) {
                 prd.done=true;
             } else {
                 prd.attenuation = prd.attenuation / RRprob;
diff --git a/zenovis/xinxinoptix/Portal.h b/zenovis/xinxinoptix/Portal.h
new file mode 100644
index 0000000000..fc2125fcec
--- /dev/null
+++ b/zenovis/xinxinoptix/Portal.h
@@ -0,0 +1,716 @@
+#pragma once
+
+#ifndef __CUDACC_RTC__
+#include <vector>
+#include <algorithm>
+#include <cassert>
+#include <stdexcept>
+#include <raiicuda.h>
+#include <zeno/vec.h>
+
+#include <zeno/utils/pfm.h>
+#include <glm/gtc/matrix_transform.hpp>
+
+#endif
+
+#include <zeno/types/LightObject.h>
+#include <Sampling.h>
+
+struct Bounds2f {
+    Vector2f pMin = Vector2f {FLT_MAX, FLT_MAX};
+    Vector2f pMax = -Vector2f {FLT_MAX, FLT_MAX};
+
+    bool contains(Vector2f p) {
+        if (p[0] < pMin[0] || p[1] < pMin[1])
+            return false;
+        if (p[0] > pMax[0] || p[1] > pMax[1])
+            return false;
+
+        return true;
+    }
+
+    float area() {
+        auto delta = pMax - pMin;
+        return delta[0] * delta[1];
+    }
+};
+
+namespace xx {
+
+template <typename T>
+struct Array2D {
+    
+#ifndef __CUDACC_RTC__
+    std::vector<T> data;
+    xinxinoptix::raii<CUdeviceptr> buffer;
+#else
+    T* data;
+#endif
+    uint32_t _x, _y;
+
+    uint32_t XSize() const { return _x; }
+    uint32_t YSize() const { return _y; }
+
+    Array2D() = default;
+
+#ifndef __CUDACC_RTC__
+
+    Array2D(uint32_t x, uint32_t y) {
+        _x = x; _y = y;
+        data.resize(_x * _y, {});
+    }
+
+    auto upload() {
+
+        size_t byte_size = sizeof(T) * data.size();
+
+        buffer.resize( byte_size );
+        cudaMemcpy((void*)buffer.handle, data.data(), byte_size, cudaMemcpyHostToDevice);
+
+        struct Dummy {
+            void* ptr;
+            uint32_t _x, _y; 
+        };
+
+        return Dummy {
+            (void*)buffer.handle, _x, _y
+        };
+    }
+
+#endif
+
+    T &operator()(uint32_t x, uint32_t y) { 
+        size_t idx = x + y * _x;
+        return data[idx];
+    }
+
+    const T &operator()(uint32_t x, uint32_t y) const { 
+        size_t idx = x + y * _x;
+        return data[idx];
+    }
+
+    T lookUV(float2 uv) const {
+
+        auto xf = uv.x * _x;
+        auto yf = uv.y * _y;
+        xf -= 0.5f; yf -= 0.5f;
+
+        auto xi = (int)floor(xf);
+        auto yi = (int)floor(yf); 
+
+        auto dx = xf - xi;
+        auto dy = yf - yi;
+
+        auto v00 = lookUp(  xi,   yi);
+        auto v10 = lookUp(1+xi,   yi); 
+        auto v01 = lookUp(  xi, 1+yi);
+        auto v11 = lookUp(1+xi, 1+yi);
+
+        return 
+        
+        (v00 * (1-dx) + v10 * dx) * (1-dy) 
+             +
+        (v01 * (1-dx) + v11 * dx) * dy;
+    }
+
+    T lookUp(int x, int y) const {
+
+        if (x<0 || x >= _x) return {};
+        if (y<0 || y >= _y) return {};
+        size_t idx = x + y * _x;
+        return data[idx];
+    }
+};
+
+};
+
+struct SummedAreaTable {
+  public:
+    // SummedAreaTable Public Methods
+    SummedAreaTable() = default;
+
+#ifndef __CUDACC_RTC__
+    SummedAreaTable(const xx::Array2D<float> &values)
+        : sum(values.XSize(), values.YSize()) {
+
+        sum(0, 0) = values(0, 0);
+        // Compute sums along first row and column
+        for (int x = 1; x < sum.XSize(); ++x)
+            sum(x, 0) = values(x, 0) + sum(x - 1, 0);
+        for (int y = 1; y < sum.YSize(); ++y)
+            sum(0, y) = values(0, y) + sum(0, y - 1);
+
+        // Compute sums for the remainder of the entries
+        for (int y = 1; y < sum.YSize(); ++y)
+            for (int x = 1; x < sum.XSize(); ++x)
+                sum(x, y) = (values(x, y) + sum(x - 1, y) + sum(x, y - 1) - sum(x - 1, y - 1));
+    }
+
+    auto upload() {
+        return sum.upload();
+    }
+
+#endif
+
+    float Integral(Bounds2f extent) const {
+        double s =  ((double)Lookup(extent.pMax[0], extent.pMax[1]) - (double)Lookup(extent.pMin[0], extent.pMax[1])) 
+                        +
+                    ((double)Lookup(extent.pMin[0], extent.pMin[1]) - (double)Lookup(extent.pMax[0], extent.pMin[1]));
+        return fmaxf(s / (sum.XSize() * sum.YSize()), 0);
+    }
+
+  private:
+    // SummedAreaTable Private Methods
+    float Lookup(float x, float y) const {
+        // Rescale $(x,y)$ to table resolution and compute integer coordinates
+        x = x * sum.XSize();
+        y = y * sum.YSize();
+
+        x = x - 0.5f;
+        y = y - 0.5f;
+
+        int x0 = (int)x; 
+        int y0 = (int)y;
+
+        float dx = x - int(x);
+        float dy = y - int(y);
+
+        // Bilinearly interpolate between surrounding table values
+        float v00 = LookupInt(x0, y0), v10 = LookupInt(x0 + 1, y0);
+        float v01 = LookupInt(x0, y0 + 1), v11 = LookupInt(x0 + 1, y0 + 1);
+        
+        return (1 - dx) * ( (1 - dy) * v00 + dy * v01 ) 
+                    + 
+                    dx  * ( (1 - dy) * v10 + dy * v11 );
+    }
+
+    float LookupInt(int x, int y) const {
+        // Return zero at lower boundaries
+        if (x <= 0 || y <= 0)
+            return 0;
+
+        // Reindex $(x,y)$ and return actual stored value
+        x = min(x-1, (int)sum.XSize() - 1);
+        y = min(y-1, (int)sum.YSize() - 1);
+        return sum(x, y);
+    }
+
+    // SummedAreaTable Private Members
+    xx::Array2D<double> sum;
+};
+
+template<typename T>
+inline T BiLinear(T* data, uint width, uint height, float2 pos) {
+
+    pos -= {0.5f, 0.5f};
+
+    auto lowX = (int)std::floor(pos.x), highX = lowX+1;
+    auto lowY = (int)std::floor(pos.y), highY = lowY+1;
+
+    auto ratioX = pos.x - lowX;
+    auto ratioY = pos.y - lowY;
+
+    auto lookUp = [&](int x, int y) {
+        if (x < 0 || x >= width ) return T{};
+        if (y < 0 || y >= height) return T{};
+
+        return *(data + (y * width + x));
+    };
+
+    auto v00 = lookUp(lowX, lowY);
+    auto v10 = lookUp(highX,lowY);
+    auto vv0 = v00 * (1-ratioX) + v10 * ratioX;
+
+    auto v01 = lookUp(lowX, highY);
+    auto v11 = lookUp(highX,highY);
+    auto vv1 = v01 * (1-ratioX) + v11 * ratioX;
+
+    return vv0 * (1-ratioY) + vv1 * ratioY;
+}
+
+struct Portal {
+    Vector3f p0, p1, p2, p3;
+    uint32_t psize;
+};
+
+struct PortalLight {
+
+    xx::Array2D<float3> image {};
+    xx::Array2D<float>  dist  {};
+    SummedAreaTable sat;
+
+    Portal portal;
+    Vector3f X,Y,Z;
+    //PortalLight() = default;
+
+#ifndef __CUDACC_RTC__
+
+    auto pack() {
+
+        auto image_dummy = image.upload();
+        auto dist_dummy = dist.upload();
+        auto sat_dummy = sat.upload();
+
+        struct Dummy {
+            
+            typeof(image_dummy) image; 
+            typeof(dist_dummy) dist;
+            typeof(sat_dummy) sat;
+            Portal portal;
+            Vector3f X,Y,Z;
+        };
+
+        return Dummy {  image_dummy, dist_dummy, sat_dummy, portal, X, Y, Z };
+    }
+
+    PortalLight(const Portal& por, float3* texture, uint tex_width, uint tex_height, glm::mat4* rotate=nullptr) : portal(por) {
+
+        Vector3f p01 = normalize(portal.p1 - portal.p0);
+        Vector3f p12 = normalize(portal.p2 - portal.p1);
+        Vector3f p32 = normalize(portal.p2 - portal.p3);
+        Vector3f p03 = normalize(portal.p3 - portal.p0);
+        // Do opposite edges have the same direction?
+        if (std::abs(dot(p01, p32) - 1) > .001 || std::abs(dot(p12, p03) - 1) > .001)
+            throw std::runtime_error("Infinite light portal isn't a planar quadrilateral");
+
+        // Sides perpendicular?
+        if (std::abs(dot(p01, p12)) > .001 || std::abs(dot(p12, p32)) > .001 ||
+            std::abs(dot(p32, p03)) > .001 || std::abs(dot(p03, p01)) > .001)
+            throw std::runtime_error("Infinite light portal isn't a planar quadrilateral");
+
+        X = p03, Y = p01, Z = -cross(X, Y);
+
+        uint pixel_count_x = por.psize, pixel_count_y = por.psize;
+        uint pixel_count = pixel_count_x * pixel_count_y;
+
+        image = xx::Array2D<float3>(pixel_count_x, pixel_count_y);
+        dist = xx::Array2D<float>(pixel_count_x, pixel_count_y);
+
+        auto luminance = [](float3 c) {
+            return dot(c, float3{0.2722287, 0.6740818, 0.0536895});
+        };
+
+        for (uint i=0; i<pixel_count_x; ++i) {
+            for (uint j=0; j<pixel_count_y; ++j) {
+
+                Vector2f uv { (i+0.5f)/pixel_count_x, (j+0.5f)/pixel_count_y };
+                float duv_dw;
+                zeno::vec3f world_dir = uv_direction(reinterpret_cast<float2&>(uv), &duv_dw);
+
+                if (rotate != nullptr && *rotate != glm::mat4(1.0f)) {
+                    glm::vec4 tmp = glm::vec4(world_dir[0], world_dir[1], world_dir[2], 0.0f);
+                    tmp = tmp * (*rotate);
+
+                    world_dir = {tmp.x, tmp.y, tmp.z};
+                }
+
+                auto suv = sphereUV(reinterpret_cast<float3&>(world_dir), true);
+                auto pos = (*(float2*)&suv) * make_float2(tex_width, tex_height);
+
+                auto pixel = BiLinear(texture, tex_width, tex_height, pos);
+                auto average = (pixel.x + pixel.y + pixel.z) / 3.0f;
+                //average = luminance(pixel);
+                //average *= std::sin(M_PIf * suv.y);
+
+                image(i, j) = pixel;
+                dist(i, j) = duv_dw * average;
+            } // j
+        } // i
+
+        sat = SummedAreaTable(dist);
+
+    #if !defined( NDEBUG ) 
+        zeno::write_pfm("portal.pfm", image.XSize(), image.YSize(), (float*)image.data.data());
+        zeno::write_pfm("dist.pfm", dist.XSize(), dist.YSize(), (float*)dist.data.data(), true);
+    #endif
+    }
+
+#endif
+
+    inline float area() {
+        auto a = length(portal.p1 - portal.p0);
+        auto b = length(portal.p2 - portal.p1);
+        return a * b;
+    }
+
+    float phi() {
+
+        float3 sum {};
+        for (uint y=0; y<image.YSize(); ++y) {
+            for (uint x=0; x<image.XSize(); ++x) {
+                auto value = image.lookUp(x, y);
+
+                float2 uv = { (x+0.5f)/image.XSize(), 
+                              (y+0.5f)/image.YSize()};
+                float duvdw;
+                uv_direction(uv, &duvdw);
+                sum += value / duvdw;
+            }
+        }
+        
+        auto average = (sum.x + sum.y + sum.z) / 3.0f;
+        average /= (image.XSize() * image.YSize());
+        return area() * average;
+    }
+
+    float3 Le(const Vector3f& ray_origin, const Vector3f& ray_dir) {
+        Bounds2f bds;
+        auto valid = ImageBounds(ray_origin, bds);
+        if (!valid) {return {};}
+
+        auto x = dot(ray_dir, X);
+        auto y = dot(ray_dir, Y);
+        auto z = dot(ray_dir, Z);
+
+        if (z <= 0) {return {};}
+
+        auto sinL = sqrt(1.0f - z * z);
+        auto angleX = asin(x / sinL);
+        auto angleY = -acos(y / sinL);
+
+         Vector2f uv = {
+            ( angleX + M_PI_2f ) / M_PIf,
+            ( angleY + M_PIf ) / M_PIf
+        };
+
+        if (bds.contains(uv)) {
+            return image.lookUp(uv[0], uv[1]);
+        } else {
+            return {};
+        }
+    }
+
+    bool ImageBounds(const Vector3f& p, Bounds2f& bounds) const {
+
+        auto v0 = normalize(portal.p0 - p);
+        auto v1 = normalize(portal.p2 - p);
+
+        auto x0 = dot(v0, X);
+        auto y0 = dot(v0, Y);
+        auto z0 = dot(v0, Z);
+
+        if (z0 <= 0) { return false; }
+
+        auto angleX0 = atan2(x0, z0);
+        auto angleY0 = atan2(y0, z0);
+
+        auto x1 = dot(v1, X);
+        auto y1 = dot(v1, Y);
+        auto z1 = dot(v1, Z);
+        
+        if (z1 <= 0) { return false; }
+
+        auto angleX1 = atan2(x1, z1);
+        auto angleY1 = atan2(y1, z1);
+
+        if (angleX0 >= angleX1 || angleY0 >= angleY1) 
+        {
+            return false;
+        } 
+
+        Vector2f uv0 = { 
+            ( angleX0 + M_PI_2f ) / M_PIf,
+            ( angleY0 + M_PI_2f ) / M_PIf
+        };
+
+        Vector2f uv1 = {
+            ( angleX1 + M_PI_2f ) / M_PIf,
+            ( angleY1 + M_PI_2f ) / M_PIf
+        };
+
+        bounds = Bounds2f{ uv0, uv1 };
+        return true; 
+    }
+
+    template <typename CDF>
+    static float SampleBisection(CDF P, const float u, float min, float max, uint n) {
+        assert(0.0<=min && min < max && max<=1.0);
+
+        while (min < max && ( (n * max) - (n * min)) > 1) {
+
+            assert(P(min) <= u);
+            assert(P(max) >= u);
+            float mid = (min + max) / 2;
+            auto PM = P(mid);
+            //PM = clamp(PM, 0.0f, 1.0f); 
+
+            if (PM > u)
+                max = mid;
+            else
+                min = mid;
+        }
+
+        // Find sample by interpolating between _min_ and _max_
+        float t = (u - P(min)) / (P(max) - P(min));
+        return clamp(pbrt::Lerp(t, min, max), min, max);
+    }
+
+    float Eval(float2 p) const {
+        float2 pi{ fminf(p.x * dist.XSize(), dist.XSize() - 1),
+                   fminf(p.y * dist.YSize(), dist.YSize() - 1) };
+        //return dist.lookUp((int)pi.x, (int)pi.y);
+        return dist.lookUV(p);
+    }
+
+    Vector2f direction_uv(Vector3f dir, float *duvdw=nullptr) {
+
+        auto x = dot(dir, X);
+        auto y = dot(dir, Y);
+        auto z = dot(dir, Z);
+
+        if (z <= 0) {return {};}
+
+        auto w = Vector3f{x, y, z};
+
+        if (duvdw)
+            *duvdw = pbrt::Sqr(M_PIf) * (1 - pbrt::Sqr(w[0])) * (1 - pbrt::Sqr(w[1])) / w[2]; 
+
+        auto sinL = sqrt(1.0f - z * z);
+        auto angleX = asin(x / sinL);
+        auto angleY = -acos(y / sinL);
+
+         Vector2f uv = {
+            ( angleX + M_PI_2f ) / M_PIf,
+            ( angleY + M_PIf ) / M_PIf
+        };
+    }
+
+    Vector3f uv_direction(float2 uv, float* duvdw=nullptr) {
+
+        float alpha = -M_PIf / 2 + uv.x * M_PIf; 
+        float beta  = -M_PIf / 2 + uv.y * M_PIf;
+        float x = tanf(alpha), y = tanf(beta);
+
+        DCHECK(!IsInf(x) && !IsInf(y));
+
+        Vector3f w = normalize(Vector3f(x, y, 1));
+        
+        if (duvdw)
+            *duvdw = pbrt::Sqr(M_PIf) * (1 - pbrt::Sqr(w[0])) * (1 - pbrt::Sqr(w[1])) / w[2];
+        
+        Vector3f dir {};
+        dir = dir + X * w[0];
+        dir = dir + Y * w[1];
+        dir = dir + Z * w[2];
+        return dir;
+    }
+
+    void sample(LightSampleRecord& lsr, const Vector3f& pos, float2 uu, float3& color) {
+        Bounds2f bds; // uv bounds
+        auto valid = ImageBounds(pos, bds);
+        if (!valid) return;
+
+        auto bIntegral = sat.Integral(bds);
+        if( bIntegral <= 0 ) return;
+
+        auto Px = [&](float x) -> float {
+            Bounds2f bx = bds;
+            bx.pMax[0] = x;
+            return sat.Integral(bx) / bIntegral;
+        };
+
+        float2 uv;
+        uv.x = SampleBisection(Px, uu.x, bds.pMin[0], bds.pMax[0], image.XSize());
+
+        uint nx = image.XSize();
+        Bounds2f bCond { 
+            { floor(uv.x * nx)/nx, bds.pMin[1] },
+            { ceil (uv.x * nx)/nx, bds.pMax[1] } };
+
+        if (bCond.pMin[0] == bCond.pMax[0])
+            bCond.pMax[0] += 1.0f / nx;
+
+        float condIntegral = sat.Integral(bCond);
+        if (condIntegral == 0) 
+            return;
+
+        auto Py = [&](float y) -> float {
+            Bounds2f by = bCond;
+            by.pMax[1] = y;
+            return sat.Integral(by) / condIntegral;
+        };
+        uv.y = SampleBisection(Py, uu.y, bds.pMin[1], bds.pMax[1], image.YSize());
+        //uv = clamp(uv, 0.0, 1.0);
+
+        float duvdw;
+        auto tmp = uv_direction(uv, &duvdw);
+
+        lsr.dir = reinterpret_cast<float3&>(tmp);
+        lsr.dist = FLT_MAX;
+        lsr.uv = uv;
+
+        // Compute PDF and return point sampled from windowed function
+        lsr.PDF = Eval(uv) / bIntegral;
+        lsr.PDF /= duvdw;
+        if(!isfinite(lsr.PDF)) {
+            lsr.PDF = 0.0;
+            return;
+        }
+
+        color = image.lookUV(uv);
+    }
+
+    float PDF(Vector3f p, Vector3f w) {
+        float duvdw;
+        auto uv = direction_uv(w, &duvdw);
+
+        Bounds2f bds;
+        bool valid = ImageBounds(p, bds);
+        if (!valid) return 0.0f;
+
+        float integ = sat.Integral(bds);
+        if (integ == 0) return 0.0f;
+
+        return Eval(reinterpret_cast<float2&>(uv)) / duvdw;
+    }
+};
+
+struct PortalLightList {
+
+#ifndef __CUDACC_RTC__
+    std::vector<PortalLight> list;
+    xinxinoptix::raii<CUdeviceptr> buffer;
+
+    std::vector<float> pdf;
+    std::vector<float> cdf;
+    xinxinoptix::raii<CUdeviceptr> pdf_buffer;
+    xinxinoptix::raii<CUdeviceptr> cdf_buffer;
+
+    xinxinoptix::raii<CUdeviceptr> dummy_buffer;
+#else
+    PortalLight *list;
+    size_t count;
+
+    float* pdf;
+    float* cdf;
+#endif
+
+    inline size_t COUNT() {
+    #ifndef __CUDACC_RTC__
+        return list.size();
+    #else
+        return count;
+    #endif
+    }
+
+#ifndef __CUDACC_RTC__
+    auto upload() {
+
+        if (list.size() == 0) {
+            *this = {};
+            return 0llu;
+        }
+
+        auto first = list.front().pack();
+        std::vector<typeof(first)> tmp;
+        tmp.reserve(list.size());
+        tmp.push_back(first);
+
+        pdf.clear(); pdf.resize(list.size());
+        cdf.clear(); cdf.resize(list.size());
+
+        auto power = list.front().phi();
+        pdf[0] = power;
+        cdf[0] = power;
+
+        for (size_t i=1; i<list.size(); ++i) {
+            auto ele = list[i].pack();
+            tmp.push_back(ele);
+
+            auto phi = list[i].phi();
+
+            power += phi;
+            pdf[i] = phi;
+            cdf[i] = power;
+        }
+
+        for (size_t i=0; i<list.size(); ++i) {
+            pdf[i] /= power;
+            cdf[i] /= power;
+        }
+        
+        size_t byte_size = sizeof(first) * list.size();
+
+        buffer.resize(byte_size);
+        cudaMemcpy((void*)buffer.handle, tmp.data(), byte_size, cudaMemcpyHostToDevice);
+
+        byte_size = sizeof(float) * list.size();
+        pdf_buffer.resize(byte_size);
+        cudaMemcpy((void*)pdf_buffer.handle, pdf.data(), byte_size, cudaMemcpyHostToDevice);
+
+        cdf_buffer.resize(byte_size);
+        cudaMemcpy((void*)cdf_buffer.handle, cdf.data(), byte_size, cudaMemcpyHostToDevice);
+
+        struct Dummy {
+            void* data; 
+            size_t count;
+            void* pdf; void* cdf;
+        };
+
+        auto result =  Dummy { (void*)buffer.handle, list.size(), 
+                                       (void*)pdf_buffer.handle, (void*)cdf_buffer.handle };
+
+        dummy_buffer.resize(sizeof(result));
+        cudaMemcpy((void*)dummy_buffer.handle, &result, sizeof(result), cudaMemcpyHostToDevice);
+
+        return dummy_buffer.handle;
+    }
+#endif
+
+};
+
+struct DistantLightList {
+
+#ifndef __CUDACC_RTC__
+    std::vector<zeno::DistantLightData> list;
+    std::vector<float> cdf;
+
+    xinxinoptix::raii<CUdeviceptr> data_buffer;
+    xinxinoptix::raii<CUdeviceptr> cdf_buffer;
+
+    xinxinoptix::raii<CUdeviceptr> dummy_buffer;
+#else
+    zeno::DistantLightData* list;
+    float* cdf;
+    uint count;
+#endif
+    inline size_t COUNT() {
+#ifndef __CUDACC_RTC__
+        return list.size();
+#else   
+        return count;
+#endif
+    }
+
+#ifndef __CUDACC_RTC__
+
+    auto upload() {
+
+        size_t byte_size  = sizeof(zeno::DistantLightData) * list.size();
+        data_buffer.resize( byte_size );
+        cudaMemcpy((void*)data_buffer.handle, list.data(), byte_size, cudaMemcpyHostToDevice);
+
+        byte_size = sizeof(float) * cdf.size();
+        cdf_buffer.resize(byte_size);
+        cudaMemcpy((void*)cdf_buffer.handle, cdf.data(), byte_size, cudaMemcpyHostToDevice);
+
+        struct Dummy {
+            void* data;
+            void* cdf;
+            size_t count;
+        };
+
+        Dummy dummy {
+            (void*)data_buffer.handle,
+            (void*)cdf_buffer.handle,
+            list.size()
+        };
+
+        dummy_buffer.resize(sizeof(dummy));
+        cudaMemcpy((void*)dummy_buffer.handle, &dummy, sizeof(dummy), cudaMemcpyHostToDevice);
+
+        return dummy_buffer.handle;     
+    }
+
+#endif
+};
\ No newline at end of file
diff --git a/zenovis/xinxinoptix/Sampling.h b/zenovis/xinxinoptix/Sampling.h
index dbffe9db65..289adb7078 100644
--- a/zenovis/xinxinoptix/Sampling.h
+++ b/zenovis/xinxinoptix/Sampling.h
@@ -7,13 +7,19 @@
 
 #ifdef __CUDACC_RTC__
     #include "zxxglslvec.h"
+    using Vector2f = vec2;
     using Vector3f = vec3;
 #else 
     #include "Host.h"
     #include <zeno/utils/vec.h>
+    using Vector2f = zeno::vec<2, float>;
     using Vector3f = zeno::vec<3, float>;
 #endif
 
+#ifndef FLT_MAX 
+#define FLT_MAX __FLT_MAX__
+#endif
+
 #ifdef __CUDACC_DEBUG__
     #define DCHECK assert
 #else
@@ -24,6 +30,22 @@
 
 #endif
 
+struct LightSampleRecord {
+    float3 p;
+    float PDF;
+
+    float3 n;
+    float NoL;
+
+    float3 dir;
+    float dist;
+
+    float2 uv;
+
+    float intensity = 1.0f;
+    bool isDelta = false;
+};
+
 namespace pbrt {
     
 template <typename T>
@@ -229,6 +251,7 @@ static __host__ __device__ __inline__ float3 sphereUV(const float3 &dir, bool in
     return float3 {u, v, 0.0f};
 } 
 
+
 static __host__ __device__ __inline__ float3 interp(float2 barys, float3 a, float3 b, float3 c)
 {
     float w0 = 1 - barys.x - barys.y;
diff --git a/zenovis/xinxinoptix/Shape.h b/zenovis/xinxinoptix/Shape.h
index 16d9668009..249f943218 100644
--- a/zenovis/xinxinoptix/Shape.h
+++ b/zenovis/xinxinoptix/Shape.h
@@ -9,22 +9,6 @@
 #include "Host.h"
 #endif
 
-struct LightSampleRecord {
-    float3 p;
-    float PDF;
-
-    float3 n;
-    float NoL;
-
-    float3 dir;
-    float dist;
-
-    float2 uv;
-
-    float intensity = 1.0f;
-    bool isDelta = false;
-};
-
 static constexpr float MinSphericalSampleArea = 3e-4f;
 static constexpr float MaxSphericalSampleArea = 6.22f;
 
diff --git a/zenovis/xinxinoptix/optixPathTracer.cpp b/zenovis/xinxinoptix/optixPathTracer.cpp
index d4251bdb72..505b340f5a 100644
--- a/zenovis/xinxinoptix/optixPathTracer.cpp
+++ b/zenovis/xinxinoptix/optixPathTracer.cpp
@@ -36,6 +36,7 @@
 #include <zeno/para/parallel_sort.h>
 #include <zeno/para/parallel_scan.h>
 #include <zeno/utils/log.h>
+#include <zeno/utils/pfm.h>
 #include <zeno/utils/zeno_p.h>
 #include <zeno/utils/fileio.h>
 #include <zeno/types/MaterialObject.h>
@@ -81,6 +82,7 @@
 
 #include "LightBounds.h"
 #include "LightTree.h"
+#include "Portal.h"
 
 #include "ChiefDesignerEXR.h"
 using namespace zeno::ChiefDesignerEXR;
@@ -269,9 +271,6 @@ struct PathTracerState
     OptixTraversableHandle         meshHandleIAS;
     raii<CUdeviceptr>              meshBufferIAS;
 
-    OptixTraversableHandle         gas_handle               = {};  // Traversable handle for triangle AS
-    raii<CUdeviceptr>              d_gas_output_buffer;  // Triangle AS memory
-
     raii<CUdeviceptr>              d_vertices;
     raii<CUdeviceptr>              d_clr;
     raii<CUdeviceptr>              d_nrm;
@@ -290,17 +289,12 @@ struct PathTracerState
     raii<CUdeviceptr>              d_uniforms;
 
     raii<OptixModule>              ptx_module;
-    raii<OptixModule>              ptx_module2;
     OptixPipelineCompileOptions    pipeline_compile_options;
     OptixPipeline                  pipeline;
 
     OptixProgramGroup              raygen_prog_group;
     OptixProgramGroup              radiance_miss_group;
     OptixProgramGroup              occlusion_miss_group;
-    OptixProgramGroup              radiance_hit_group;
-    OptixProgramGroup              occlusion_hit_group;
-    OptixProgramGroup              radiance_hit_group2;
-    OptixProgramGroup              occlusion_hit_group2;
 
     raii<CUstream>                       stream;
     raii<CUdeviceptr> accum_buffer_p;
@@ -311,7 +305,14 @@ struct PathTracerState
     raii<CUdeviceptr> accum_buffer_s;
     raii<CUdeviceptr> accum_buffer_t;
     raii<CUdeviceptr> accum_buffer_b;
-    raii<CUdeviceptr> lightsbuf_p;
+
+    raii<CUdeviceptr> finite_lights_ptr;
+
+    PortalLightList  plights;
+    DistantLightList dlights;
+
+    //std::vector<Portal> portals; 
+    
     raii<CUdeviceptr> sky_cdf_p;
     raii<CUdeviceptr> sky_start;
     Params                         params;
@@ -1301,55 +1302,6 @@ void updateRootIAS()
 	state.params.handle = state.rootHandleIAS;
 }
 
-static void buildMeshAccel( PathTracerState& state )
-{
-    //
-    // copy mesh data to device
-    //
-    const size_t vertices_size_in_bytes = g_vertices.size() * sizeof( Vertex );
-    CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &state.d_vertices.reset() ), vertices_size_in_bytes ) );
-    CUDA_CHECK( cudaMemcpy(
-                reinterpret_cast<void*>( (CUdeviceptr&)state.d_vertices ),
-                g_vertices.data(), vertices_size_in_bytes,
-                cudaMemcpyHostToDevice
-                ) );
-
-    const size_t mat_indices_size_in_bytes = g_mat_indices.size() * sizeof( uint32_t );
-    CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &state.d_mat_indices.reset() ), mat_indices_size_in_bytes ) );
-    CUDA_CHECK( cudaMemcpy(
-                reinterpret_cast<void*>( (CUdeviceptr)state.d_mat_indices ),
-                g_mat_indices.data(),
-                mat_indices_size_in_bytes,
-                cudaMemcpyHostToDevice
-                ) );
-
-    // // Build triangle GAS // // One per SBT record for this build input
-    std::vector<uint32_t> triangle_input_flags(//MAT_COUNT
-        g_mtlidlut.size(),
-        OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL);
-
-    OptixBuildInput triangle_input                           = {};
-    triangle_input.type                                      = OPTIX_BUILD_INPUT_TYPE_TRIANGLES;
-    triangle_input.triangleArray.vertexFormat                = OPTIX_VERTEX_FORMAT_FLOAT3;
-    triangle_input.triangleArray.vertexStrideInBytes         = sizeof( Vertex );
-    triangle_input.triangleArray.numVertices                 = static_cast<uint32_t>( g_vertices.size() );
-    triangle_input.triangleArray.vertexBuffers               = g_vertices.empty() ? nullptr : & state.d_vertices;
-    triangle_input.triangleArray.flags                       = triangle_input_flags.data();
-    triangle_input.triangleArray.numSbtRecords               = g_vertices.empty() ? 1 : g_mtlidlut.size();
-    triangle_input.triangleArray.sbtIndexOffsetBuffer        = state.d_mat_indices;
-    triangle_input.triangleArray.sbtIndexOffsetSizeInBytes   = sizeof( uint32_t );
-    triangle_input.triangleArray.sbtIndexOffsetStrideInBytes = sizeof( uint32_t );
-
-    OptixAccelBuildOptions accel_options = {};
-    accel_options.buildFlags             = OPTIX_BUILD_FLAG_ALLOW_COMPACTION | OPTIX_BUILD_FLAG_ALLOW_RANDOM_VERTEX_ACCESS | OPTIX_BUILD_FLAG_ALLOW_RANDOM_INSTANCE_ACCESS;
-    accel_options.operation              = OPTIX_BUILD_OPERATION_BUILD;
-
-    buildXAS(state.context, accel_options, triangle_input, state.d_gas_output_buffer, state.d_gas_output_buffer);
-
-    state.d_vertices.reset();
-    state.d_mat_indices.reset();
-}
-
 static void createSBT( PathTracerState& state )
 {
         state.d_raygen_record.reset();
@@ -1357,7 +1309,6 @@ static void createSBT( PathTracerState& state )
         state.d_hitgroup_records.reset();
         state.d_callable_records.reset();
 
-        state.d_gas_output_buffer.reset();
         state.accum_buffer_p.reset();
         state.albedo_buffer_p.reset();
         state.normal_buffer_p.reset();
@@ -1742,7 +1693,10 @@ void optixinit( int argc, char* argv[] )
     auto cur_path = std::string(_pgmptr);
     cur_path = cur_path.substr(0, cur_path.find_last_of("\\"));
 #endif
-    OptixUtil::sky_tex = cur_path + "/hdr/Panorama.hdr";
+
+    OptixUtil::default_sky_tex = cur_path + "/hdr/Panorama.hdr";
+    OptixUtil::sky_tex = OptixUtil::default_sky_tex;
+
     OptixUtil::addTexture(OptixUtil::sky_tex.value());
     xinxinoptix::update_hdr_sky(0, {0, 0, 0}, 0.8);
 }
@@ -2105,6 +2059,14 @@ void unload_light(){
     triangleLightCoords.clear();
     triangleLightNormals.clear();
 
+    state.dlights = {};
+    state.plights = {};
+
+    state.params.dlights_ptr = 0llu;
+    state.params.plights_ptr = 0llu;
+
+    OptixUtil::portal_delayed.reset();
+
     std::cout << "Lights unload done. \n"<< std::endl;
 }
 
@@ -2158,6 +2120,62 @@ void show_background(bool enable) {
     state.params.show_background = enable;
 }
 
+void updatePortalLights(const std::vector<Portal>& portals) {
+
+    auto &tex = OptixUtil::g_tex[OptixUtil::sky_tex.value()];
+
+    auto& pll = state.plights;
+    auto& pls = pll.list;
+    pls.clear();
+    pls.reserve(max(portals.size(), 0) );
+
+    glm::mat4 rotation = glm::mat4(1.0f);
+    rotation = glm::rotate(rotation, glm::radians(state.params.sky_rot_y), glm::vec3(0,1,0));
+    rotation = glm::rotate(rotation, glm::radians(state.params.sky_rot_x), glm::vec3(1,0,0));
+    rotation = glm::rotate(rotation, glm::radians(state.params.sky_rot_z), glm::vec3(0,0,1));
+    rotation = glm::rotate(rotation, glm::radians(state.params.sky_rot), glm::vec3(0,1,0));
+    
+    glm::mat4* rotation_ptr = nullptr;
+    if ( glm::mat4(1.0f) != rotation ) {
+        rotation_ptr = &rotation;
+    }
+
+    for (auto& portal : portals) {
+        auto pl = PortalLight(portal, (float3*)tex->rawData.data(), tex->width, tex->height, rotation_ptr);
+        pls.push_back(std::move(pl));
+    }
+
+    state.params.plights_ptr = (void*)pll.upload();
+}
+
+void updateDistantLights(std::vector<zeno::DistantLightData>& dldl) 
+{
+    if (dldl.empty()) {
+        state.dlights = {};
+        state.params.dlights_ptr = 0u;
+        return;
+    }
+
+    float power = 0.0f;
+
+    std::vector<float> cdf; cdf.reserve(dldl.size());
+
+    for (auto& dld : dldl) {
+        auto ppp = dld.color * dld.intensity;
+        power += (ppp[0] + ppp[1] + ppp[2]) / 3.0f;
+        cdf.push_back(power);
+    }
+
+    for(auto& c : cdf) {
+        c /= power;
+    }
+
+    state.dlights.list = dldl;
+    state.dlights.cdf = cdf;
+
+    state.params.dlights_ptr = (void*)state.dlights.upload();
+}
+
 void update_procedural_sky(
     zeno::vec2f sunLightDir,
     float sunLightSoftness,
@@ -2358,7 +2376,7 @@ static void buildLightTrianglesGAS( PathTracerState& state, std::vector<float3>&
 
 void buildLightTree() {
     camera_changed = true;
-    state.lightsbuf_p.reset();
+    state.finite_lights_ptr.reset();
 
     state.params.lightTreeSampler = 0llu;
     state.params.triangleLightCoordsBuffer = 0llu;
@@ -2544,13 +2562,13 @@ void buildLightTree() {
     buildLightTrianglesGAS(state, lightsWrapper._triangleLightGeo, lightsWrapper.lightTrianglesGasBuffer, lightsWrapper.lightTrianglesGas);
 
     CUDA_CHECK( cudaMalloc(
-        reinterpret_cast<void**>( &state.lightsbuf_p.reset() ),
+        reinterpret_cast<void**>( &state.finite_lights_ptr.reset() ),
         sizeof( GenericLight ) * std::max(lightsWrapper.g_lights.size(),(size_t)1)
         ) );
 
-    state.params.lights = (GenericLight*)(CUdeviceptr)state.lightsbuf_p;
+    state.params.lights = (GenericLight*)(CUdeviceptr)state.finite_lights_ptr;
     CUDA_CHECK( cudaMemcpy(
-                reinterpret_cast<void*>( (CUdeviceptr)state.lightsbuf_p ),
+                reinterpret_cast<void*>( (CUdeviceptr)state.finite_lights_ptr ),
                 lightsWrapper.g_lights.data(), sizeof( GenericLight ) * lightsWrapper.g_lights.size(),
                 cudaMemcpyHostToDevice
                 ) );
@@ -2729,28 +2747,39 @@ OptixUtil::_compile_group.wait();
     theTimer.tock("Done Optix Shader Compile:");
 
     if (OptixUtil::sky_tex.has_value()) {
-        state.params.sky_texture = OptixUtil::g_tex[OptixUtil::sky_tex.value()]->texture;
-        state.params.skynx = OptixUtil::sky_nx_map[OptixUtil::sky_tex.value()];
-        state.params.skyny = OptixUtil::sky_ny_map[OptixUtil::sky_tex.value()];
-        state.params.envavg = OptixUtil::sky_avg_map[OptixUtil::sky_tex.value()];
+
+        auto &tex = OptixUtil::g_tex[OptixUtil::sky_tex.value()];
+        if (tex.get() == 0) {
+            tex = OptixUtil::g_tex[OptixUtil::default_sky_tex];
+        }
+        
+        if (tex->texture == state.params.sky_texture) return;
+
+        state.params.sky_texture = tex->texture;
+        state.params.skynx = tex->width;
+        state.params.skyny = tex->height;
+        state.params.envavg = tex->average;
+
         CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &state.sky_cdf_p.reset() ),
-                              sizeof(float2)*OptixUtil::sky_cdf_map[OptixUtil::sky_tex.value()].size() ) );
+                            sizeof(float2)*tex->cdf.size() ) );
         CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &state.sky_start.reset() ),
-                              sizeof(int)*OptixUtil::sky_start_map[OptixUtil::sky_tex.value()].size() ) );
+                              sizeof(int)*tex->start.size() ) );
+
         cudaMemcpy(reinterpret_cast<char *>((CUdeviceptr)state.sky_cdf_p),
-                   OptixUtil::sky_cdf_map[OptixUtil::sky_tex.value()].data(),
-                   sizeof(float)*OptixUtil::sky_cdf_map[OptixUtil::sky_tex.value()].size(),
+                   tex->cdf.data(),
+                   sizeof(float)*tex->cdf.size(),
                    cudaMemcpyHostToDevice);
-        cudaMemcpy(reinterpret_cast<char *>((CUdeviceptr)state.sky_cdf_p)+sizeof(float)*OptixUtil::sky_cdf_map[OptixUtil::sky_tex.value()].size(),
-                   OptixUtil::sky_pdf_map[OptixUtil::sky_tex.value()].data(),
-                   sizeof(float)*OptixUtil::sky_pdf_map[OptixUtil::sky_tex.value()].size(),
+        cudaMemcpy(reinterpret_cast<char *>((CUdeviceptr)state.sky_cdf_p) + sizeof(float)*tex->cdf.size(),
+                   tex->pdf.data(),
+                   sizeof(float)*tex->pdf.size(),
                    cudaMemcpyHostToDevice);
         cudaMemcpy(reinterpret_cast<char *>((CUdeviceptr)state.sky_start),
-                   OptixUtil::sky_start_map[OptixUtil::sky_tex.value()].data(),
-                   sizeof(int)*OptixUtil::sky_start_map[OptixUtil::sky_tex.value()].size(),
+                   tex->start.data(),
+                   sizeof(int)*tex->start.size(),
                    cudaMemcpyHostToDevice);
+
         state.params.skycdf = reinterpret_cast<float *>((CUdeviceptr)state.sky_cdf_p);
-        state.params.sky_start = reinterpret_cast<int *>((CUdeviceptr)state.sky_start);
+        state.params.sky_start = reinterpret_cast<int*>((CUdeviceptr)state.sky_start);
 
     } else {
         state.params.skynx = 0;
@@ -2764,19 +2793,6 @@ void optixupdateend() {
         OptixUtil::createPipeline();
 
     printf("Pipeline created \n");
-        //static bool hadOnce = false;
-        //if (hadOnce) {
-    //OPTIX_CHECK( optixPipelineDestroy( state.pipeline ) );
-    //state.raygen_prog_group ) );
-    //state.radiance_miss_group ) );
-    //state.occlusion_miss_group ) );
-    //OPTIX_CHECK( optixProgramGroupDestroy( state.radiance_hit_group ) );
-    //OPTIX_CHECK( optixProgramGroupDestroy( state.occlusion_hit_group ) );
-    //OPTIX_CHECK( optixProgramGroupDestroy( state.radiance_hit_group2 ) );
-    //OPTIX_CHECK( optixProgramGroupDestroy( state.occlusion_hit_group2 ) );
-    //OPTIX_CHECK( optixModuleDestroy( state.ptx_module ) );
-    //OPTIX_CHECK( optixDeviceContextDestroy( state.context ) );
-        //} hadOnce = true;
 
         state.pipeline_compile_options = OptixUtil::pipeline_compile_options;
         state.pipeline = OptixUtil::pipeline;
@@ -3767,13 +3783,6 @@ void set_perspective_by_focal_length(float const *U, float const *V, float const
 void set_outside_random_number(int32_t outside_random_number) {
     state.params.outside_random_number = outside_random_number;
 }
-static void write_pfm(std::string& path, int w, int h, const float *rgb) {
-    std::string header = zeno::format("PF\n{} {}\n-1.0\n", w, h);
-    std::vector<char> data(header.size() + w * h * sizeof(zeno::vec3f));
-    memcpy(data.data(), header.data(), header.size());
-    memcpy(data.data() + header.size(), rgb, w * h * sizeof(zeno::vec3f));
-    zeno::file_put_binary(data, path);
-}
 
 void *optixgetimg_extra(std::string name) {
     if (name == "diffuse") {
@@ -3897,13 +3906,13 @@ void optixrender(int fbo, int samples, bool denoise, bool simpleRender) {
                     //SaveEXR(_albedo_buffer, w, h, 4, 0, (path+".albedo.exr").c_str(), nullptr);
                     auto a_path = path + ".albedo.pfm";
                     std::string native_a_path = zeno::create_directories_when_write_file(a_path);
-                    write_pfm(native_a_path, w, h, _albedo_buffer);
+                    zeno::write_pfm(native_a_path.c_str(), w, h, _albedo_buffer);
 
                     const float* _normal_buffer = reinterpret_cast<float*>(state.normal_buffer_p.handle);
                     //SaveEXR(_normal_buffer, w, h, 4, 0, (path+".normal.exr").c_str(), nullptr);
                     auto n_path = path + ".normal.pfm";
                     std::string native_n_path = zeno::create_directories_when_write_file(n_path);
-                    write_pfm(native_n_path, w, h, _normal_buffer);
+                    zeno::write_pfm(native_n_path.c_str(), w, h, _normal_buffer);
                 }
             }
         }
@@ -3930,7 +3939,29 @@ void *optixgetimg(int &w, int &h) {
     //sutil::saveImage( outfile, buffer, false );
 //}
 
-void optixcleanup() {
+void optixCleanup() {
+
+    state.dlights = {};
+    state.params.dlights_ptr = 0u;
+
+    state.plights = {};
+    state.params.plights_ptr = 0u;
+
+    lightsWrapper.reset();
+    state.finite_lights_ptr.reset();
+    
+    state.params.sky_strength = 1.0f;
+    state.params.sky_texture;
+
+    auto sky_path = OptixUtil::default_sky_tex;
+    auto sky_tex = OptixUtil::g_tex[sky_path];
+
+    OptixUtil::g_tex = { {sky_path, sky_tex} };
+    //OptixUtil::g_tex.at(sky_path) = sky_tex;
+    OptixUtil::sky_tex = OptixUtil::default_sky_tex;
+}
+
+void optixDestroy() {
     using namespace OptixUtil;
     try {
         CUDA_SYNC_CHECK();
diff --git a/zenovis/xinxinoptix/optixPathTracer.h b/zenovis/xinxinoptix/optixPathTracer.h
index d26f167511..7df84d652e 100644
--- a/zenovis/xinxinoptix/optixPathTracer.h
+++ b/zenovis/xinxinoptix/optixPathTracer.h
@@ -196,6 +196,9 @@ struct Params
     uint32_t firstSoloSphereOffset;
     void* sphereInstAuxLutBuffer;
 
+    void*    dlights_ptr;
+    void*    plights_ptr;
+
     float skyLightProbablity() {
 
         if (sky_strength <= 0.0f)
diff --git a/zenovis/xinxinoptix/proceduralSky.h b/zenovis/xinxinoptix/proceduralSky.h
index fd8e4d01e7..ad95192d29 100644
--- a/zenovis/xinxinoptix/proceduralSky.h
+++ b/zenovis/xinxinoptix/proceduralSky.h
@@ -292,30 +292,6 @@ static __inline__ __device__ vec3 proceduralSky(
     return col;
 }
 
-static __inline__ __device__ vec3 hdrSky2(
-    vec3 dir
-){
-  dir = dir
-            .rotY(to_radians(params.sky_rot_y))
-            .rotX(to_radians(params.sky_rot_x))
-            .rotZ(to_radians(params.sky_rot_z))
-            .rotY(to_radians(params.sky_rot));
-            
-  vec3 uv = sphereUV(dir, true);
-  vec3 col = vec3(0);
-  for(int jj=-2;jj<=2;jj++)
-  {
-    for(int ii=-2;ii<=2;ii++)
-    {
-      float dx = (float)ii / (float)(params.skynx);
-      float dy = (float)jj / (float)(params.skyny);
-      col = col + (vec3)texture2D(params.sky_texture, vec2(uv[0] + dx, uv[1] + dy)) * params.sky_strength;
-    }
-  }
-
-  return col/9.0f;
-}
-
 static __inline__ __device__ vec3 hdrSky(
         vec3 dir, float upperBound,  float isclamp, float &pdf
 ){
@@ -363,10 +339,7 @@ static __inline__ __device__ vec3 colorTemperatureToRGB(float temperatureInKelvi
 
     return retColor;
 }
-static __inline__ __device__ vec3 envSky2(vec3 dir)
-{
-  return hdrSky2(dir);
-}
+
 static __inline__ __device__ vec3 envSky(
     vec3 dir,
     vec3 sunLightDir,
diff --git a/zenovis/xinxinoptix/xinxinoptixapi.h b/zenovis/xinxinoptix/xinxinoptixapi.h
index 826145e270..6194369c91 100644
--- a/zenovis/xinxinoptix/xinxinoptixapi.h
+++ b/zenovis/xinxinoptix/xinxinoptixapi.h
@@ -8,6 +8,9 @@
 
 #include "optixSphere.h"
 #include "zeno/utils/vec.h"
+#include "zeno/types/LightObject.h"
+
+#include "Portal.h"
 
 enum ShaderMaker {
     Mesh = 0,
@@ -30,7 +33,9 @@ namespace xinxinoptix {
 
 std::set<std::string> uniqueMatsForMesh();
 
-void optixcleanup();
+void optixCleanup();
+
+void optixDestroy();
 void optixrender(int fbo = 0, int samples = 1, bool denoise = false, bool simpleRender = false);
 void *optixgetimg(int &w, int &h);
 void optixinit(int argc, char* argv[]);
@@ -102,6 +107,9 @@ void update_procedural_sky(zeno::vec2f sunLightDir, float sunLightSoftness, zeno
 void update_hdr_sky(float sky_rot, zeno::vec3f sky_rot3d, float sky_strength);
 void using_hdr_sky(bool enable);
 void show_background(bool enable);
+
+void updatePortalLights(const std::vector<Portal>& portals);
+void updateDistantLights(std::vector<zeno::DistantLightData>& dldl);
 // void optixUpdateUniforms(std::vector<float4> & inConstants);
 void optixUpdateUniforms(void *inConstants, std::size_t size);
 }
diff --git a/zenovis/xinxinoptix/zxxglslvec.h b/zenovis/xinxinoptix/zxxglslvec.h
index 28f59ac447..699081020f 100644
--- a/zenovis/xinxinoptix/zxxglslvec.h
+++ b/zenovis/xinxinoptix/zxxglslvec.h
@@ -92,6 +92,13 @@ struct vec3{
 
 struct vec2{
     float x, y;
+
+    __forceinline__ __device__ float& operator[](unsigned int index) {
+        auto ptr= &this->x;
+        ptr += index;
+        return *ptr;
+    }
+    
     __forceinline__ __device__ vec2(const float2 &_v)
     {
         x = _v.x;