
#include <miopen/hipoc_kernel.hpp>
#include <miopen/errors.hpp>
#include <hip/hip_hcc.h>

namespace miopen {

void HIPOCKernelInvoke::run(void* args, std::size_t size) const
{
    HipEventPtr start = nullptr;
    HipEventPtr stop = nullptr;
    void *config[] = {
        HIP_LAUNCH_PARAM_BUFFER_POINTER, args,
        HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
        HIP_LAUNCH_PARAM_END
    };
    if (callback)
    {
        start = make_hip_event();
        stop = make_hip_event();

        hipEventRecord(start.get(), stream);
    }

    // std::cerr << "Launch kernel: " << name << std::endl;
    auto status = hipHccModuleLaunchKernel(fun, gdims[0], gdims[1], gdims[2], ldims[0], ldims[1], ldims[2], 0, stream, nullptr, (void**)&config);
    if(status != hipSuccess) MIOPEN_THROW_HIP_STATUS(status, "Failed to launch kernel");

    if (callback)
    {
        hipEventRecord(stop.get(), stream);
        hipEventSynchronize(stop.get());
        callback(start.get(), stop.get());
    }
}

HIPOCKernelInvoke HIPOCKernel::Invoke(hipStream_t stream, std::function<void(hipEvent_t, hipEvent_t)> callback)
{
    return HIPOCKernelInvoke{stream, fun, ldims, gdims, name, callback};
}
}


#include <miopen/hipoc_kernel.hpp>
#include <miopen/errors.hpp>
#include <hip/hip_hcc.h>

namespace miopen {

void HIPOCKernelInvoke::run(void* args, std::size_t size) const
{
    HipEventPtr start = nullptr;
    HipEventPtr stop = nullptr;
    void *config[] = {
        HIP_LAUNCH_PARAM_BUFFER_POINTER, args,
        HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
        HIP_LAUNCH_PARAM_END
    };
    if (callback)
    {
        start = make_hip_event();
        stop = make_hip_event();
    }

    // std::cerr << "Launch kernel: " << name << std::endl;
    auto status = hipHccModuleLaunchKernel(fun, gdims[0], gdims[1], gdims[2], ldims[0], ldims[1], ldims[2], 0, stream, nullptr, (void**)&config, start.get(), stop.get());
    if(status != hipSuccess) MIOPEN_THROW_HIP_STATUS(status, "Failed to launch kernel");

    if (callback)
    {
        hipEventSynchronize(stop.get());
        callback(start.get(), stop.get());
    }
}

HIPOCKernelInvoke HIPOCKernel::Invoke(hipStream_t stream, std::function<void(hipEvent_t, hipEvent_t)> callback)
{
    return HIPOCKernelInvoke{stream, fun, ldims, gdims, name, callback};
}
}

Use more precise kernel timers for HIP.
Use more precise kernel timers for HIP.

Use new hipHccModuleLaunchKernel API which optionally accepts start /
stop events.  This allows host code to precisely time the kernel start
and stop times.  The traditional CUDA/HIP approach involves enqueueing
additional events and the time for the GPU to process those events is
included in the measurement.  The diff in this patch shows a succinct
comparison of the two approaches since we replaced the traditional
HIP/CUDA code with the new more surgical one.

This requires the following HIP checkin that supports the new API:

    commit 6656d33f75bbc13c48e6e24b02ec43dc43af1fac
    Author: Ben Sander <ben.sander@amd.com>
    Date:   Thu Apr 6 21:00:00 2017 -0500

-ben
