                               thrust::device_vector<MortonBound>::iterator boundsBegin,
                               thrust::device_vector<MortonBound>::iterator boundsEnd);
 
+struct InvalidAxis {
+    __host__ __device__
+    inline bool operator()(const MortonBound b) const {
+        return (b.min & 0xE0000000) >= 0xC0000000;
+    }
+};
+
 void MortonDacrtNodes::Create(RayContainer& rayContainer, SpheresGeometry& spheres) {
     
     // TestMortonEncoding();

                                                        RawPointer(intermediateBounds),
                                                        thrust::raw_pointer_cast(&*boundsBegin), boundsSize);

    std::cout << "intermediateBounds\n" << intermediateBounds << std::endl;
    
    std::cout << "bounds:\n";
    for (int i = 0; i < boundsSize; ++i) {
        std::cout << i << ": " << boundsBegin[i];
        if (i < boundsSize-1)
            std::cout << "\n";
    }
    std::cout << std::endl;
    
    ReduceMinMaxMortonByAxisPass2<<<1, 128>>>(RawPointer(intermediateBounds), intermediateBounds.size(),
                                              thrust::raw_pointer_cast(&*boundsBegin), boundsSize);

    std::cout << "bounds:\n";
    for (int i = 0; i < boundsSize; ++i) {
        std::cout << i << ": " << boundsBegin[i];
        if (i < boundsSize-1)
            std::cout << "\n";
    }
    std::cout << std::endl;
    
}

void TestMortonEncoding() {


    __host__ __device__ inline MortonCode& operator=(const unsigned int rhs) { code = rhs; return *this; }
    __host__ __device__ inline void operator+=(const MortonCode rhs) { code += rhs; }
    __host__ __device__ inline MortonCode operator&(const unsigned int rhs) { return MortonCode(code & rhs); }
    __host__ __device__ inline operator unsigned int() const { return code; }
    
    __host__ __device__ inline MortonCode WithoutAxis() const { return MortonCode(code & 0x1FFFFFFF); }
