// Copyright 2021 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
//    list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
//    this list of conditions and the following disclaimer in the documentation
//    and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
//    contributors may be used to endorse or promote products derived from
//    this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

#include "dawn/native/IndirectDrawValidationEncoder.h"

#include <algorithm>
#include <cstdint>
#include <cstdlib>
#include <limits>
#include <memory>
#include <utility>
#include <vector>

#include "dawn/common/Constants.h"
#include "dawn/common/Math.h"
#include "dawn/native/BindGroup.h"
#include "dawn/native/BindGroupLayout.h"
#include "dawn/native/CommandEncoder.h"
#include "dawn/native/ComputePassEncoder.h"
#include "dawn/native/ComputePipeline.h"
#include "dawn/native/Device.h"
#include "dawn/native/InternalPipelineStore.h"
#include "dawn/native/Queue.h"
#include "dawn/native/utils/WGPUHelpers.h"
#include "partition_alloc/pointers/raw_ptr.h"

namespace dawn::native {

namespace {
// NOTE: This must match the workgroup_size attribute on the compute entry point below.
constexpr uint64_t kWorkgroupSize = 64;

// Bitmasks for BatchInfo::flags
constexpr uint32_t kDuplicateBaseVertexInstance = 1;
constexpr uint32_t kIndexedDraw = 2;
constexpr uint32_t kValidationEnabled = 4;
constexpr uint32_t kIndirectFirstInstanceEnabled = 8;
constexpr uint32_t kUseFirstIndexToEmulateIndexBufferOffset = 16;

// Equivalent to the IndirectDraw struct defined in the shader below.
struct IndirectDraw {
    uint32_t indirectOffset;
    uint32_t numIndexBufferElementsLow;
    uint32_t numIndexBufferElementsHigh;
    uint32_t indexOffsetAsNumElements;
};
static_assert(sizeof(IndirectDraw) == sizeof(uint32_t) * 4);
static_assert(alignof(IndirectDraw) == alignof(uint32_t));

// Equivalent to the BatchInfo struct defined in the shader below.
struct BatchInfo {
    uint32_t numDraws;
    uint32_t flags;
};

// The size, in bytes, of the IndirectDraw struct defined in the shader below.
constexpr uint32_t kIndirectDrawByteSize = sizeof(uint32_t) * 4;

// TODO(https://crbug.com/dawn/1108): Propagate validation feedback from this shader in
// various failure modes.
static const char sRenderValidationShaderSource[] = R"(

            const kNumDrawIndirectParams = 4u;

            const kIndexCountEntry = 0u;
            const kFirstIndexEntry = 2u;

            // Bitmasks for BatchInfo::flags
            const kDuplicateBaseVertexInstance = 1u;
            const kIndexedDraw = 2u;
            const kValidationEnabled = 4u;
            const kIndirectFirstInstanceEnabled = 8u;
            const kUseFirstIndexToEmulateIndexBufferOffset = 16u;

            struct IndirectDraw {
                indirectOffset: u32,
                numIndexBufferElementsLow: u32,
                numIndexBufferElementsHigh: u32,
                indexOffsetAsNumElements: u32,
            }

            struct BatchInfo {
                numDraws: u32,
                flags: u32,
                draws: array<IndirectDraw>,
            }

            struct IndirectParams {
                data: array<u32>,
            }

            @group(0) @binding(0) var<storage, read> batch: BatchInfo;
            @group(0) @binding(1) var<storage, read_write> inputParams: IndirectParams;
            @group(0) @binding(2) var<storage, read_write> outputParams: IndirectParams;

            fn numIndirectParamsPerDrawCallInput() -> u32 {
                var numParams = kNumDrawIndirectParams;
                // Indexed Draw has an extra parameter (firstIndex)
                if (bool(batch.flags & kIndexedDraw)) {
                    numParams = numParams + 1u;
                }
                return numParams;
            }

            fn numIndirectParamsPerDrawCallOutput() -> u32 {
                var numParams = numIndirectParamsPerDrawCallInput();
                // 2 extra parameter for duplicated first/baseVertex and firstInstance
                if (bool(batch.flags & kDuplicateBaseVertexInstance)) {
                    numParams = numParams + 2u;
                }
                return numParams;
            }

            fn fail(drawIndex: u32) {
                let numParams = numIndirectParamsPerDrawCallOutput();
                let index = drawIndex * numParams;
                for(var i = 0u; i < numParams; i = i + 1u) {
                    outputParams.data[index + i] = 0u;
                }
            }

