Generated struct named gaussianParticle_RawParameters_0 mismatches its usage
#6
Assignees
Comments
|
Could you please :
This will give me the version of the slang compiler setup. |
|
I'm sorry, due to information security regulations, I am unable to upload files. However, I can confirm that the version of slangc is
struct gaussianParticle_RawParameters_0
{
float3 position_1;
float density_1;
float4 quaternion_0;
float3 scale_1;
float padding_0;
};patch diff --git a/threedgut_tracer/include/3dgut/kernels/cuda/models/shRadiativeGaussianParticles.cuh b/threedgut_tracer/include/3dgut/kernels/cuda/models/shRadiativeGaussianParticles.cuh
index 3d5c6a8..ab2a45a 100644
--- a/threedgut_tracer/include/3dgut/kernels/cuda/models/shRadiativeGaussianParticles.cuh
+++ b/threedgut_tracer/include/3dgut/kernels/cuda/models/shRadiativeGaussianParticles.cuh
@@ -69,15 +69,15 @@ struct ShRadiativeGaussianVolumetricFeaturesParticles : Params, public ExtParams
}
__forceinline__ __device__ tcnn::vec3 fetchPosition(uint32_t particleIdx) const {
- return *(reinterpret_cast<const tcnn::vec3*>(&m_densityRawParameters.ptr[particleIdx].position_0));
+ return *(reinterpret_cast<const tcnn::vec3*>(&m_densityRawParameters.ptr[particleIdx].position_1));
}
__forceinline__ __device__ const tcnn::vec3& position(const TDensityParameters& parameters) const {
- return *(reinterpret_cast<const tcnn::vec3*>(¶meters.position_1));
+ return *(reinterpret_cast<const tcnn::vec3*>(¶meters.position_0));
}
__forceinline__ __device__ const tcnn::vec3& scale(const TDensityParameters& parameters) const {
- return *(reinterpret_cast<const tcnn::vec3*>(¶meters.scale_1));
+ return *(reinterpret_cast<const tcnn::vec3*>(¶meters.scale_0));
}
__forceinline__ __device__ const tcnn::mat3& rotation(const TDensityParameters& parameters) const {
@@ -86,7 +86,7 @@ struct ShRadiativeGaussianVolumetricFeaturesParticles : Params, public ExtParams
}
__forceinline__ __device__ const float& opacity(const TDensityParameters& parameters) const {
- return parameters.density_1;
+ return parameters.density_0;
}
__forceinline__ __device__ bool densityHit(const tcnn::vec3& rayOrigin,
@@ -433,45 +433,45 @@ struct ShRadiativeGaussianVolumetricFeaturesParticles : Params, public ExtParams
// Perform warp reduction
#pragma unroll
for (int mask = 1; mask < warpSize; mask *= 2) {
- densityRawParameters.position_0.x += __shfl_xor_sync(0xffffffff, densityRawParameters.position_0.x, mask);
- densityRawParameters.position_0.y += __shfl_xor_sync(0xffffffff, densityRawParameters.position_0.y, mask);
- densityRawParameters.position_0.z += __shfl_xor_sync(0xffffffff, densityRawParameters.position_0.z, mask);
- densityRawParameters.density_0 += __shfl_xor_sync(0xffffffff, densityRawParameters.density_0, mask);
+ densityRawParameters.position_1.x += __shfl_xor_sync(0xffffffff, densityRawParameters.position_1.x, mask);
+ densityRawParameters.position_1.y += __shfl_xor_sync(0xffffffff, densityRawParameters.position_1.y, mask);
+ densityRawParameters.position_1.z += __shfl_xor_sync(0xffffffff, densityRawParameters.position_1.z, mask);
+ densityRawParameters.density_1 += __shfl_xor_sync(0xffffffff, densityRawParameters.density_1, mask);
densityRawParameters.quaternion_0.x += __shfl_xor_sync(0xffffffff, densityRawParameters.quaternion_0.x, mask);
densityRawParameters.quaternion_0.y += __shfl_xor_sync(0xffffffff, densityRawParameters.quaternion_0.y, mask);
densityRawParameters.quaternion_0.z += __shfl_xor_sync(0xffffffff, densityRawParameters.quaternion_0.z, mask);
densityRawParameters.quaternion_0.w += __shfl_xor_sync(0xffffffff, densityRawParameters.quaternion_0.w, mask);
- densityRawParameters.scale_0.x += __shfl_xor_sync(0xffffffff, densityRawParameters.scale_0.x, mask);
- densityRawParameters.scale_0.y += __shfl_xor_sync(0xffffffff, densityRawParameters.scale_0.y, mask);
- densityRawParameters.scale_0.z += __shfl_xor_sync(0xffffffff, densityRawParameters.scale_0.z, mask);
+ densityRawParameters.scale_1.x += __shfl_xor_sync(0xffffffff, densityRawParameters.scale_1.x, mask);
+ densityRawParameters.scale_1.y += __shfl_xor_sync(0xffffffff, densityRawParameters.scale_1.y, mask);
+ densityRawParameters.scale_1.z += __shfl_xor_sync(0xffffffff, densityRawParameters.scale_1.z, mask);
}
// First thread in the warp performs the atomic add
if ((tileThreadIdx & (warpSize - 1)) == 0) {
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_0.x, densityRawParameters.position_0.x);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_0.y, densityRawParameters.position_0.y);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_0.z, densityRawParameters.position_0.z);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].density_0, densityRawParameters.density_0);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_1.x, densityRawParameters.position_1.x);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_1.y, densityRawParameters.position_1.y);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_1.z, densityRawParameters.position_1.z);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].density_1, densityRawParameters.density_1);
atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].quaternion_0.x, densityRawParameters.quaternion_0.x);
atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].quaternion_0.y, densityRawParameters.quaternion_0.y);
atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].quaternion_0.z, densityRawParameters.quaternion_0.z);
atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].quaternion_0.w, densityRawParameters.quaternion_0.w);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_0.x, densityRawParameters.scale_0.x);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_0.y, densityRawParameters.scale_0.y);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_0.z, densityRawParameters.scale_0.z);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_1.x, densityRawParameters.scale_1.x);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_1.y, densityRawParameters.scale_1.y);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_1.z, densityRawParameters.scale_1.z);
}
} else {
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_0.x, densityRawParameters.position_0.x);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_0.y, densityRawParameters.position_0.y);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_0.z, densityRawParameters.position_0.z);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].density_0, densityRawParameters.density_0);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_1.x, densityRawParameters.position_1.x);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_1.y, densityRawParameters.position_1.y);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].position_1.z, densityRawParameters.position_1.z);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].density_1, densityRawParameters.density_1);
atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].quaternion_0.x, densityRawParameters.quaternion_0.x);
atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].quaternion_0.y, densityRawParameters.quaternion_0.y);
atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].quaternion_0.z, densityRawParameters.quaternion_0.z);
atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].quaternion_0.w, densityRawParameters.quaternion_0.w);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_0.x, densityRawParameters.scale_0.x);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_0.y, densityRawParameters.scale_0.y);
- atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_0.z, densityRawParameters.scale_0.z);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_1.x, densityRawParameters.scale_1.x);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_1.y, densityRawParameters.scale_1.y);
+ atomicAdd(&m_densityRawParameters.gradPtr[particleIdx].scale_1.z, densityRawParameters.scale_1.z);
}
}
diff --git a/threedgut_tracer/include/3dgut/kernels/cuda/renderers/gutKBufferRenderer.cuh b/threedgut_tracer/include/3dgut/kernels/cuda/renderers/gutKBufferRenderer.cuh
index 23536c5..e38d088 100644
--- a/threedgut_tracer/include/3dgut/kernels/cuda/renderers/gutKBufferRenderer.cuh
+++ b/threedgut_tracer/include/3dgut/kernels/cuda/renderers/gutKBufferRenderer.cuh
@@ -340,10 +340,10 @@ struct GUTKBufferRenderer : Params {
}
DensityRawParameters densityRawParametersGrad;
- densityRawParametersGrad.density_0 = 0.0f;
- densityRawParametersGrad.position_0 = make_float3(0.0f);
+ densityRawParametersGrad.density_1 = 0.0f;
+ densityRawParametersGrad.position_1 = make_float3(0.0f);
densityRawParametersGrad.quaternion_0 = make_float4(0.0f);
- densityRawParametersGrad.scale_0 = make_float3(0.0f);
+ densityRawParametersGrad.scale_1 = make_float3(0.0f);
TFeaturesVec featuresGrad = TFeaturesVec::zero(); |
|
Thanks a lot for the feedback. Slang version installed in the conda environment (through slangtorch) is If the Slang version is the issue (to be validated), we should refactor the code to alias this struct. |
|
I was able to recreate this error with from within the conda environment |
|
I am now certain that the problem is due to a mismatch of the slangc version. However, the 1.3.4 version of slangtorch does not work for me, as the accompanying slangc version is still 2025.6.1 on my system. |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Hi,
While run the code, I noticed a discrepancy between the definition and the actual usage of the struct
gaussianParticle_RawParameters_0. This mismatch is causing a compilation error during the training of 3DGUT.Steps to Reproduce
Expected Behavior
In compiling native code stage, build extension module
lib3dgut_ccshould be successful.Actual Behavior
Possible Solution
The text was updated successfully, but these errors were encountered: