///////////////////////////////////////////////////////////////////////////////
//
// The MIT License
//
// Copyright (c) 2006 Scientific Computing and Imaging Institute,
// University of Utah (USA)
//
// License for the specific language governing rights and limitations under
// 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.
//
///////////////////////////////////////////////////////////////////////////////

#ifndef ELVIS_CORE_INNER_INTEGRAL_ADAPTIVE_TRAPEZOIDAL_CU
#define ELVIS_CORE_INNER_INTEGRAL_ADAPTIVE_TRAPEZOIDAL_CU

namespace ElVis
{
  //    /// Adaptive trapezoidal integration using a runtime stack to handle
  //    recursion.
  //    /// Value are kept for subsequent sampling.
  //    template<typename T, unsigned int n>
  //    struct InnerIntegralAdaptiveTrapezoidal
  //    {
  //        public:
  //            struct StackPoint
  //            {
  //                template<typename FieldFunc>
  //                __device__
  //                void Evaluate(const TransferFunction* densityFunc,
  //                              TransferFunctionChannel channel,
  //                    const FieldFunc& fieldFunc)
  //                {
  //                    T s = fieldFunc(TVal);
  //                    F = densityFunc->Sample(channel, s);
  //                }

  //                __device__ void Reset()
  //                {
  //                    TVal = MAKE_FLOAT(1e30);
  //                }

  //                __device__ bool IsUninitialized() const
  //                {
  //                    return TVal == MAKE_FLOAT(1e30);
  //                }

  //                __device__ StackPoint& operator=(const StackPoint& rhs)
  //                {
  //                    TVal = rhs.TVal;
  //                    F = rhs.F;
  //                    return *this;
  //                }

  //                T TVal;
  //                T F;
  //            };

  //            struct StackEntry
  //            {
  //                __device__ void CalculateMidpointT()
  //                {
  //                    Mid().TVal = Left().TVal + (Right().TVal -
  //                    Left().TVal)/2.0;
  //                }

  //                __device__ void SetT(const T& t0, const T& t1)
  //                {
  //                    Left().TVal = t0;
  //                    Right().TVal = t1;
  //                    CalculateMidpointT();
  //                }

  //                __device__ T GetH() const
  //                {
  //                    return Right().TVal - Left().TVal;
  //                }

  //                template<typename FieldFunc>
  //                __device__ void EvaluateAll(const TransferFunction*
  //                densityFunc,
  //                                           TransferFunctionChannel channel,
  //                    const FieldFunc& fieldFunc)
  //                {
  //                    for(unsigned int i = 0; i < 3; ++i)
  //                    {
  //                        points[i].Evaluate(densityFunc, channel, fieldFunc);
  //                    }
  //                }

  //                __device__ StackPoint& Left()  { return points[0]; }
  //                __device__ StackPoint& Mid()  { return points[1]; }
  //                __device__ StackPoint& Right() { return points[2]; }

  //                __device__ const StackPoint& Left() const  { return
  //                points[0]; }
  //                __device__ const StackPoint& Mid() const  { return
  //                points[1]; }
  //                __device__ const StackPoint& Right() const { return
  //                points[2]; }

  //                StackPoint points[3];
  //            };

  //            template<typename FieldFunctionType>
  //            __device__ void Integrate(const T& t0, const T& t1, const
  //            TransferFunction* transferFunction,
  //                TransferFunctionChannel channel,
  //                const FieldFunctionType& fieldFunction, const T&
  //                globalEpsilon,
  //                const T& globalIntegralEstimate,
  //                                     const T& maxFunctionValue, bool&
  //                                     reachedMaxRecursion,
  //                                     bool traceEnabled)
  //            {
  //                if( traceEnabled )
  //                {
  //                    printf("Global Epsilon %f, globalIntegralEstimate %f,
  //                    maxValue %f\n", globalEpsilon, globalIntegralEstimate,
  //                    maxFunctionValue);
  //                }

  //                const unsigned int maxRecursion = n;
  //                reachedMaxRecursion = false;
  //                StackEntry stack[maxRecursion];

  //                stack[0].SetT(t0, t1);
  //                stack[0].EvaluateAll(transferFunction, channel,
  //                fieldFunction);

  //                stack[1].Left() = stack[0].Left();
  //                stack[1].Mid().Reset();
  //                stack[1].Right() = stack[0].Mid();

  //                unsigned int minimumDepth = 2;

  //                int i = 1;
  //                t[0] = t0;
  //                f[0] = stack[0].Left().F;
  //                I[0] = 0.0;
  //                adaptiveIndex = 0;
  //                while( i > 0 )
  //                {
  //                    reachedMaxRecursion |= (i == maxRecursion-1);
  //                    if( stack[i].Mid().IsUninitialized() )
  //                    {
  //                        bool needToSubdivide = false;

  //                        stack[i].CalculateMidpointT();
  //                        stack[i].Mid().Evaluate(transferFunction, channel,
  //                        fieldFunction);

  //                        if( i < minimumDepth )
  //                        {
  //                            needToSubdivide = true;
  //                        }
  //                        else
  //                        {
  //                            T I0 = stack[i].GetH()/MAKE_FLOAT(2.0) *
  //                            (stack[i].Left().F + stack[i].Right().F);
  //                            T I1 = stack[i].GetH()/MAKE_FLOAT(4.0) *
  //                            (stack[i].Left().F + 2.0*stack[i].Mid().F +
  //                            stack[i].Right().F);
  //                            T localEpsilon =
  //                            globalEpsilon*globalIntegralEstimate *
  //                            (stack[i].GetH()/stack[0].GetH());

  //                            if( traceEnabled )
  //                            {
  //                                printf("Level %d, Interval (%f, %f, %f),
  //                                values (%f, %f, %f) I0 = %f, I1 = %f,
  //                                localEpsilon = %f\n", i,
  //                                stack[i].Left().TVal, stack[i].Mid().TVal,
  //                                stack[i].Right().TVal,
  //                                       stack[i].Left().F, stack[i].Mid().F,
  //                                       stack[i].Right().F, I0, I1,
  //                                       localEpsilon);
  //                            }

  //                            ElVisFloat h = stack[i].GetH()/MAKE_FLOAT(2.0);

  //                            if( stack[i].Left().F == MAKE_FLOAT(0.0) &&
  //                                stack[i].Mid().F == MAKE_FLOAT(0.0) &&
  //                                stack[i].Right().F == MAKE_FLOAT(0.0) )
  //                            {
  //                                ElVis::Interval<ElVisFloat> range =
  //                                fieldFunction.EstimateRange(stack[i].Left().TVal,
  //                                stack[i].Right().TVal);

  //                                ElVisFloat maxValue =
  //                                transferFunction->GetMaxValue(channel,
  //                                range);
  //                                T maxSegmentError =
  //                                (maxFunctionValue*h)/globalIntegralEstimate;
  //                                T updatedSegmentError =
  //                                (maxValue*h)/globalIntegralEstimate;

  //                                if( traceEnabled )
  //                                {
  //                                    printf("All 3 values are 0.  Scalar
  //                                    range is (%f, %f), maxSegmentError %f,
  //                                    updatedSegmentError %f\n",
  //                                    range.GetLow(), range.GetHigh(),
  //                                           maxSegmentError,
  //                                           updatedSegmentError);
  //                                }

  //                                if( updatedSegmentError > localEpsilon && i
  //                                < maxRecursion-1 )
  //                                {
  //                                    needToSubdivide = true;
  //                                }
  //                            }
  //                            else if( stack[i].Left().F == MAKE_FLOAT(0.0) ||
  //                                stack[i].Mid().F == MAKE_FLOAT(0.0) ||
  //                                stack[i].Right().F == MAKE_FLOAT(0.0) )
  //                            {
  //                                // If any of the samples are 0, then we know
  //                                there is a breakpoint somewhere and we
  //                                should subdivide.
  //                                T maxSegmentError =
  //                                (maxFunctionValue*h)/globalIntegralEstimate;

