#include "../../rocprim/include/rocprim/device/{{algo_type}}_{{ algo_name }}.hpp"
#include "../../rocprim/include/rocprim/device/detail/device_config_helper.hpp"
#include "../utils/tuning_utils.hpp"
#include <hip/hip_runtime.h>
#include <chrono>
#include <stdio.h>

{% block type_definitions %}
using key_type = {{ key_type }};
{% if value_type %}
using value_type = {{ value_type }};
{% endif %}
{% endblock %}

{% block tuning_config %}
using tuning_config = rocprim::{{ algo_name }}_config<{{ config_params }}>;
{% endblock %}

{% block extern_c_wrapper %}
extern "C" float {{algo_type}}_{{ algo_name }}_wrapper(
    {% block arguments required %}{% endblock %}
)
{
    hipStream_t stream = 0;

    {% block setup %}{% endblock %}

    {# Temporary storage section #}
    void* d_temporary_storage = nullptr;
    size_t temporary_storage_bytes = 0;
    {% block storage_query %}{% endblock %}
    HIP_CHECK(hipMalloc(&d_temporary_storage, temporary_storage_bytes));

    hipEvent_t start, stop;
    HIP_CHECK(hipEventCreate(&start));
    HIP_CHECK(hipEventCreate(&stop));

    HIP_CHECK(hipEventRecord(start, stream));
    {% block algorithm_call required %}{% endblock %}
    HIP_CHECK(hipEventRecord(stop, stream));
    HIP_CHECK(hipEventSynchronize(stop));

    float milliseconds;
    HIP_CHECK(hipEventElapsedTime(&milliseconds, start, stop));

    HIP_CHECK(hipEventDestroy(start));
    HIP_CHECK(hipEventDestroy(stop));
    HIP_CHECK(hipFree(d_temporary_storage));

    return milliseconds;
}
{% endblock %}