Skip to content

Commit

Permalink
Fix isosurfaces on the GPU
Browse files Browse the repository at this point in the history
  • Loading branch information
@szellmann
szellmann committed Nov 2, 2024
1 parent a891421 commit 9682d1a
Show file tree
Hide file tree
Showing 4 changed files with 68 additions and 88 deletions.
13 changes: 0 additions & 13 deletions DeviceCopyableObjects.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1371,9 +1371,6 @@ struct Array
const void *data{nullptr};
size_t len{0};
TypeInfo typeInfo;

VSNRAY_FUNC bool empty() const
{ return len==0; }
};

enum class Attribute
Expand Down Expand Up @@ -1592,16 +1589,6 @@ struct Geometry
Array index;
Array normal;
Array tangent;

VSNRAY_FUNC
inline bool isValid() const
{
if (type == ISOSurface) {
return !primitives.empty() && as<dco::ISOSurface>(0).numValues > 0;
}
// TODO..
return true;
}
};

// Sampler //
Expand Down
10 changes: 4 additions & 6 deletions scene/VisionarayScene.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -86,7 +86,6 @@ void VisionaraySceneImpl::commit()
if (!dco::validHandle(geomID)) continue;

const dco::Geometry &geom = deviceState()->dcos.geometries[geomID];
if (!geom.isValid()) continue;

switch (geom.type) {
case dco::Geometry::Triangle:
Expand Down Expand Up @@ -136,7 +135,6 @@ void VisionaraySceneImpl::commit()
if (!dco::validHandle(geomID)) continue;

const dco::Geometry &geom = deviceState()->dcos.geometries[geomID];
if (!geom.isValid()) continue;

binned_sah_builder builder;

Expand Down Expand Up @@ -199,7 +197,6 @@ void VisionaraySceneImpl::commit()
if (!dco::validHandle(geomID)) continue;

const dco::Geometry &geom = deviceState()->dcos.geometries[geomID];
if (!geom.isValid()) continue;

dco::BLS bls;
bls.blsID = m_BLSs.alloc(bls);
Expand Down Expand Up @@ -321,12 +318,13 @@ void VisionaraySceneImpl::attachInstance(
void VisionaraySceneImpl::attachGeometry(
dco::Geometry geom, unsigned geomID, unsigned userID)
{
if (!geom.isValid())
return;

#if defined(WITH_CUDA) || defined(WITH_HIP)
m_gpuScene->attachGeometry(geom, geomID, userID);
#else

if (geom.primitives.len == 0)
return;

m_geometries.set(geomID, geom.geomID);
m_objIds.set(geomID, userID);

Expand Down
130 changes: 64 additions & 66 deletions scene/VisionaraySceneGPU.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -151,7 +151,6 @@ void VisionaraySceneGPU::commit()
if (!dco::validHandle(geomID)) continue;

dco::Geometry geom = deviceState()->dcos.geometries[geomID];
if (!geom.isValid()) continue;

switch (geom.type) {
case dco::Geometry::Triangle:
Expand Down Expand Up @@ -201,7 +200,6 @@ void VisionaraySceneGPU::commit()
if (!dco::validHandle(geomID)) continue;

const dco::Geometry &geom = deviceState()->dcos.geometries[geomID];
if (!geom.isValid()) continue;

binned_sah_builder builder;

Expand Down Expand Up @@ -306,7 +304,6 @@ void VisionaraySceneGPU::commit()
if (!dco::validHandle(geomID)) continue;

const dco::Geometry &geom = deviceState()->dcos.geometries[geomID];
if (!geom.isValid()) continue;

dco::BLS bls;
bls.blsID = m_impl->parent->m_BLSs.alloc(bls);
Expand Down Expand Up @@ -408,75 +405,76 @@ void VisionaraySceneGPU::dispatch()
void VisionaraySceneGPU::attachGeometry(
dco::Geometry geom, unsigned geomID, unsigned userID)
{
if (geom.primitives.len == 0)
return;

m_impl->parent->m_geometries.set(geomID, geom.geomID);
m_impl->parent->m_objIds.set(geomID, userID);

// Patch geomID into scene primitives
// (first copy to CPU, patch there, then copy back...)
if (geom.primitives.len > 0) {
if (geom.type == dco::Geometry::Triangle) {
std::vector<basic_triangle<3,float>> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::Quad) {
std::vector<basic_triangle<3,float>> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::Sphere) {
std::vector<basic_sphere<float>> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::Cone) {
std::vector<dco::Cone> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::Cylinder) {
std::vector<basic_cylinder<float>> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::BezierCurve) {
std::vector<dco::BezierCurve> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::ISOSurface) {
dco::ISOSurface iso;
CUDA_SAFE_CALL(cudaMemcpy(&iso, geom.primitives.data,
sizeof(iso), cudaMemcpyDeviceToHost));
iso.isoID = geomID;
iso.geomID = geomID;
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, &iso,
sizeof(iso), cudaMemcpyHostToDevice));
if (geom.type == dco::Geometry::Triangle) {
std::vector<basic_triangle<3,float>> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::Quad) {
std::vector<basic_triangle<3,float>> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::Sphere) {
std::vector<basic_sphere<float>> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::Cone) {
std::vector<dco::Cone> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::Cylinder) {
std::vector<basic_cylinder<float>> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::BezierCurve) {
std::vector<dco::BezierCurve> hostData(geom.primitives.len);
CUDA_SAFE_CALL(cudaMemcpy(hostData.data(), geom.primitives.data,
hostData.size() * sizeof(hostData[0]), cudaMemcpyDeviceToHost));
for (size_t i=0;i<geom.primitives.len;++i) {
hostData[i].geom_id = geomID;
}
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, hostData.data(),
hostData.size() * sizeof(hostData[0]), cudaMemcpyHostToDevice));
} else if (geom.type == dco::Geometry::ISOSurface) {
dco::ISOSurface iso;
CUDA_SAFE_CALL(cudaMemcpy(&iso, geom.primitives.data,
sizeof(iso), cudaMemcpyDeviceToHost));
iso.isoID = geomID;
iso.geomID = geomID;
CUDA_SAFE_CALL(cudaMemcpy((void *)geom.primitives.data, &iso,
sizeof(iso), cudaMemcpyHostToDevice));
}


Expand Down
3 changes: 0 additions & 3 deletions scene/VisionaraySceneGPU.hip
View file
Original file line number Diff line number Diff line change
Expand Up @@ -166,7 +166,6 @@ void VisionaraySceneGPU::commit()
if (!dco::validHandle(geomID)) continue;

dco::Geometry geom = deviceState()->dcos.geometries[geomID];
if (!geom.isValid()) continue;

switch (geom.type) {
case dco::Geometry::Triangle:
Expand Down Expand Up @@ -216,7 +215,6 @@ void VisionaraySceneGPU::commit()
if (!dco::validHandle(geomID)) continue;

const dco::Geometry &geom = deviceState()->dcos.geometries[geomID];
if (!geom.isValid()) continue;

binned_sah_builder builder;

Expand Down Expand Up @@ -321,7 +319,6 @@ void VisionaraySceneGPU::commit()
if (!dco::validHandle(geomID)) continue;

const dco::Geometry &geom = deviceState()->dcos.geometries[geomID];
if (!geom.isValid()) continue;

dco::BLS bls;
bls.blsID = m_impl->parent->m_BLSs.alloc(bls);
Expand Down

0 comments on commit 9682d1a

Please sign in to comment.