  //                                ElVis::Interval<ElVisFloat> range =
  //                                fieldFunction.EstimateRange(stack[i].Left().TVal,
  //                                stack[i].Right().TVal);

  //                                ElVisFloat maxValue =
  //                                transferFunction->GetMaxValue(channel,
  //                                range);
  //                                T updatedSegmentError =
  //                                (maxValue*h)/globalIntegralEstimate;

  //                                if( traceEnabled )
  //                                {
  //                                    printf("At least one value is 0.  Scalar
  //                                    range is (%f, %f), maxSegmentError %f,
  //                                    updatedSegmentError %f\n",
  //                                    range.GetLow(), range.GetHigh(),
  //                                           maxSegmentError,
  //                                           updatedSegmentError);
  //                                }

  //                                if( traceEnabled )
  //                                {
  //                                    printf("One of the samples is 0,
  //                                    maxSegmentError = %f, localEpsilon =
  //                                    %f\n", maxSegmentError, localEpsilon);
  //                                }
  //                                if(updatedSegmentError > localEpsilon && i <
  //                                maxRecursion-1 )
  //                                {
  //                                    needToSubdivide = true;
  //                                }
  //                            }
  //                            else
  //                            {
  //                                T errorEstimate =
  //                                fabs(I0-I1)/globalIntegralEstimate;
  //                                if( traceEnabled )
  //                                {
  //                                    printf("No samples 0, errorEstimate =
  //                                    %f, localEpsilon = %f\n", errorEstimate,
  //                                    localEpsilon);
  //                                }
  //                                if( errorEstimate > localEpsilon && i <
  //                                maxRecursion-1 )
  //                                {
  //                                    needToSubdivide = true;
  //                                }
  //                            }
  //                        }

  //                        if( traceEnabled )
  //                        {
  //                            printf("Subdividing = %d\n", needToSubdivide? 1
  //                            : 0);
  //                        }

  //                        if( needToSubdivide )
  //                        {
  //                            stack[i+1].Left() = stack[i].Left();
  //                            stack[i+1].Mid().Reset();
  //                            stack[i+1].Right() = stack[i].Mid();
  //                            i = i + 1;
  //                        }
  //                        else
  //                        {
  //                            T prevValue = I[adaptiveIndex];
  //                            T h = stack[i].GetH()/MAKE_FLOAT(4.0);
  //                            T mid_f = stack[i].Mid().F;
  //                            T right_f = stack[i].Right().F;

  //                            t[adaptiveIndex+1] = stack[i].Mid().TVal;
  //                            t[adaptiveIndex+2] = stack[i].Right().TVal;
  //                            f[adaptiveIndex+1] = mid_f;
  //                            f[adaptiveIndex+2] = right_f;

  //                            T leftContribution = h * (stack[i].Left().F +
  //                            mid_f);
  //                            T rightContribution = h * (mid_f + right_f);

  //                            I[adaptiveIndex+1] = prevValue +
  //                            leftContribution;
  //                            I[adaptiveIndex+2] = prevValue +
  //                            leftContribution+rightContribution;

  //                            if( traceEnabled )
  //                            {
  //                                printf("Integral Value at %f = %f\n",
  //                                t[adaptiveIndex+1], I[adaptiveIndex+1]);
  //                                printf("Integral Value at %f = %f\n",
  //                                t[adaptiveIndex+2], I[adaptiveIndex+2]);
  //                            }
  //                            adaptiveIndex += 2;
  //                        }
  //                    }
  //                    else
  //                    {
  //                        if( stack[i].Right().TVal == stack[i-1].Mid().TVal )
  //                        {
  //                            // We just finished traversing the left side,
  //                            now go to
  //                            // the right.
  //                            stack[i].Left() = stack[i-1].Mid();
  //                            stack[i].Mid().Reset();
  //                            stack[i].Right() = stack[i-1].Right();
  //                        }
  //                        else
  //                        {
  //                            // We finished this branch.  Remove it and go up
  //                            to
  //                            // the next one.
  //                            i = i-1;
  //                        }
  //                    }
  //                }

