Skip to content

Commit d27461b

Browse files
wangkepfewangkepfe
committed
wip
1 parent bf0de2c commit d27461b

File tree

10 files changed

+166
-837
lines changed

10 files changed

+166
-837
lines changed

‎data/perf/performance_report.txt‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -139,3 +139,11 @@
139139
2025-07-30 17:20:37.549 | 64 | 3840x2160 | 50.09 | 6.77 | 0.18 | 0.48 | 23.77 | 0.09 | 23.94 | default run
140140
2025-07-30 23:53:10.988 | 64 | 3840x2160 | 43.72 | 6.67 | 0.20 | 0.53 | 23.03 | 0.10 | 18.16 | default run
141141
2025-07-31 18:40:44.927 | 64 | 3840x2160 | 42.86 | 6.64 | 0.20 | 0.47 | 22.64 | 0.08 | 17.84 | default run
142+
2025-08-06 14:51:03.266 | 64 | 3840x2160 | 29.08 | 7.61 | 0.24 | 0.49 | 7.59 | 0.09 | 19.02 | default run
143+
2025-08-06 14:53:04.882 | 64 | 3840x2160 | 26.35 | 7.68 | 0.22 | 0.49 | 0.25 | 0.09 | 23.55 | default run
144+
2025-08-06 14:53:47.715 | 64 | 3840x2160 | 29.22 | 7.25 | 0.20 | 0.50 | 7.69 | 0.08 | 19.11 | default run
145+
2025-08-06 15:37:41.198 | 64 | 3840x2160 | 7.93 | 7.64 | 0.22 | 0.51 | 1.56 | 0.09 | 3.88 | default run
146+
2025-08-06 15:38:21.300 | 64 | 3840x2160 | 8.18 | 9.24 | 0.29 | 0.49 | 1.56 | 0.09 | 3.91 | default run
147+
2025-08-06 15:38:43.664 | 64 | 3840x2160 | 7.89 | 7.48 | 0.23 | 0.46 | 1.58 | 0.08 | 3.92 | default run
148+
2025-08-06 15:39:07.172 | 64 | 3840x2160 | 8.04 | 8.20 | 0.23 | 0.47 | 1.56 | 0.08 | 3.88 | default run
149+
2025-08-06 15:42:13.133 | 64 | 3840x2160 | 8.16 | 7.64 | 0.23 | 0.51 | 1.61 | 0.08 | 4.05 | default run

‎renderer/core/OptixRenderer.cpp‎

Lines changed: 59 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -187,6 +187,7 @@ void OptixRenderer::render()
187187

188188
// Handle light-related parameters safely to prevent null pointer access
189189
m_systemParameter.lights = scene.m_lights;
190+
m_systemParameter.numLights = scene.m_numLights;
190191
m_systemParameter.lightAliasTable = scene.d_lightAliasTable;
191192
m_systemParameter.instanceLightMapping = scene.d_instanceLightMapping;
192193
m_systemParameter.numInstancedLightMesh = scene.numInstancedLightMesh;
@@ -198,13 +199,22 @@ void OptixRenderer::render()
198199

199200
CUDA_CHECK(cudaMemcpyAsync((void *)m_d_systemParameter, &m_systemParameter, sizeof(SystemParameter), cudaMemcpyHostToDevice, backend.getCudaStream()));
200201

202+
// Debug: Check for errors before OptixLaunch
203+
CUDA_CHECK(cudaPeekAtLastError());
204+
201205
// Only launch OptixLaunch if we have a valid top object (scene is not empty)
202206
if (m_systemParameter.topObject != 0) {
203207
OPTIX_CHECK(m_api.optixLaunch(m_pipeline, backend.getCudaStream(), (CUdeviceptr)m_d_systemParameter, sizeof(SystemParameter), &m_sbt, m_width, m_height, 1));
208+
209+
// Check immediately after OptixLaunch with device sync
210+
CUDA_CHECK(cudaDeviceSynchronize());
211+
cudaError_t err = cudaPeekAtLastError();
212+
if (err != cudaSuccess) {
213+
std::cout << "ERROR: OptixLaunch kernels failed with CUDA error: " << cudaGetErrorString(err) << std::endl;
214+
}
204215
}
205216

