import os
os.environ['PYTORCH_ROCM_ARCH']='gfx942'  # Maybe 'gfx942:xnack-' later?
os.environ["CXX"] = "clang++"

# This script provides a template for using load_inline to run a HIP kernel for MI300
from torch.utils.cpp_extension import load_inline
from task import input_t, output_t

CPP_WRAPPER = """
void fp8_mm(torch::Tensor a, torch::Tensor b, torch::Tensor a_scale, torch::Tensor b_scale, torch::Tensor c);
"""

HIP_SRC = r"""
#include <hip/hip_runtime.h>
#include <hip/amd_detail/amd_hip_fp8.h>
#include <hip/amd_detail/amd_hip_bf16.h>
#include <rocwmma/rocwmma.hpp>

// Helper for checking HIP errors (optional, good for standalone)
#define HIP_CHECK(cmd)                                                         \\
    do {                                                                       \\
        hipError_t e = cmd;                                                    \\
        if (e != hipSuccess) {                                                 \\
            printf("HIP error %s:%d '%s' (%d)\\n", __FILE__, __LINE__,          \\
                   hipGetErrorString(e), e);                                   \\
            abort();                                                           \\
        }                                                                      \\
    } while (0)

// Typedefs for HIP datatypes
typedef __hip_fp8_e4m3_fnuz fp8_t;
typedef __hip_bfloat16 bf16_t;

// Constants from problem description
constexpr int SCALE_BLOCK_DIM_K_CONST = 128;
constexpr int SCALE_BLOCK_DIM_N_CONST = 128; // For b_scale indexing

// rocWMMA Parameters
constexpr uint32_t MFMA_M_TILE_M = 32u; // MFMA output tile M dimension
constexpr uint32_t MFMA_N_TILE_N = 32u; // MFMA output tile N dimension
constexpr uint32_t MFMA_K_TILE_K = 16u; // K-slice per MFMA operation

// Thread block configuration (one wave = 64 threads for gfx942)
constexpr uint32_t TBLOCK_X_DIM = rocwmma::Constants::AMDGCN_WAVE_SIZE_64;
constexpr uint32_t TBLOCK_Y_DIM = 1u; 
constexpr uint32_t TOTAL_THREADS_PER_BLOCK = TBLOCK_X_DIM * TBLOCK_Y_DIM;

// Number of elements per thread for different fragment types
constexpr int NUM_FLOAT_PER_THREAD_ACC = 16; // For 32x32x16 fp8 MFMA, float accumulator

// rocWMMA Fragment Types
using namespace rocwmma;

using GlobalLayoutA = col_major; 
using GlobalLayoutB = col_major; 
using GlobalLayoutC = row_major; 

// NOTE: The MatrixA_MfmaFrag and MatrixB_MfmaFrag definitions depend on MFMA_M_TILE_M, MFMA_N_TILE_N, MFMA_K_TILE_K,
// which are fixed. The Accumulator_MfmaFrag also uses these fixed MFMA tile dimensions.
// The templated TB_M, TB_N parameters will define how many of these base MFMA operations are done.
using MatrixA_MfmaFrag = fragment<matrix_a, MFMA_M_TILE_M, MFMA_N_TILE_N, MFMA_K_TILE_K, fp8_t, col_major>;
using MatrixB_MfmaFrag = fragment<matrix_b, MFMA_M_TILE_M, MFMA_N_TILE_N, MFMA_K_TILE_K, fp8_t, row_major>;
using Accumulator_MfmaFrag = fragment<accumulator, MFMA_M_TILE_M, MFMA_N_TILE_N, MFMA_K_TILE_K, float, GlobalLayoutC>;

// --- Helper Device Function: Load a tile from Global to LDS (Vectorized) ---
template <uint32_t TB_M, uint32_t TB_N, uint32_t TB_K>
__device__ void load_gmem_tile_to_lds_vectorized(
    fp8_t* lds_a_target, const fp8_t* __restrict__ global_a_ptr,
    fp8_t* lds_b_target, const fp8_t* __restrict__ global_b_ptr,
    int current_k_block_base,
    int M_param, int N_out_param, int K_param, 
    uint32_t block_c_base_m_dim, uint32_t block_c_base_n_dim, 
    uint32_t wave_lane_id)
{
    constexpr uint32_t VECTOR_SIZE_FP8 = 4; 
    static_assert(TB_M % VECTOR_SIZE_FP8 == 0, "TB_M must be a multiple of VECTOR_SIZE_FP8");
    static_assert(TB_N % VECTOR_SIZE_FP8 == 0, "TB_N must be a multiple of VECTOR_SIZE_FP8");

    // Load A: Global A (M x K, col-major) to LDS A (TB_K x TB_M)
    // LDS A stores TB_K rows, each of TB_M elements.
    constexpr uint32_t VECTORS_PER_K_SLICE_A = TB_M / VECTOR_SIZE_FP8;
    constexpr uint32_t TOTAL_A_VECTORS_IN_LDS = VECTORS_PER_K_SLICE_A * TB_K;
    static_assert(TOTAL_A_VECTORS_IN_LDS % TOTAL_THREADS_PER_BLOCK == 0, "Total A vectors in LDS not evenly divisible by threads per block.");
    constexpr uint32_t A_VECTORS_PER_THREAD = TOTAL_A_VECTORS_IN_LDS / TOTAL_THREADS_PER_BLOCK;

    for (uint32_t i = 0; i < A_VECTORS_PER_THREAD; ++i) {
        uint32_t flat_vector_idx = wave_lane_id * A_VECTORS_PER_THREAD + i;
        
        uint32_t k_idx_in_tile = flat_vector_idx / VECTORS_PER_K_SLICE_A;
        uint32_t vec_m_idx     = flat_vector_idx % VECTORS_PER_K_SLICE_A;
        
        uint32_t gmem_a_k = current_k_block_base + k_idx_in_tile;
        uint32_t m_start_in_block = vec_m_idx * VECTOR_SIZE_FP8; 
        uint32_t gmem_a_m_start = block_c_base_m_dim + m_start_in_block;

        if (gmem_a_k < K_param && gmem_a_m_start < M_param) { 
            const int* gmem_src_ptr = reinterpret_cast<const int*>(
                &global_a_ptr[gmem_a_k * M_param + gmem_a_m_start]);
            
            int* lds_dst_ptr = reinterpret_cast<int*>(
                &lds_a_target[k_idx_in_tile * TB_M + m_start_in_block]);
            
            *lds_dst_ptr = *gmem_src_ptr;
        }
    }

    // Load B: Global B (N x K, col-major) to LDS B (TB_K x TB_N)
    constexpr uint32_t VECTORS_PER_K_SLICE_B = TB_N / VECTOR_SIZE_FP8;
    constexpr uint32_t TOTAL_B_VECTORS_IN_LDS = VECTORS_PER_K_SLICE_B * TB_K;
    static_assert(TOTAL_B_VECTORS_IN_LDS % TOTAL_THREADS_PER_BLOCK == 0, "Total B vectors in LDS not evenly divisible by threads per block.");
    constexpr uint32_t B_VECTORS_PER_THREAD = TOTAL_B_VECTORS_IN_LDS / TOTAL_THREADS_PER_BLOCK;
    
    for (uint32_t i = 0; i < B_VECTORS_PER_THREAD; ++i) {
        uint32_t flat_vector_idx = wave_lane_id * B_VECTORS_PER_THREAD + i;

        uint32_t k_idx_in_tile = flat_vector_idx / VECTORS_PER_K_SLICE_B;
        uint32_t vec_n_idx     = flat_vector_idx % VECTORS_PER_K_SLICE_B;
        
        uint32_t gmem_b_k = current_k_block_base + k_idx_in_tile;
        uint32_t n_start_in_block = vec_n_idx * VECTOR_SIZE_FP8; 
        uint32_t gmem_b_n_start = block_c_base_n_dim + n_start_in_block;

        if (gmem_b_k < K_param && gmem_b_n_start < N_out_param) { 
            const int* gmem_src_ptr = reinterpret_cast<const int*>(
                &global_b_ptr[gmem_b_k * N_out_param + gmem_b_n_start]);
            
            int* lds_dst_ptr = reinterpret_cast<int*>(
                &lds_b_target[k_idx_in_tile * TB_N + n_start_in_block]);
            
            *lds_dst_ptr = *gmem_src_ptr;
        }
    }
}

// --- Helper Device Function: Compute MMAs on data in LDS ---
template <uint32_t TB_M, uint32_t TB_N, uint32_t TB_K>
__device__ void compute_lds_tile_mma(
    const fp8_t* lds_a_current, 
    const fp8_t* lds_b_current,
    Accumulator_MfmaFrag current_k_block_acc_frags[TB_M / MFMA_M_TILE_M][TB_N / MFMA_N_TILE_N]) // Output param
{
    constexpr uint32_t BLOCKS_PER_TILE_M_VAL = TB_M / MFMA_M_TILE_M;
    constexpr uint32_t BLOCKS_PER_TILE_N_VAL = TB_N / MFMA_N_TILE_N;

    static_assert(TB_M % MFMA_M_TILE_M == 0, "TB_M must be a multiple of MFMA_M_TILE_M");
    static_assert(TB_N % MFMA_N_TILE_N == 0, "TB_N must be a multiple of MFMA_N_TILE_N");
    static_assert(TB_K % MFMA_K_TILE_K == 0, "TB_K must be a multiple of MFMA_K_TILE_K");

    for (uint32_t r_frag_idx = 0; r_frag_idx < BLOCKS_PER_TILE_M_VAL; ++r_frag_idx) {
        for (uint32_t c_frag_idx = 0; c_frag_idx < BLOCKS_PER_TILE_N_VAL; ++c_frag_idx) {
            fill_fragment(current_k_block_acc_frags[r_frag_idx][c_frag_idx], 0.0f);
        }
    }
    
    for (int k_inner_offset = 0; k_inner_offset < TB_K; k_inner_offset += MFMA_K_TILE_K) {
        MatrixA_MfmaFrag a_mfma_sub_frags[BLOCKS_PER_TILE_M_VAL]; 
        MatrixB_MfmaFrag b_mfma_sub_frags[BLOCKS_PER_TILE_N_VAL]; 

        for (uint32_t m_tile_idx = 0; m_tile_idx < BLOCKS_PER_TILE_M_VAL; ++m_tile_idx) {
            const fp8_t* lds_a_sub_tile_ptr = &lds_a_current[(k_inner_offset * TB_M) + (m_tile_idx * MFMA_M_TILE_M)];
            load_matrix_sync(a_mfma_sub_frags[m_tile_idx], lds_a_sub_tile_ptr, TB_M);
        }

        for (uint32_t n_tile_idx = 0; n_tile_idx < BLOCKS_PER_TILE_N_VAL; ++n_tile_idx) {
            const fp8_t* lds_b_sub_tile_ptr = &lds_b_current[(k_inner_offset * TB_N) + (n_tile_idx * MFMA_N_TILE_N)];
            load_matrix_sync(b_mfma_sub_frags[n_tile_idx], lds_b_sub_tile_ptr, TB_N);
        }
        
        for (uint32_t m_frag_idx = 0; m_frag_idx < BLOCKS_PER_TILE_M_VAL; ++m_frag_idx) {
            for (uint32_t n_frag_idx = 0; n_frag_idx < BLOCKS_PER_TILE_N_VAL; ++n_frag_idx) {
                mma_sync(current_k_block_acc_frags[m_frag_idx][n_frag_idx],
                         a_mfma_sub_frags[m_frag_idx],    
                         b_mfma_sub_frags[n_frag_idx],    
                         current_k_block_acc_frags[m_frag_idx][n_frag_idx]);
            }
        }
    } 
}

template <
    uint32_t TB_M, // Tile Block M dimension for C
    uint32_t TB_N, // Tile Block N dimension for C
    uint32_t TB_K  // Tile Block K dimension (outer loop over K)
>
__global__ void __launch_bounds__(TOTAL_THREADS_PER_BLOCK)
    custom_kernel_rocwmma_pipelined(
    const fp8_t* __restrict__ global_a_ptr,
    const fp8_t* __restrict__ global_b_ptr,
    const float* __restrict__ global_a_scale_ptr,
    const float* __restrict__ global_b_scale_ptr,
    bf16_t* __restrict__ global_c_ptr,
    int M_param, int N_out_param, int K_param)
{
    constexpr uint32_t MAX_LDS_BYTES = 32768; // Example: 32KB limit
    constexpr uint32_t LDS_USAGE_BYTES = (TB_M * TB_K + TB_N * TB_K) * 2 * sizeof(fp8_t);
    static_assert(LDS_USAGE_BYTES <= MAX_LDS_BYTES, "Selected tile dimensions exceed LDS limit.");

    static_assert(TB_M % MFMA_M_TILE_M == 0, "TB_M must be a multiple of MFMA_M_TILE_M");
    static_assert(TB_N % MFMA_N_TILE_N == 0, "TB_N must be a multiple of MFMA_N_TILE_N");
    static_assert(TB_K % SCALE_BLOCK_DIM_K_CONST == 0, "TB_K must be a multiple of SCALE_BLOCK_DIM_K_CONST (128)");


    constexpr uint32_t BLOCKS_PER_TILE_M_VAL = TB_M / MFMA_M_TILE_M;
    constexpr uint32_t BLOCKS_PER_TILE_N_VAL = TB_N / MFMA_N_TILE_N;

    __shared__ fp8_t lds_a_ping[TB_M * TB_K]; 
    __shared__ fp8_t lds_a_pong[TB_M * TB_K]; 
    __shared__ fp8_t lds_b_ping[TB_N * TB_K]; 
    __shared__ fp8_t lds_b_pong[TB_N * TB_K]; 
    
    fp8_t* lds_a_double_buffers[2] = {lds_a_ping, lds_a_pong};
    fp8_t* lds_b_double_buffers[2] = {lds_b_ping, lds_b_pong};
    int buffer_selector = 0; 

    uint32_t block_c_start_row = blockIdx.x * TB_M;
    uint32_t block_c_start_col = blockIdx.y * TB_N;
    
    uint32_t wave_lane_id = threadIdx.x; 

    Accumulator_MfmaFrag final_accum_frags[BLOCKS_PER_TILE_M_VAL][BLOCKS_PER_TILE_N_VAL];
    for (uint32_t i = 0; i < BLOCKS_PER_TILE_M_VAL; ++i) {
        for (uint32_t j = 0; j < BLOCKS_PER_TILE_N_VAL; ++j) {
            fill_fragment(final_accum_frags[i][j], 0.0f);
        }
    }
    
    if (K_param == 0) { return; } 

    load_gmem_tile_to_lds_vectorized<TB_M, TB_N, TB_K>(
        lds_a_double_buffers[buffer_selector], global_a_ptr,
        lds_b_double_buffers[buffer_selector], global_b_ptr,
        0, 
        M_param, N_out_param, K_param,
        block_c_start_row, block_c_start_col,
        wave_lane_id);

    int num_k_outer_iterations = K_param / TB_K;

    for (int k_iter_idx = 0; k_iter_idx < num_k_outer_iterations; ++k_iter_idx) {
        int current_k_block_base = k_iter_idx * TB_K;

        fp8_t* lds_a_for_compute = lds_a_double_buffers[buffer_selector];
        fp8_t* lds_b_for_compute = lds_b_double_buffers[buffer_selector];

        int next_k_block_base = current_k_block_base + TB_K;
        if (next_k_block_base < K_param) { 
            load_gmem_tile_to_lds_vectorized<TB_M, TB_N, TB_K>(
                lds_a_double_buffers[1 - buffer_selector], global_a_ptr,
                lds_b_double_buffers[1 - buffer_selector], global_b_ptr,
                next_k_block_base,
                M_param, N_out_param, K_param,
                block_c_start_row, block_c_start_col,
                wave_lane_id);
        }

        synchronize_workgroup();

        Accumulator_MfmaFrag current_k_block_accum_frags[BLOCKS_PER_TILE_M_VAL][BLOCKS_PER_TILE_N_VAL];
        
        compute_lds_tile_mma<TB_M, TB_N, TB_K>(lds_a_for_compute, lds_b_for_compute, current_k_block_accum_frags);
        
        uint32_t k_scale_block_idx = current_k_block_base / SCALE_BLOCK_DIM_K_CONST;
        uint32_t b_scale_num_n_blocks = (N_out_param + SCALE_BLOCK_DIM_N_CONST - 1) / SCALE_BLOCK_DIM_N_CONST;

        for (uint32_t m_frag_idx = 0; m_frag_idx < BLOCKS_PER_TILE_M_VAL; ++m_frag_idx) {
            for (uint32_t n_frag_idx = 0; n_frag_idx < BLOCKS_PER_TILE_N_VAL; ++n_frag_idx) {
                
                auto& source_accum_frag = current_k_block_accum_frags[m_frag_idx][n_frag_idx];
                auto& dest_final_accum_frag = final_accum_frags[m_frag_idx][n_frag_idx];

                for (int p_idx = 0; p_idx < NUM_FLOAT_PER_THREAD_ACC; ++p_idx) {
                    uint32_t p_div_4 = p_idx / 4;
                    uint32_t p_mod_4 = p_idx % 4;
                    uint32_t lane_div_32 = wave_lane_id / 32; 
                    uint32_t lane_mod_32 = wave_lane_id % 32; 
                    
                    uint32_t m_coord_in_mfma_tile = (8 * p_div_4) + (4 * lane_div_32) + p_mod_4;
                    uint32_t n_coord_in_mfma_tile = lane_mod_32;

                    uint32_t global_m_coord = block_c_start_row + m_frag_idx * MFMA_M_TILE_M + m_coord_in_mfma_tile;
                    uint32_t global_n_coord = block_c_start_col + n_frag_idx * MFMA_N_TILE_N + n_coord_in_mfma_tile;

                    float scale_a_val;
                    // Boundary checks are for extreme robustness; with correct launch parameters and M,N divisibility,
                    // global_m_coord and global_n_coord should always be in bounds.
                    if (global_m_coord < M_param) {
                         scale_a_val = global_a_scale_ptr[global_m_coord + k_scale_block_idx * M_param];
                    } else { continue; }

                    float scale_b_val;
                    uint32_t n_b_scale_block_idx = global_n_coord / SCALE_BLOCK_DIM_N_CONST;
                     if (global_n_coord < N_out_param && n_b_scale_block_idx < b_scale_num_n_blocks) {
                        scale_b_val = global_b_scale_ptr[n_b_scale_block_idx + k_scale_block_idx * b_scale_num_n_blocks];
                    } else { continue; }
                    
                    float combined_scale = scale_a_val * scale_b_val;
                    dest_final_accum_frag.x[p_idx] += source_accum_frag.x[p_idx] * combined_scale;
                }
            }
        }
        
        if (next_k_block_base < K_param) { 
             synchronize_workgroup();
        }

        buffer_selector = 1 - buffer_selector;
    } 
    
    for (uint32_t m_frag_idx = 0; m_frag_idx < BLOCKS_PER_TILE_M_VAL; ++m_frag_idx) {
        for (uint32_t n_frag_idx = 0; n_frag_idx < BLOCKS_PER_TILE_N_VAL; ++n_frag_idx) {
            
            auto& result_frag_to_store = final_accum_frags[m_frag_idx][n_frag_idx];

            for (int p_idx = 0; p_idx < NUM_FLOAT_PER_THREAD_ACC; ++p_idx) {
                uint32_t p_div_4 = p_idx / 4;
                uint32_t p_mod_4 = p_idx % 4;
                uint32_t lane_div_32 = wave_lane_id / 32;
                uint32_t lane_mod_32 = wave_lane_id % 32;
                uint32_t m_coord_in_mfma_tile = (8 * p_div_4) + (4 * lane_div_32) + p_mod_4;
                uint32_t n_coord_in_mfma_tile = lane_mod_32;

                uint32_t global_m_coord = block_c_start_row + m_frag_idx * MFMA_M_TILE_M + m_coord_in_mfma_tile;
                uint32_t global_n_coord = block_c_start_col + n_frag_idx * MFMA_N_TILE_N + n_coord_in_mfma_tile;

                if (global_m_coord < M_param && global_n_coord < N_out_param) {
                    global_c_ptr[global_m_coord * N_out_param + global_n_coord] = static_cast<bf16_t>(result_frag_to_store.x[p_idx]);
                }
            }
        }
    }
}

// Entry-point from PyTorch (this function signature is fixed, and unchangeable)
#include <torch/types.h> 

void fp8_mm(torch::Tensor a, torch::Tensor b, torch::Tensor a_scale, torch::Tensor b_scale, torch::Tensor c) {
  int M = a.size(0);
  int K = a.size(1); 
  int N_out = b.size(0); 

  // Define the tile configuration for this specific compilation/launch.
  // To test other configurations, change these values and recompile.
  // These must be multiples of MFMA tile sizes and respect LDS constraints.
  // BLOCK_K_OUTER must also be multiple of SCALE_BLOCK_DIM_K_CONST (128) as per kernel static_assert
  constexpr uint32_t LAUNCH_BLOCK_M = 64u; 
  constexpr uint32_t LAUNCH_BLOCK_N = 64u; 
  constexpr uint32_t LAUNCH_BLOCK_K_OUTER = 128u;


  
  
  uint32_t b_scale_expected_dim0 = (N_out + SCALE_BLOCK_DIM_N_CONST -1) / SCALE_BLOCK_DIM_N_CONST;



  dim3 grid_dim(M / LAUNCH_BLOCK_M, N_out / LAUNCH_BLOCK_N);
  dim3 block_dim(TBLOCK_X_DIM, TBLOCK_Y_DIM); 
  
  hipStream_t stream = nullptr; 
  // hipStream_t stream = c10::hip::getCurrentHIPStreamMasqueradingAsCUDA(); // If using PyTorch's stream

  custom_kernel_rocwmma_pipelined<LAUNCH_BLOCK_M, LAUNCH_BLOCK_N, LAUNCH_BLOCK_K_OUTER><<<grid_dim, block_dim, 0, stream>>>(
    reinterpret_cast<const fp8_t*>(a.data_ptr()), 
    reinterpret_cast<const fp8_t*>(b.data_ptr()), 
    a_scale.data_ptr<float>(), 
    b_scale.data_ptr<float>(), 
    reinterpret_cast<bf16_t*>(c.data_ptr()), 
    M, N_out, K);
    
  hipError_t last_error = hipGetLastError();
  if (last_error != hipSuccess) { 
      std::string err_msg = "HIP kernel launch failed: ";
      err_msg += hipGetErrorString(last_error);
  }
}
"""