  //            }

  //            __device__ T SampleInnerIntegral(T t_i, T sample,
  //            TransferFunctionChannel channel, const TransferFunction*
  //            densityFunc) const
  //            {
  //                if( t_i < t[0] ||
  //                    t_i > t[adaptiveIndex] )
  //                {
  //                    return MAKE_FLOAT(0.0);
  //                }

  //                if( t_i == t[0] ) return MAKE_FLOAT(0.0);
  //                if( t_i == t[adaptiveIndex] ) return I[adaptiveIndex];

  //                const T* a = &(t[0]);
  //                const T* b = &(t[adaptiveIndex]);
  //                while(b-a > 1 )
  //                {
  //                    const T* mid = (b-a)/2 + a;
  //                    if( *mid == t_i )
  //                    {
  //                        return I[mid-a];
  //                    }
  //                    if( t_i < *mid )
  //                    {
  //                        b = mid;
  //                    }
  //                    else
  //                    {
  //                        a = mid;
  //                    }
  //                }

  //                T baseline = I[a-t];
  //                T segment = (t_i-*a)/MAKE_FLOAT(2.0) * ( f[a-t] +
  //                densityFunc->Sample(channel, sample));
  //                return baseline+segment;
  //            }

  //            __device__ T OverallValue() const
  //            {
  //                return I[adaptiveIndex];
  //            }

  //            static const unsigned int arraySize = (0x01 << n) + 1;

  //            T t[arraySize];
  //            T f[arraySize];
  //            T I[arraySize];
  //        private:
  //            unsigned int adaptiveIndex;
  //    };
}

#endif

	.file	"tmpxft_002e6cfb_00000000-6_InnerIntegralAdaptiveTrapezoidal.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

///////////////////////////////////////////////////////////////////////////////
//
// The MIT License
//
// Copyright (c) 2006 Scientific Computing and Imaging Institute,
// University of Utah (USA)
//
// License for the specific language governing rights and limitations under
// 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.
//
///////////////////////////////////////////////////////////////////////////////

#ifndef ELVIS_CORE_INNER_INTEGRAL_ADAPTIVE_TRAPEZOIDAL_CU
#define ELVIS_CORE_INNER_INTEGRAL_ADAPTIVE_TRAPEZOIDAL_CU

namespace ElVis
{
  //    /// Adaptive trapezoidal integration using a runtime stack to handle
  //    recursion.
  //    /// Value are kept for subsequent sampling.
  //    template<typename T, unsigned int n>
  //    struct InnerIntegralAdaptiveTrapezoidal
  //    {
  //        public:
  //            struct StackPoint
  //            {
  //                template<typename FieldFunc>
  //                __device__
  //                void Evaluate(const TransferFunction* densityFunc,
  //                              TransferFunctionChannel channel,
  //                    const FieldFunc& fieldFunc)
  //                {
  //                    T s = fieldFunc(TVal);
  //                    F = densityFunc->Sample(channel, s);
  //                }

  //                __device__ void Reset()
  //                {
  //                    TVal = MAKE_FLOAT(1e30);
  //                }

  //                __device__ bool IsUninitialized() const
  //                {
  //                    return TVal == MAKE_FLOAT(1e30);
  //                }

  //                __device__ StackPoint& operator=(const StackPoint& rhs)
  //                {
  //                    TVal = rhs.TVal;
  //                    F = rhs.F;
  //                    return *this;
  //                }

  //                T TVal;
  //                T F;
  //            };

  //            struct StackEntry
  //            {
  //                __device__ void CalculateMidpointT()
  //                {
  //                    Mid().TVal = Left().TVal + (Right().TVal -
  //                    Left().TVal)/2.0;
  //                }