206217
CUDA_CHECK(cudaStreamSynchronize(backend.getCudaStream()));
207-
208218
CUDA_CHECK(cudaDeviceSynchronize());
209219
CUDA_CHECK(cudaPeekAtLastError());
210220

@@ -317,6 +327,10 @@ void OptixRenderer::updateAnimatedEntities(CUstream cudaStream, float currentTim
317327
unsigned int sbtIndex = entitySbtOffset + rayType;
318328
m_sbtRecordGeometryInstanceData[sbtIndex].data.indices = (Int3 *)geometry.indices;
319329
m_sbtRecordGeometryInstanceData[sbtIndex].data.attributes = (VertexAttributes *)geometry.attributes;
330+
331+
// CRITICAL FIX: Add bounds information to prevent illegal memory access
332+
m_sbtRecordGeometryInstanceData[sbtIndex].data.numIndices = geometry.numIndices / 3; // Convert to triangle count
333+
m_sbtRecordGeometryInstanceData[sbtIndex].data.numAttributes = geometry.numAttributes;
320334
}
321335
}
322336
}
@@ -592,6 +606,10 @@ void OptixRenderer::update()
592606
{
593607
m_sbtRecordGeometryInstanceData[sbtIndex + rayType].data.indices = (Int3 *)geometry.indices;
594608
m_sbtRecordGeometryInstanceData[sbtIndex + rayType].data.attributes = (VertexAttributes *)geometry.attributes;
609+
610+
// CRITICAL FIX: Add bounds information to prevent illegal memory access
611+
m_sbtRecordGeometryInstanceData[sbtIndex + rayType].data.numIndices = geometry.numIndices / 3; // Convert to triangle count
612+
m_sbtRecordGeometryInstanceData[sbtIndex + rayType].data.numAttributes = geometry.numAttributes;
595613
}
596614
}
597615
}
@@ -704,6 +722,10 @@ void OptixRenderer::update()
704722
{
705723
m_sbtRecordGeometryInstanceData[sbtIndex + rayType].data.indices = (Int3 *)geometry.indices;
706724
m_sbtRecordGeometryInstanceData[sbtIndex + rayType].data.attributes = (VertexAttributes *)geometry.attributes;
725+
726+
// CRITICAL FIX: Add bounds information to prevent illegal memory access
727+
m_sbtRecordGeometryInstanceData[sbtIndex + rayType].data.numIndices = geometry.numIndices / 3; // Convert to triangle count
728+
m_sbtRecordGeometryInstanceData[sbtIndex + rayType].data.numAttributes = geometry.numAttributes;
707729
}
708730
}
709731
}
@@ -1467,6 +1489,12 @@ void OptixRenderer::init()
14671489
{
14681490
unsigned int geometryIndex = chunkIndex * scene.uninstancedGeometryCount + objectId;
14691491

1492+
// Skip SBT creation if geometry has null pointers (empty geometry)
1493+
assert(geometryIndex < m_geometries.size());
1494+
if (m_geometries[geometryIndex].indices == nullptr || m_geometries[geometryIndex].attributes == nullptr) {
1495+
continue; // Skip this geometry, no SBT records created
1496+
}
1497+
14701498
for (unsigned int rayType = 0; rayType < numTypesOfRays; ++rayType)
14711499
{
14721500
if (rayType == 0)
@@ -1478,10 +1506,16 @@ void OptixRenderer::init()
14781506
memcpy(m_sbtRecordGeometryInstanceData[sbtRecordIndex].header, m_sbtRecordHitShadow.header, OPTIX_SBT_RECORD_HEADER_SIZE);
14791507
}
14801508

1481-
assert(geometryIndex < m_geometries.size());
14821509
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.indices = (Int3 *)m_geometries[geometryIndex].indices;
14831510
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.attributes = (VertexAttributes *)m_geometries[geometryIndex].attributes;
1484-
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.materialIndex = objectId;
1511+
1512+
// CRITICAL FIX: Add bounds information to prevent illegal memory access
1513+
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.numIndices = m_geometries[geometryIndex].numIndices / 3; // Convert to triangle count
1514+
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.numAttributes = m_geometries[geometryIndex].numAttributes;
1515+
1516+
// Map objectId to material index (objectId 1->16 maps to material index 0->15)
1517+
unsigned int materialIndex = (objectId >= 1) ? (objectId - 1) : 0;
1518+
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.materialIndex = materialIndex;
14851519
sbtRecordIndex++;
14861520
}
14871521
}
@@ -1495,6 +1529,12 @@ void OptixRenderer::init()
14951529
// Instanced geometries are stored after all chunk-based geometries
14961530
unsigned int geometryIndex = scene.numChunks * scene.uninstancedGeometryCount + (objectId - GetInstancedObjectIdBegin());
14971531