            fn set_pass(drawIndex: u32) {
                let numInputParams = numIndirectParamsPerDrawCallInput();
                var outIndex = drawIndex * numIndirectParamsPerDrawCallOutput();
                let inIndex = batch.draws[drawIndex].indirectOffset;

                // The first 2 parameter is reserved for the duplicated first/baseVertex and firstInstance

                if (bool(batch.flags & kDuplicateBaseVertexInstance)) {
                    // first/baseVertex and firstInstance are always last two parameters
                    let dupIndex = inIndex + numInputParams - 2u;
                    outputParams.data[outIndex] = inputParams.data[dupIndex];
                    outputParams.data[outIndex + 1u] = inputParams.data[dupIndex + 1u];

                    outIndex = outIndex + 2u;
                }

                for(var i = 0u; i < numInputParams; i = i + 1u) {
                    outputParams.data[outIndex + i] = inputParams.data[inIndex + i];
                }

                if (bool(batch.flags & kUseFirstIndexToEmulateIndexBufferOffset)) {
                    outputParams.data[outIndex + kFirstIndexEntry] += batch.draws[drawIndex].indexOffsetAsNumElements;
                }
            }

            @compute @workgroup_size(64, 1, 1)
            fn main(@builtin(global_invocation_id) id : vec3u) {
                if (id.x >= batch.numDraws) {
                    return;
                }

                if(!bool(batch.flags & kValidationEnabled)) {
                    set_pass(id.x);
                    return;
                }

                let inputIndex = batch.draws[id.x].indirectOffset;
                if(!bool(batch.flags & kIndirectFirstInstanceEnabled)) {
                    // firstInstance is always the last parameter
                    let firstInstance = inputParams.data[inputIndex + numIndirectParamsPerDrawCallInput() - 1u];
                    if (firstInstance != 0u) {
                        fail(id.x);
                        return;
                    }
                }

                if (!bool(batch.flags & kIndexedDraw)) {
                    set_pass(id.x);
                    return;
                }

                let numIndexBufferElementsHigh = batch.draws[id.x].numIndexBufferElementsHigh;

                if (numIndexBufferElementsHigh >= 2u) {
                    // firstIndex and indexCount are both u32. The maximum possible sum of these
                    // values is 0x1fffffffe, which is less than 0x200000000. Nothing to validate.
                    set_pass(id.x);
                    return;
                }

                let numIndexBufferElementsLow = batch.draws[id.x].numIndexBufferElementsLow;

                let firstIndex = inputParams.data[inputIndex + kFirstIndexEntry];
                if (numIndexBufferElementsHigh == 0u &&
                    numIndexBufferElementsLow < firstIndex) {
                    fail(id.x);
                    return;
                }

                // Note that this subtraction may underflow, but only when
                // numIndexBufferElementsHigh is 1u. The result is still correct in that case.
                let maxIndexCount = numIndexBufferElementsLow - firstIndex;
                let indexCount = inputParams.data[inputIndex + kIndexCountEntry];
                if (indexCount > maxIndexCount) {
                    fail(id.x);
                    return;
                }
                set_pass(id.x);
            }
        )";

ResultOrError<ComputePipelineBase*> GetOrCreateRenderValidationPipeline(DeviceBase* device) {
    InternalPipelineStore* store = device->GetInternalPipelineStore();

    if (store->renderValidationPipeline == nullptr) {
        // If we need to apply the index buffer offset to the first index then
        // we can't handle buffers larger than 4gig otherwise we'll overflow first_index
        // which is a 32bit value.
        //
        // When a buffer is less than 4gig the largest index buffer offset you can pass to
        // SetIndexBuffer is 0xffff_fffe. Otherwise you'll get a validation error. This
        // is converted to count of indices and so at most 0x7fff_ffff.
        //
        // The largest valid first_index would be 0x7fff_ffff. Anything larger will fail
        // the validation used in this compute shader and the validated indirect buffer
        // will have 0,0,0,0,0.
        //
        // Adding 0x7fff_ffff + 0x7fff_ffff does not overflow so as long as we keep
        // maxBufferSize < 4gig we're safe.
        DAWN_ASSERT(!device->ShouldApplyIndexBufferOffsetToFirstIndex() ||
                    device->GetLimits().v1.maxBufferSize < 0x100000000u);

        // Create compute shader module if not cached before.
        if (store->renderValidationShader == nullptr) {
            DAWN_TRY_ASSIGN(store->renderValidationShader,
                            utils::CreateShaderModule(device, sRenderValidationShaderSource));
        }

        Ref<BindGroupLayoutBase> bindGroupLayout;
        DAWN_TRY_ASSIGN(
            bindGroupLayout,
            utils::MakeBindGroupLayout(
                device,
                {
                    {0, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::ReadOnlyStorage},
                    {1, wgpu::ShaderStage::Compute, kInternalStorageBufferBinding},
                    {2, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage},
                },
                /* allowInternalBinding */ true));

        Ref<PipelineLayoutBase> pipelineLayout;
        DAWN_TRY_ASSIGN(pipelineLayout, utils::MakeBasicPipelineLayout(device, bindGroupLayout));

        ComputePipelineDescriptor computePipelineDescriptor = {};
        computePipelineDescriptor.layout = pipelineLayout.Get();
        computePipelineDescriptor.compute.module = store->renderValidationShader.Get();
        computePipelineDescriptor.compute.entryPoint = "main";

        DAWN_TRY_ASSIGN(store->renderValidationPipeline,
                        device->CreateComputePipeline(&computePipelineDescriptor));
    }

    return store->renderValidationPipeline.Get();
}

size_t GetBatchDataSize(uint32_t numDraws) {
    return sizeof(BatchInfo) + (numDraws * kIndirectDrawByteSize);
}

}  // namespace