  //                __device__ void SetT(const T& t0, const T& t1)
  //                {
  //                    Left().TVal = t0;
  //                    Right().TVal = t1;
  //                    CalculateMidpointT();
  //                }

  //                __device__ T GetH() const
  //                {
  //                    return Right().TVal - Left().TVal;
  //                }

  //                template<typename FieldFunc>
  //                __device__ void EvaluateAll(const TransferFunction*
  //                densityFunc,
  //                                           TransferFunctionChannel channel,
  //                    const FieldFunc& fieldFunc)
  //                {
  //                    for(unsigned int i = 0; i < 3; ++i)
  //                    {
  //                        points[i].Evaluate(densityFunc, channel, fieldFunc);
  //                    }
  //                }

  //                __device__ StackPoint& Left()  { return points[0]; }
  //                __device__ StackPoint& Mid()  { return points[1]; }
  //                __device__ StackPoint& Right() { return points[2]; }

  //                __device__ const StackPoint& Left() const  { return
  //                points[0]; }
  //                __device__ const StackPoint& Mid() const  { return
  //                points[1]; }
  //                __device__ const StackPoint& Right() const { return
  //                points[2]; }

  //                StackPoint points[3];
  //            };

  //            template<typename FieldFunctionType>
  //            __device__ void Integrate(const T& t0, const T& t1, const
  //            TransferFunction* transferFunction,
  //                TransferFunctionChannel channel,
  //                const FieldFunctionType& fieldFunction, const T&
  //                globalEpsilon,
  //                const T& globalIntegralEstimate,
  //                                     const T& maxFunctionValue, bool&
  //                                     reachedMaxRecursion,
  //                                     bool traceEnabled)
  //            {
  //                if( traceEnabled )
  //                {
  //                    printf("Global Epsilon %f, globalIntegralEstimate %f,
  //                    maxValue %f\n", globalEpsilon, globalIntegralEstimate,
  //                    maxFunctionValue);
  //                }

  //                const unsigned int maxRecursion = n;
  //                reachedMaxRecursion = false;
  //                StackEntry stack[maxRecursion];

  //                stack[0].SetT(t0, t1);
  //                stack[0].EvaluateAll(transferFunction, channel,
  //                fieldFunction);

  //                stack[1].Left() = stack[0].Left();
  //                stack[1].Mid().Reset();
  //                stack[1].Right() = stack[0].Mid();

  //                unsigned int minimumDepth = 2;

  //                int i = 1;
  //                t[0] = t0;
  //                f[0] = stack[0].Left().F;
  //                I[0] = 0.0;
  //                adaptiveIndex = 0;
  //                while( i > 0 )
  //                {
  //                    reachedMaxRecursion |= (i == maxRecursion-1);
  //                    if( stack[i].Mid().IsUninitialized() )
  //                    {
  //                        bool needToSubdivide = false;

  //                        stack[i].CalculateMidpointT();
  //                        stack[i].Mid().Evaluate(transferFunction, channel,
  //                        fieldFunction);

  //                        if( i < minimumDepth )
  //                        {
  //                            needToSubdivide = true;
  //                        }
  //                        else
  //                        {
  //                            T I0 = stack[i].GetH()/MAKE_FLOAT(2.0) *
  //                            (stack[i].Left().F + stack[i].Right().F);
  //                            T I1 = stack[i].GetH()/MAKE_FLOAT(4.0) *
  //                            (stack[i].Left().F + 2.0*stack[i].Mid().F +
  //                            stack[i].Right().F);
  //                            T localEpsilon =
  //                            globalEpsilon*globalIntegralEstimate *
  //                            (stack[i].GetH()/stack[0].GetH());

  //                            if( traceEnabled )
  //                            {
  //                                printf("Level %d, Interval (%f, %f, %f),
  //                                values (%f, %f, %f) I0 = %f, I1 = %f,
  //                                localEpsilon = %f\n", i,
  //                                stack[i].Left().TVal, stack[i].Mid().TVal,
  //                                stack[i].Right().TVal,
  //                                       stack[i].Left().F, stack[i].Mid().F,
  //                                       stack[i].Right().F, I0, I1,
  //                                       localEpsilon);
  //                            }