1532+
// Skip SBT creation if geometry has null pointers (empty geometry)
1533+
assert(geometryIndex < m_geometries.size());
1534+
if (m_geometries[geometryIndex].indices == nullptr || m_geometries[geometryIndex].attributes == nullptr) {
1535+
continue; // Skip this geometry, no SBT records created
1536+
}
1537+
14981538
for (unsigned int rayType = 0; rayType < numTypesOfRays; ++rayType)
14991539
{
15001540
if (rayType == 0)
@@ -1506,10 +1546,16 @@ void OptixRenderer::init()
15061546
memcpy(m_sbtRecordGeometryInstanceData[sbtRecordIndex].header, m_sbtRecordHitShadow.header, OPTIX_SBT_RECORD_HEADER_SIZE);
15071547
}
15081548

1509-
assert(geometryIndex < m_geometries.size());
15101549
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.indices = (Int3 *)m_geometries[geometryIndex].indices;
15111550
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.attributes = (VertexAttributes *)m_geometries[geometryIndex].attributes;
1512-
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.materialIndex = objectId;
1551+
1552+
// CRITICAL FIX: Add bounds information to prevent illegal memory access
1553+
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.numIndices = m_geometries[geometryIndex].numIndices / 3; // Convert to triangle count
1554+
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.numAttributes = m_geometries[geometryIndex].numAttributes;
1555+
1556+
// Map objectId to material index (objectId 1->16 maps to material index 0->15)
1557+
unsigned int materialIndex = (objectId >= 1) ? (objectId - 1) : 0;
1558+
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.materialIndex = materialIndex;
15131559
sbtRecordIndex++;
15141560
}
15151561
}
@@ -1538,6 +1584,11 @@ void OptixRenderer::init()
15381584
assert(geometryIndex < m_geometries.size());
15391585
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.indices = (Int3 *)m_geometries[geometryIndex].indices;
15401586
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.attributes = (VertexAttributes *)m_geometries[geometryIndex].attributes;
1587+
1588+
// CRITICAL FIX: Add bounds information to prevent illegal memory access
1589+
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.numIndices = m_geometries[geometryIndex].numIndices / 3; // Convert to triangle count
1590+
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.numAttributes = m_geometries[geometryIndex].numAttributes;
1591+
15411592
// Use the correct material index for entities from the material array
15421593
unsigned int entityMaterialIndex = m_entityMaterialStartIndex + static_cast<unsigned int>(entity->getType());
15431594
m_sbtRecordGeometryInstanceData[sbtRecordIndex].data.materialIndex = entityMaterialIndex;
@@ -1592,6 +1643,9 @@ void OptixRenderer::init()
15921643
{
15931644
CUDA_CHECK(cudaMalloc((void **)&m_systemParameter.materialParameters, sizeof(MaterialParameter) * m_materialParameters.size()));
15941645
CUDA_CHECK(cudaMemcpyAsync((void *)m_systemParameter.materialParameters, m_materialParameters.data(), sizeof(MaterialParameter) * m_materialParameters.size(), cudaMemcpyHostToDevice, Backend::Get().getCudaStream()));
1646+
1647+
// Set the number of material parameters for bounds checking
1648+
m_systemParameter.numMaterialParameters = (unsigned int)m_materialParameters.size();
15951649

15961650
CUDA_CHECK(cudaMalloc((void **)&m_d_systemParameter, sizeof(SystemParameter)));
15971651
CUDA_CHECK(cudaMemcpyAsync((void *)m_d_systemParameter, &m_systemParameter, sizeof(SystemParameter), cudaMemcpyHostToDevice, Backend::Get().getCudaStream()));

‎renderer/core/Scene.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -73,6 +73,7 @@ class Scene
7373
std::unordered_map<int, std::array<float, 12>> instanceTransformMatrices;
7474

7575
LightInfo *m_lights = nullptr;
76+
unsigned int m_numLights = 0; // Number of lights in m_lights array
7677
AliasTable lightAliasTable;
7778
AliasTable *d_lightAliasTable = nullptr;
7879
float accumulatedLocalLightLuminance = 0.0f;

‎renderer/shaders/AliasTable.h‎

Lines changed: 12 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,12 @@ class AliasTable
3333
// Device function to perform sampling.
3434
__device__ unsigned int sample(float u, float &pmf) const
3535
{
36+
// Handle empty alias table case
37+
if (len == 0 || bins == nullptr) {
38+
pmf = 0.0f;
39+
return 0;
40+
}
41+
3642
int offset = min(int(u * len), int(len - 1));
3743
float up = min(u * len - offset, 0.999999f);
3844

@@ -51,7 +57,12 @@ class AliasTable
5157

5258
// Host/device inline functions.
5359
__host__ __device__ unsigned int size() const { return len; }
54-
__device__ float PMF(unsigned int index) const { return bins[index].p; }
60+
__device__ float PMF(unsigned int index) const {
61+
if (len == 0 || bins == nullptr || index >= len) {
62+
return 0.0f;
63+
}
64+
return bins[index].p;
65+
}
5566

5667
private:
5768
AliasTableBin *bins = nullptr;

‎renderer/shaders/RayGen.cu‎

Lines changed: 12 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -56,9 +56,18 @@ __device__ __inline__ bool TraceNextPath(
5656
if (optixHitObjectIsHit())
5757
{
5858
const GeometryInstanceData *instanceData = reinterpret_cast<const GeometryInstanceData *>(optixHitObjectGetSbtDataPointer());
59-
const MaterialParameter &parameters = sysParam.materialParameters[instanceData->materialIndex];
60-
int materialId = parameters.materialId;
61-
hint = materialId;
59+
60+
// Defensive null pointer check
61+
if (instanceData && instanceData->indices && instanceData->attributes) {
62+
// Bounds check for material parameter access
63+
unsigned int materialIndex = instanceData->materialIndex;
64+
if (materialIndex >= sysParam.numMaterialParameters) {
65+
materialIndex = 0; // Default to first material if out of bounds
66+
}
67+
const MaterialParameter &parameters = sysParam.materialParameters[materialIndex];
68+
int materialId = parameters.materialId;
69+
hint = materialId;
70+
}
6271
}
6372
optixReorder(hint, NumOfBitsMaxBsdfIndex);
6473
optixInvoke(payload.x, payload.y);

‎renderer/shaders/Restir.h‎

Lines changed: 9 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -386,9 +386,15 @@ INL_DEVICE LightSample GetLightSampleFromReservoir(const DIReservoir &reservoir,
386386
}
387387
else
388388
{
389-
LightInfo lightInfo = sysParam.lights[lightIndex];
390-
TriangleLight triLight = TriangleLight::Create(lightInfo);
391-
ls = triLight.calcSample(uv, surface.pos);
389+
// Bounds check for light array access
390+
if (lightIndex >= sysParam.numLights) {
391+
ls = LightSample{}; // Return empty light sample
392+
}
393+
else {
394+
LightInfo lightInfo = sysParam.lights[lightIndex];
395+
TriangleLight triLight = TriangleLight::Create(lightInfo);
396+
ls = triLight.calcSample(uv, surface.pos);
397+
}
392398
}
393399
return ls;
394400
}

‎renderer/shaders/SystemParameter.h‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -68,6 +68,7 @@ struct SystemParameter
6868
SurfObj UIBuffer;
6969

7070
MaterialParameter *materialParameters;
71+
unsigned int numMaterialParameters; // Number of materials in materialParameters array
7172
InstanceLightMapping *instanceLightMapping;
7273
unsigned int numInstancedLightMesh;
7374

@@ -83,6 +84,7 @@ struct SystemParameter
8384

8485
LightInfo *lights;
8586
AliasTable *lightAliasTable;
87+
unsigned int numLights; // Number of lights in the lights array
8688
float accumulatedLocalLightLuminance;
8789

8890
Float3 *edgeToHighlight;
@@ -111,6 +113,9 @@ struct GeometryInstanceData
111113
{
112114
Int3 *indices;
113115
VertexAttributes *attributes;
116+
117+
unsigned int numIndices; // Number of indices (count of Int3, not individual ints)
118+
unsigned int numAttributes; // Number of vertex attributes
114119

115120
int materialIndex;
116121
};

0 commit comments

Comments
 (0)