uint32_t ComputeMaxDrawCallsPerIndirectValidationBatch(const CombinedLimits& limits) {
    const uint64_t batchDrawCallLimitByDispatchSize =
        static_cast<uint64_t>(limits.v1.maxComputeWorkgroupsPerDimension) * kWorkgroupSize;
    const uint64_t batchDrawCallLimitByStorageBindingSize =
        (limits.v1.maxStorageBufferBindingSize - sizeof(BatchInfo)) / kIndirectDrawByteSize;
    return static_cast<uint32_t>(
        std::min({batchDrawCallLimitByDispatchSize, batchDrawCallLimitByStorageBindingSize,
                  uint64_t(std::numeric_limits<uint32_t>::max())}));
}

MaybeError EncodeIndirectDrawValidationCommands(DeviceBase* device,
                                                CommandEncoder* commandEncoder,
                                                RenderPassResourceUsageTracker* usageTracker,
                                                IndirectDrawMetadata* indirectDrawMetadata) {
    DAWN_ASSERT(device->IsLockedByCurrentThreadIfNeeded());
    // Since encoding validation commands may create new objects, verify that the device is alive.
    // TODO(dawn:1199): This check is obsolete if device loss causes device.destroy().
    //   - This function only happens within the context of a TryEncode which would catch the
    //     same issue if device loss implied device.destroy().
    DAWN_TRY(device->ValidateIsAlive());

    struct Batch {
        raw_ptr<const IndirectDrawMetadata::IndirectValidationBatch> metadata;
        uint64_t dataBufferOffset;
        uint64_t dataSize;
        uint64_t inputIndirectOffset;
        uint64_t inputIndirectSize;
        uint64_t outputParamsOffset;
        uint64_t outputParamsSize;
        raw_ptr<BatchInfo, AllowPtrArithmetic> batchInfo;
    };

    struct Pass {
        uint32_t flags;
        raw_ptr<BufferBase> inputIndirectBuffer;
        IndirectDrawMetadata::DrawType drawType;
        uint64_t outputParamsSize = 0;
        uint64_t batchDataSize = 0;
        std::unique_ptr<void, void (*)(void*)> batchData{nullptr, std::free};
        std::vector<Batch> batches;
    };

    // First stage is grouping all batches into passes. We try to pack as many batches into a
    // single pass as possible. Batches can be grouped together as long as they're validating
    // data from the same indirect buffer and draw type, but they may still be split into
    // multiple passes if the number of draw calls in a pass would exceed some (very high)
    // upper bound.
    uint64_t outputParamsSize = 0;
    std::vector<Pass> passes;
    IndirectDrawMetadata::IndexedIndirectBufferValidationInfoMap& bufferInfoMap =
        *indirectDrawMetadata->GetIndexedIndirectBufferValidationInfo();
    if (bufferInfoMap.empty()) {
        return {};
    }

    const uint64_t maxStorageBufferBindingSize = device->GetLimits().v1.maxStorageBufferBindingSize;
    const uint32_t minStorageBufferOffsetAlignment =
        device->GetLimits().v1.minStorageBufferOffsetAlignment;

    const bool applyIndexBufferOffsetToFirstIndex =
        device->ShouldApplyIndexBufferOffsetToFirstIndex();

    for (auto& [config, validationInfo] : bufferInfoMap) {
        const uint64_t indirectDrawCommandSize =
            config.drawType == IndirectDrawMetadata::DrawType::Indexed ? kDrawIndexedIndirectSize
                                                                       : kDrawIndirectSize;

        uint64_t outputIndirectSize = indirectDrawCommandSize;
        if (config.duplicateBaseVertexInstance) {
            outputIndirectSize += 2 * sizeof(uint32_t);
        }

        for (const IndirectDrawMetadata::IndirectValidationBatch& batch :
             validationInfo.GetBatches()) {
            const uint64_t minOffsetFromAlignedBoundary =
                batch.minOffset % minStorageBufferOffsetAlignment;
            const uint64_t minOffsetAlignedDown = batch.minOffset - minOffsetFromAlignedBoundary;

            Batch newBatch;
            newBatch.metadata = &batch;
            newBatch.dataSize = GetBatchDataSize(batch.draws.size());
            newBatch.inputIndirectOffset = minOffsetAlignedDown;
            newBatch.inputIndirectSize =
                batch.maxOffset + indirectDrawCommandSize - minOffsetAlignedDown;

            newBatch.outputParamsSize = batch.draws.size() * outputIndirectSize;
            newBatch.outputParamsOffset = Align(outputParamsSize, minStorageBufferOffsetAlignment);
            outputParamsSize = newBatch.outputParamsOffset + newBatch.outputParamsSize;
            if (outputParamsSize > maxStorageBufferBindingSize) {
                return DAWN_INTERNAL_ERROR("Too many drawIndexedIndirect calls to validate");
            }

            Pass* currentPass = passes.empty() ? nullptr : &passes.back();
            if (currentPass &&
                reinterpret_cast<uintptr_t>(currentPass->inputIndirectBuffer.get()) ==
                    config.inputIndirectBufferPtr &&
                currentPass->drawType == config.drawType) {
                uint64_t nextBatchDataOffset =
                    Align(currentPass->batchDataSize, minStorageBufferOffsetAlignment);
                uint64_t newPassBatchDataSize = nextBatchDataOffset + newBatch.dataSize;
                if (newPassBatchDataSize <= maxStorageBufferBindingSize) {
                    // We can fit this batch in the current pass.
                    newBatch.dataBufferOffset = nextBatchDataOffset;
                    currentPass->batchDataSize = newPassBatchDataSize;
                    currentPass->batches.push_back(newBatch);
                    continue;
                }
            }

            // We need to start a new pass for this batch.
            newBatch.dataBufferOffset = 0;

            Pass newPass{};
            newPass.inputIndirectBuffer = validationInfo.GetIndirectBuffer();
            newPass.drawType = config.drawType;
            newPass.batchDataSize = newBatch.dataSize;
            newPass.batches.push_back(newBatch);
            newPass.flags = 0;
            if (config.duplicateBaseVertexInstance) {
                newPass.flags |= kDuplicateBaseVertexInstance;
            }
            if (config.drawType == IndirectDrawMetadata::DrawType::Indexed) {
                newPass.flags |= kIndexedDraw;

                if (applyIndexBufferOffsetToFirstIndex) {
                    newPass.flags |= kUseFirstIndexToEmulateIndexBufferOffset;
                }
            }
            if (device->IsValidationEnabled()) {
                newPass.flags |= kValidationEnabled;
            }
            if (device->HasFeature(Feature::IndirectFirstInstance)) {
                newPass.flags |= kIndirectFirstInstanceEnabled;
            }
            passes.push_back(std::move(newPass));
        }
    }

    auto* const store = device->GetInternalPipelineStore();
    ScratchBuffer& outputParamsBuffer = store->scratchIndirectStorage;
    ScratchBuffer& batchDataBuffer = store->scratchStorage;

    uint64_t requiredBatchDataBufferSize = 0;
    for (const Pass& pass : passes) {
        requiredBatchDataBufferSize = std::max(requiredBatchDataBufferSize, pass.batchDataSize);
    }
    DAWN_TRY(batchDataBuffer.EnsureCapacity(requiredBatchDataBufferSize));

    DAWN_TRY(outputParamsBuffer.EnsureCapacity(outputParamsSize));
    // We swap the indirect buffer used so we need to explicitly add the usage.
    usageTracker->BufferUsedAs(outputParamsBuffer.GetBuffer(), wgpu::BufferUsage::Indirect);

    // Now we allocate and populate host-side batch data to be copied to the GPU.
    for (Pass& pass : passes) {
        // We use std::malloc here because it guarantees maximal scalar alignment.
        pass.batchData = {std::malloc(pass.batchDataSize), std::free};
        memset(pass.batchData.get(), 0, pass.batchDataSize);
        uint8_t* batchData = static_cast<uint8_t*>(pass.batchData.get());
        for (Batch& batch : pass.batches) {
            batch.batchInfo = new (&batchData[batch.dataBufferOffset]) BatchInfo();
            batch.batchInfo->numDraws = static_cast<uint32_t>(batch.metadata->draws.size());
            batch.batchInfo->flags = pass.flags;

            IndirectDraw* indirectDraw = reinterpret_cast<IndirectDraw*>(batch.batchInfo.get() + 1);
            uint64_t outputParamsOffset = batch.outputParamsOffset;
            for (auto& draw : batch.metadata->draws) {
                // The shader uses this to index an array of u32, hence the division by 4 bytes.
                indirectDraw->indirectOffset =
                    static_cast<uint32_t>((draw.inputBufferOffset - batch.inputIndirectOffset) / 4);
                // The index buffer elements are 64 bit values, and so need to be set as a
                // low uint32_t and a high uint32_t.
                indirectDraw->numIndexBufferElementsLow =
                    static_cast<uint32_t>(draw.numIndexBufferElements & 0xFFFFFFFF);
                indirectDraw->numIndexBufferElementsHigh =
                    static_cast<uint32_t>((draw.numIndexBufferElements >> 32) & 0xFFFFFFFF);

                // This is only used in the GL backend.
                indirectDraw->indexOffsetAsNumElements = draw.indexBufferOffsetInElements;
                indirectDraw++;

                draw.cmd->indirectBuffer = outputParamsBuffer.GetBuffer();
                draw.cmd->indirectOffset = outputParamsOffset;
                if (pass.flags & kIndexedDraw) {
                    outputParamsOffset += kDrawIndexedIndirectSize;
                } else {
                    outputParamsOffset += kDrawIndirectSize;
                }
                if (pass.flags & kDuplicateBaseVertexInstance) {
                    // Add the extra offset for the duplicated base vertex and instance.
                    outputParamsOffset += 2 * sizeof(uint32_t);
                }
            }
        }
    }

    ComputePipelineBase* pipeline;
    DAWN_TRY_ASSIGN(pipeline, GetOrCreateRenderValidationPipeline(device));

    Ref<BindGroupLayoutBase> layout;
    DAWN_TRY_ASSIGN(layout, pipeline->GetBindGroupLayout(0));

    BindGroupEntry bindings[3];
    BindGroupEntry& bufferDataBinding = bindings[0];
    bufferDataBinding.binding = 0;
    bufferDataBinding.buffer = batchDataBuffer.GetBuffer();

    BindGroupEntry& inputIndirectBinding = bindings[1];
    inputIndirectBinding.binding = 1;

    BindGroupEntry& outputParamsBinding = bindings[2];
    outputParamsBinding.binding = 2;
    outputParamsBinding.buffer = outputParamsBuffer.GetBuffer();

    BindGroupDescriptor bindGroupDescriptor = {};
    bindGroupDescriptor.layout = layout.Get();
    bindGroupDescriptor.entryCount = 3;
    bindGroupDescriptor.entries = bindings;

    // Finally, we can now encode our validation and duplication passes. Each pass first does a
    // two WriteBuffer to get batch and pass data over to the GPU, followed by a single compute
    // pass. The compute pass encodes a separate SetBindGroup and Dispatch command for each
    // batch.
    for (const Pass& pass : passes) {
        commandEncoder->APIWriteBuffer(batchDataBuffer.GetBuffer(), 0,
                                       static_cast<const uint8_t*>(pass.batchData.get()),
                                       pass.batchDataSize);

        Ref<ComputePassEncoder> passEncoder = commandEncoder->BeginComputePass();
        passEncoder->APISetPipeline(pipeline);

        inputIndirectBinding.buffer = pass.inputIndirectBuffer;

        for (const Batch& batch : pass.batches) {
            bufferDataBinding.offset = batch.dataBufferOffset;
            bufferDataBinding.size = batch.dataSize;
            inputIndirectBinding.offset = batch.inputIndirectOffset;
            inputIndirectBinding.size = batch.inputIndirectSize;
            outputParamsBinding.offset = batch.outputParamsOffset;
            outputParamsBinding.size = batch.outputParamsSize;

            Ref<BindGroupBase> bindGroup;
            DAWN_TRY_ASSIGN(bindGroup, device->CreateBindGroup(&bindGroupDescriptor));

            const uint32_t numDrawsRoundedUp =
                (batch.batchInfo->numDraws + kWorkgroupSize - 1) / kWorkgroupSize;
            passEncoder->APISetBindGroup(0, bindGroup.Get());
            passEncoder->APIDispatchWorkgroups(numDrawsRoundedUp);
        }

        passEncoder->APIEnd();
    }

    return {};
}

}  // namespace dawn::native