  //                            ElVisFloat h = stack[i].GetH()/MAKE_FLOAT(2.0);

  //                            if( stack[i].Left().F == MAKE_FLOAT(0.0) &&
  //                                stack[i].Mid().F == MAKE_FLOAT(0.0) &&
  //                                stack[i].Right().F == MAKE_FLOAT(0.0) )
  //                            {
  //                                ElVis::Interval<ElVisFloat> range =
  //                                fieldFunction.EstimateRange(stack[i].Left().TVal,
  //                                stack[i].Right().TVal);

  //                                ElVisFloat maxValue =
  //                                transferFunction->GetMaxValue(channel,
  //                                range);
  //                                T maxSegmentError =
  //                                (maxFunctionValue*h)/globalIntegralEstimate;
  //                                T updatedSegmentError =
  //                                (maxValue*h)/globalIntegralEstimate;

  //                                if( traceEnabled )
  //                                {
  //                                    printf("All 3 values are 0.  Scalar
  //                                    range is (%f, %f), maxSegmentError %f,
  //                                    updatedSegmentError %f\n",
  //                                    range.GetLow(), range.GetHigh(),
  //                                           maxSegmentError,
  //                                           updatedSegmentError);
  //                                }

  //                                if( updatedSegmentError > localEpsilon && i
  //                                < maxRecursion-1 )
  //                                {
  //                                    needToSubdivide = true;
  //                                }
  //                            }
  //                            else if( stack[i].Left().F == MAKE_FLOAT(0.0) ||
  //                                stack[i].Mid().F == MAKE_FLOAT(0.0) ||
  //                                stack[i].Right().F == MAKE_FLOAT(0.0) )
  //                            {
  //                                // If any of the samples are 0, then we know
  //                                there is a breakpoint somewhere and we
  //                                should subdivide.
  //                                T maxSegmentError =
  //                                (maxFunctionValue*h)/globalIntegralEstimate;

  //                                ElVis::Interval<ElVisFloat> range =
  //                                fieldFunction.EstimateRange(stack[i].Left().TVal,
  //                                stack[i].Right().TVal);

  //                                ElVisFloat maxValue =
  //                                transferFunction->GetMaxValue(channel,
  //                                range);
  //                                T updatedSegmentError =
  //                                (maxValue*h)/globalIntegralEstimate;

  //                                if( traceEnabled )
  //                                {
  //                                    printf("At least one value is 0.  Scalar
  //                                    range is (%f, %f), maxSegmentError %f,
  //                                    updatedSegmentError %f\n",
  //                                    range.GetLow(), range.GetHigh(),
  //                                           maxSegmentError,
  //                                           updatedSegmentError);
  //                                }

  //                                if( traceEnabled )
  //                                {
  //                                    printf("One of the samples is 0,
  //                                    maxSegmentError = %f, localEpsilon =
  //                                    %f\n", maxSegmentError, localEpsilon);
  //                                }
  //                                if(updatedSegmentError > localEpsilon && i <
  //                                maxRecursion-1 )
  //                                {
  //                                    needToSubdivide = true;
  //                                }
  //                            }
  //                            else
  //                            {
  //                                T errorEstimate =
  //                                fabs(I0-I1)/globalIntegralEstimate;
  //                                if( traceEnabled )
  //                                {
  //                                    printf("No samples 0, errorEstimate =
  //                                    %f, localEpsilon = %f\n", errorEstimate,
  //                                    localEpsilon);
  //                                }
  //                                if( errorEstimate > localEpsilon && i <
  //                                maxRecursion-1 )
  //                                {
  //                                    needToSubdivide = true;
  //                                }
  //                            }
  //                        }