module = load_inline(
    name='mdda',
    cpp_sources=[CPP_WRAPPER],
    cuda_sources=[HIP_SRC],
    functions=['fp8_mm'],
    verbose=True,
    extra_cuda_cflags=["--offload-arch=gfx942", "-std=c++20"],
)

def custom_kernel(data: input_t) -> output_t:
    a, b, a_scale, b_scale, c = data
    module.fp8_mm(a, b, a_scale, b_scale, c)
    return c


python submission.py
[92mSuccessfully preprocessed all matching files.[0m

/home/runner/.cache/torch_extensions/py310_cpu/mdda/main.cpp -> /home/runner/.cache/torch_extensions/py310_cpu/mdda/main.cpp [skipped, no changes]
/home/runner/.cache/torch_extensions/py310_cpu/mdda/cuda.cu -> /home/runner/.cache/torch_extensions/py310_cpu/mdda/hip.hip [ok]
Total number of unsupported CUDA function calls: 0


Total number of replaced kernel launches: 1
[1/3] clang++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=mdda -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1016\" -isystem /usr/local/lib/python3.10/dist-packages/torch/include -isystem /usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -isystem /usr/local/lib/python3.10/dist-packages/torch/include/THH -isystem /opt/rocm-6.3.1/include -isystem /usr/include/python3.10 -fPIC -std=c++17 -c /home/runner/.cache/torch_extensions/py310_cpu/mdda/main.cpp -o main.o -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -fPIC
[2/3] /opt/rocm-6.3.1/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=mdda -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1016\" -isystem /usr/local/lib/python3.10/dist-packages/torch/include -isystem /usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -isystem /usr/local/lib/python3.10/dist-packages/torch/include/THH -isystem /opt/rocm-6.3.1/include -isystem /usr/include/python3.10 -fPIC -std=c++17 -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -fPIC -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -DHIP_ENABLE_WARP_SYNC_BUILTINS=1 --offload-arch=gfx942 -fno-gpu-rdc --offload-arch=gfx942 -std=c++20 -c /home/runner/.cache/torch_extensions/py310_cpu/mdda/hip.hip -o hip.cuda.o 
[3/3] clang++ main.o hip.cuda.o -shared -L/usr/local/lib/python3.10/dist-packages/torch/lib -lc10 -lc10_hip -ltorch_cpu -ltorch_hip -ltorch -ltorch_python -L/opt/rocm-6.3.1/lib -lamdhip64 -o mdda.so

python eval.py benchmark /tmp/tmpo440lm8p
[92mSuccessfully preprocessed all matching files.[0m

/home/runner/.cache/torch_extensions/py310_cpu/mdda/main.cpp -> /home/runner/.cache/torch_extensions/py310_cpu/mdda/main.cpp [skipped, no changes]
/home/runner/.cache/torch_extensions/py310_cpu/mdda/cuda.cu -> /home/runner/.cache/torch_extensions/py310_cpu/mdda/hip.hip [skipped, already hipified]
Total number of unsupported CUDA function calls: 0


Total number of replaced kernel launches: 1
[1/1] clang++ hip.cuda.o main.o -shared -L/usr/local/lib/python3.10/dist-packages/torch/lib -lc10 -lc10_hip -ltorch_cpu -ltorch_hip -ltorch -ltorch_python -L/opt/rocm-6.3.1/lib -lamdhip64 -o mdda.so
