Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
24 changes: 0 additions & 24 deletions include/madrona/math.inl
Original file line number Diff line number Diff line change
Expand Up @@ -756,30 +756,6 @@ Quat Quat::fromAngularVec(Vector3 v)
Quat Quat::fromBasis(Vector3 a, Vector3 b, Vector3 c)
{
//Modified from glm::quat_cast
#if 0
===============================================================================
The MIT License
-------------------------------------------------------------------------------
Copyright (c) 2005 - G-Truc Creation

Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
#endif

float four_x_squared_minus1 = a.x - b.y - c.z;
float four_y_squared_minus1 = b.y - a.x - c.z;
Expand Down
153 changes: 0 additions & 153 deletions include/madrona/mesh_bvh.inl
Original file line number Diff line number Diff line change
Expand Up @@ -166,153 +166,6 @@ bool MeshBVH::traceRay(math::Vector3 ray_o,
return ray_hit;
}

#if 0
bool MeshBVH::traceRay(math::Vector3 ray_o,
math::Vector3 ray_d,
float *out_hit_t,
math::Vector3 *out_hit_normal,
void* shared,
TraversalStack *stack,
float t_max) const
{
using namespace math;
constexpr float diveps = 0.0000001f;

Diag3x3 inv_d = Diag3x3::fromVec(ray_d).inv();

RayIsectTxfm tri_isect_txfm = computeRayIsectTxfm(ray_o, ray_d, inv_d);

uint32_t previous_stack_size = stack->size;

stack->push(0);

#ifdef SHARED_STACK
const int32_t mwgpu_warp_id = threadIdx.x / 32;
const int32_t mwgpu_warp_lane = threadIdx.x % 32;
const int32_t num_smem_bytes_per_warp =
(mwGPU::SharedMemStorage::numBytesPerWarp()/4)*4;

auto sharedMem = ((char*)shared) + mwgpu_warp_id * num_smem_bytes_per_warp +
SHARED_STACK_SIZE * sizeof(int32_t) * mwgpu_warp_lane;
int32_t* shared_stack = (int32_t*)sharedMem;
shared_stack[0] = 0;
#endif

bool ray_hit = false;
Vector3 closest_hit_normal = Vector3{0,0,0};

while (stack->size > previous_stack_size) {
int32_t node_idx = stack->pop();
const QBVHNode &node = nodes[node_idx];

float rayXInv = copysignf(ray_d.x == 0 ? 1/diveps : 1/ray_d.x,ray_d.x);
float rayYInv = copysignf(ray_d.y == 0 ? 1/diveps : 1/ray_d.y,ray_d.y);
float rayZInv = copysignf(ray_d.z == 0 ? 1/diveps : 1/ray_d.z,ray_d.z);
//NVIDIA's method, transform for ray plane to quantized space. Shift to IEEE exponent bits.

#ifdef MADRONA_GPU_MODE
float dirQuantX = __uint_as_float((node.expX + 127) << 23) * rayXInv;
float dirQuantY = __uint_as_float((node.expY + 127) << 23) * rayYInv;
float dirQuantZ = __uint_as_float((node.expZ + 127) << 23) * rayZInv;
#else
float dirQuantX = std::bit_cast<float>((node.expX + 127) << 23) * rayXInv;
float dirQuantY = std::bit_cast<float>((node.expY + 127) << 23) * rayYInv;
float dirQuantZ = std::bit_cast<float>((node.expZ + 127) << 23) * rayZInv;
#endif

float originQuantX = (node.minX - ray_o.x) * rayXInv;
float originQuantY = (node.minY - ray_o.y) * rayYInv;
float originQuantZ = (node.minZ - ray_o.z) * rayZInv;

for (CountT i = 0; i < MeshBVH::nodeWidth; i++) {
if (!node.hasChild(i)) {
continue; // Technically this could be break?
};


float t_near_x = node.qMinX[i] * dirQuantX + originQuantX;
float t_near_y = node.qMinY[i] * dirQuantY + originQuantY;
float t_near_z = node.qMinZ[i] * dirQuantZ + originQuantZ;

float t_far_x = node.qMaxX[i] * dirQuantX + originQuantX;
float t_far_y = node.qMaxY[i] * dirQuantY + originQuantY;
float t_far_z = node.qMaxZ[i] * dirQuantZ + originQuantZ;
/*
madrona::math::AABB child_aabb {
.pMin = {
node.minX + std::bit_cast<float>((node.expX + 127) << 23) * node.qMinX[i],
node.minY + std::bit_cast<float>((node.expY + 127) << 23) * node.qMinY[i],
node.minZ + std::bit_cast<float>((node.expZ + 127) << 23) * node.qMinZ[i],
},
.pMax = {
node.minX + std::bit_cast<float>((node.expX + 127) << 23) * node.qMaxX[i],
node.minY + std::bit_cast<float>((node.expY + 127) << 23) * node.qMaxY[i],
node.minZ + std::bit_cast<float>((node.expZ + 127) << 23) * node.qMaxZ[i]
},
};

float t_near_x = (child_aabb[tri_isect_txfm.nearX] -
tri_isect_txfm.oNear.x) *
tri_isect_txfm.invDirNear.x;
float t_near_y = (child_aabb[tri_isect_txfm.nearY] -
tri_isect_txfm.oNear.y) *
tri_isect_txfm.invDirNear.y;
float t_near_z = (child_aabb[tri_isect_txfm.nearZ] -
tri_isect_txfm.oNear.z) *
tri_isect_txfm.invDirNear.z;

float t_far_x = (child_aabb[tri_isect_txfm.farX] -
tri_isect_txfm.oFar.x) *
tri_isect_txfm.invDirFar.x;
float t_far_y = (child_aabb[tri_isect_txfm.farY] -
tri_isect_txfm.oFar.y) *
tri_isect_txfm.invDirFar.y;
float t_far_z = (child_aabb[tri_isect_txfm.farZ] -
tri_isect_txfm.oFar.z) *
tri_isect_txfm.invDirFar.z;
float t_near = fmaxf(t_near_x, fmaxf(t_near_y,
fmaxf(t_near_z, 0.f)));
float t_far = fminf(t_far_x, fminf(t_far_y,
fminf(t_far_z, t_max)));

*/

float t_near = fmaxf(fminf(t_near_x,t_far_x), fmaxf(fminf(t_near_y,t_far_y),
fmaxf(fminf(t_near_z,t_far_z), 0.f)));
float t_far = fminf(fmaxf(t_far_x,t_near_x), fminf(fmaxf(t_far_y,t_near_y),
fminf(fmaxf(t_far_z,t_near_z), t_max)));

if (t_near <= t_far) {
if (node.isLeaf(i)) {
int32_t leaf_idx = node.leafIDX(i);

float hit_t;
Vector3 leaf_hit_normal;
bool leaf_hit = traceRayLeaf(leaf_idx, node.triSize[i], tri_isect_txfm,
ray_o, t_max, &hit_t, &leaf_hit_normal);

if (leaf_hit) {
ray_hit = true;
t_max = hit_t;
closest_hit_normal = leaf_hit_normal;
}
} else {
// assert(stack->size < 32);
stack->push(node.childrenIdx[i]);
}
}
}
}

if (!ray_hit) {
return false;
}

*out_hit_t = t_max;
*out_hit_normal = closest_hit_normal;
return ray_hit;
}
#endif

bool MeshBVH::traceRayLeaf(int32_t leaf_idx,
int32_t num_tris,
Expand Down Expand Up @@ -482,12 +335,6 @@ bool MeshBVH::rayTriangleIntersection(
// normalize U, V, W, and T
const float rcpDet = 1.0f / det;

#if 0
hit.u = U * rcpDet;
hit.v = V * rcpDet;
hit.w = W * rcpDet;
hit.t = T * rcpDet;
#endif

*out_hit_t = T * rcpDet;
*bary_out = Vector3{U,V,W} * rcpDet;
Expand Down
10 changes: 0 additions & 10 deletions include/madrona/rand.inl
Original file line number Diff line number Diff line change
Expand Up @@ -209,16 +209,6 @@ constexpr float bitsToFloat01(uint32_t rand_bits)
// implementations that generate in the range (0, 1]
return (rand_bits >> 8_u32) * 0x1p-24f;

#if 0
constexpr uint32_t exponent = 0x3f800000;
uint32_t raw = (exponent | (rand_bits >> 9)) - 1;

#ifdef MADRONA_GPU_MODE
return __uint_as_float(raw);
#else
return std::bit_cast<float>(raw);
#endif
#endif
}

}
Expand Down
6 changes: 0 additions & 6 deletions include/madrona/state.inl
Original file line number Diff line number Diff line change
Expand Up @@ -282,12 +282,6 @@ ComponentID StateManager::componentID() const
template <typename ArchetypeT>
ArchetypeID StateManager::archetypeID() const
{
#if 0
uint32_t id = TypeTracker::typeID<ArchetypeID>();

assert(id != TypeTracker::unassignedTypeID &&
"Trying to access an unregistered archetype!");
#endif
return ArchetypeID {
TypeTracker::typeID<ArchetypeT>(),
};
Expand Down
65 changes: 0 additions & 65 deletions include/madrona/system.hpp

This file was deleted.

50 changes: 0 additions & 50 deletions include/madrona/system.inl

This file was deleted.

4 changes: 0 additions & 4 deletions src/bridge/sim.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -79,10 +79,6 @@ static void setupRenderTasks(TaskGraphBuilder &builder,
Span<const TaskGraphNodeID> deps,
bool update_mats = false)
{
#if 0
builder.addToGraph<ParallelForNode<
Engine, printTransforms, Position, Rotation, Scale, ObjectID>>(deps);
#endif
RenderingSystem::setupTasks(builder, deps, update_mats);
}

Expand Down
7 changes: 0 additions & 7 deletions src/common/table.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -28,13 +28,6 @@ Table::Table(const TypeInfo *component_types, CountT num_components,
for (int i = 0; i < (int)num_components; i++) {
const TypeInfo &type = component_types[i];

#if 0
// 3rd argument is offsetting the start from the page aligned boundary
// to avoid everything mapping to the same cache sets. Should revisit -
// maybe add a random offset for each Table as well?
columns_.emplace_back(type.numBytes, type.alignment,
MADRONA_CACHE_LINE * (i + 1), ICfg::maxRowsPerTable);
#endif

size_t column_bytes_per_row = (size_t)type.numBytes;
columns_[i] = malloc(
Expand Down
10 changes: 0 additions & 10 deletions src/core/system.cpp

This file was deleted.

7 changes: 0 additions & 7 deletions src/mw/cuda_exec.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -646,13 +646,6 @@ static GPUCompileResults compileCode(
}
};

#if 0
// Don't need the device runtime without dynamic parallelism

checkLinker(cuLinkAddFile(linker, CU_JIT_INPUT_LIBRARY,
MADRONA_CUDADEVRT_PATH,
0, nullptr, nullptr));
#endif

nvJitLinkInputType linker_input_type;
if (opt_mode == CompileConfig::OptMode::LTO) {
Expand Down
Loading
Loading