  //                        if( traceEnabled )
  //                        {
  //                            printf("Subdividing = %d\n", needToSubdivide? 1
  //                            : 0);
  //                        }

  //                        if( needToSubdivide )
  //                        {
  //                            stack[i+1].Left() = stack[i].Left();
  //                            stack[i+1].Mid().Reset();
  //                            stack[i+1].Right() = stack[i].Mid();
  //                            i = i + 1;
  //                        }
  //                        else
  //                        {
  //                            T prevValue = I[adaptiveIndex];
  //                            T h = stack[i].GetH()/MAKE_FLOAT(4.0);
  //                            T mid_f = stack[i].Mid().F;
  //                            T right_f = stack[i].Right().F;

  //                            t[adaptiveIndex+1] = stack[i].Mid().TVal;
  //                            t[adaptiveIndex+2] = stack[i].Right().TVal;
  //                            f[adaptiveIndex+1] = mid_f;
  //                            f[adaptiveIndex+2] = right_f;

  //                            T leftContribution = h * (stack[i].Left().F +
  //                            mid_f);
  //                            T rightContribution = h * (mid_f + right_f);

  //                            I[adaptiveIndex+1] = prevValue +
  //                            leftContribution;
  //                            I[adaptiveIndex+2] = prevValue +
  //                            leftContribution+rightContribution;

  //                            if( traceEnabled )
  //                            {
  //                                printf("Integral Value at %f = %f\n",
  //                                t[adaptiveIndex+1], I[adaptiveIndex+1]);
  //                                printf("Integral Value at %f = %f\n",
  //                                t[adaptiveIndex+2], I[adaptiveIndex+2]);
  //                            }
  //                            adaptiveIndex += 2;
  //                        }
  //                    }
  //                    else
  //                    {
  //                        if( stack[i].Right().TVal == stack[i-1].Mid().TVal )
  //                        {
  //                            // We just finished traversing the left side,
  //                            now go to
  //                            // the right.
  //                            stack[i].Left() = stack[i-1].Mid();
  //                            stack[i].Mid().Reset();
  //                            stack[i].Right() = stack[i-1].Right();
  //                        }
  //                        else
  //                        {
  //                            // We finished this branch.  Remove it and go up
  //                            to
  //                            // the next one.
  //                            i = i-1;
  //                        }
  //                    }
  //                }

  //            }

  //            __device__ T SampleInnerIntegral(T t_i, T sample,
  //            TransferFunctionChannel channel, const TransferFunction*
  //            densityFunc) const
  //            {
  //                if( t_i < t[0] ||
  //                    t_i > t[adaptiveIndex] )
  //                {
  //                    return MAKE_FLOAT(0.0);
  //                }

  //                if( t_i == t[0] ) return MAKE_FLOAT(0.0);
  //                if( t_i == t[adaptiveIndex] ) return I[adaptiveIndex];

  //                const T* a = &(t[0]);
  //                const T* b = &(t[adaptiveIndex]);
  //                while(b-a > 1 )
  //                {
  //                    const T* mid = (b-a)/2 + a;
  //                    if( *mid == t_i )
  //                    {
  //                        return I[mid-a];
  //                    }
  //                    if( t_i < *mid )
  //                    {
  //                        b = mid;
  //                    }
  //                    else
  //                    {
  //                        a = mid;
  //                    }
  //                }

  //                T baseline = I[a-t];
  //                T segment = (t_i-*a)/MAKE_FLOAT(2.0) * ( f[a-t] +
  //                densityFunc->Sample(channel, sample));
  //                return baseline+segment;
  //            }

  //            __device__ T OverallValue() const
  //            {
  //                return I[adaptiveIndex];
  //            }

  //            static const unsigned int arraySize = (0x01 << n) + 1;

  //            T t[arraySize];
  //            T f[arraySize];
  //            T I[arraySize];
  //        private:
  //            unsigned int adaptiveIndex;
  //    };
}

#endif

	.text
	.file	"InnerIntegralAdaptiveTrapezoidal.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
