[Feature]: Support of pipeline

[yiakwy-xpu-ml-framework-team]

Suggestion Description

The latest kernel implementation in GPGPU arch for data center relies on pipeline features to overlap memory accessing latencies over computation:

Here is an example : flashinfer-ai/flashinfer#491

Hence we need to have a simple yet effective pipeline interface to work best with current HW specification even though hareware may not support in its best (NO TMA for example).

API to be supported while work in its best way in AMD MI30X series chips:

template <hip::thread_scope Scope>
class hip::pipeline {
public:
  pipeline() = delete;

  __host__ __device__ ~pipeline();

  pipeline& operator=(pipeline const&) = delete;

  __host__ __device__ void producer_acquire();

  __host__ __device__ void producer_commit();

  __host__ __device__ void consumer_wait();

  template <typename Rep, typename Period>
  __host__ __device__ bool consumer_wait_for(cuda::std::chrono::duration<Rep, Period> const& duration);

  template <typename Clock, typename Duration>
  __host__ __device__
  bool consumer_wait_until(cuda::std::chrono::time_point<Clock, Duration> const& time_point);

  __host__ __device__ void consumer_release();

  __host__ __device__ bool quit();
};

Operating System

Ubuntu 22.04

GPU

MI300/MI308

ROCm Component

ROCm 6.2/6.3/6.4