# Diagnostic report

Level: Low | ranking measurement: {"path": "trajectory/20260926_105204_iter-49/output.txt", "valid": true, "reference_and_protocol_verified": true, "COMPILED": "True", "CORRECT": "True", "RUNTIME": "4.036014080047607", "REF_RUNTIME": "44.953800201416016", "SPEEDUP": "11.138167337831922"}
HIP call-group costs are diagnostic, not per-device-kernel ranking or hardware floors.

Call 2: 1.781450 ms | runtime_native
High: null | Native: {"file": "kernel_2.cpp", "entry_lines": [10]}
Tensor bindings: [{"value": "input_0", "shape": [128, 512, 768], "dtype": "float32"}, {"value": "value_5", "shape": [2304, 768], "dtype": "float16"}, {"value": "input_2", "shape": [2304], "dtype": "float32"}, {"value": "value_7", "shape": [128, 512, 2304], "dtype": "float16"}]
Native sites: 10: extern "C" __global__ void __launch_bounds__(256) main_kernel(float* __restrict__ B, half_t* __restrict__ W, float* __restrict__ X, half_t* __restrict__ Y) {
11:   extern __shared__ __align__(1024) uchar buf_dyn_shmem[];
22:     __syncthreads();
44:       __syncthreads();
57:       *((tl_v8f32*)c + ((i_4 * 8) + j_1)) = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(

Call 3: 1.301007 ms | runtime_native
High: null | Native: {"file": "kernel_3.cpp", "entry_lines": [10]}
Tensor bindings: [{"value": "value_7", "shape": [128, 512, 2304], "dtype": "float16"}, {"value": "value_8", "shape": [128, 512, 768], "dtype": "float16"}]
Native sites: 10: extern "C" __global__ void __launch_bounds__(256) main_kernel(half_t* __restrict__ Q, half_t* __restrict__ Y) {
11:   extern __shared__ __align__(1024) uchar buf_dyn_shmem[];
47:     __syncthreads();
77:       *((tl_v8f32*)s + j_1) = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(
138:     __syncthreads();
146:       __syncthreads();
159:       *((tl_v8f32*)o + j_3) = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(

Call 4: 0.633884 ms | runtime_native
High: null | Native: {"file": "kernel_4.cpp", "entry_lines": [10]}
Tensor bindings: [{"value": "value_8", "shape": [128, 512, 768], "dtype": "float16"}, {"value": "value_6", "shape": [768, 768], "dtype": "float16"}, {"value": "input_4", "shape": [768], "dtype": "float32"}, {"value": "value_9", "shape": [128, 512, 768], "dtype": "float32"}]
Native sites: 10: extern "C" __global__ void __launch_bounds__(256) main_kernel(float* __restrict__ B, half_t* __restrict__ W, half_t* __restrict__ X, float* __restrict__ Y) {
11:   extern __shared__ __align__(1024) uchar buf_dyn_shmem[];
22:     __syncthreads();
34:       __syncthreads();
47:       *((tl_v8f32*)c + ((i_4 * 4) + j_1)) = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(

Call 0: 0.026760 ms | runtime_native
High: null | Native: {"file": "kernel_0.cpp", "entry_lines": [10]}
Tensor bindings: [{"value": "input_1", "shape": [2304, 768], "dtype": "float32"}, {"value": "value_5", "shape": [2304, 768], "dtype": "float16"}]
Native sites: 10: extern "C" __global__ void __launch_bounds__(256) main_kernel(float* __restrict__ X, half_t* __restrict__ Y) {

Call 1: 0.013760 ms | runtime_native
High: null | Native: {"file": "kernel_1.cpp", "entry_lines": [10]}
Tensor bindings: [{"value": "input_3", "shape": [768, 768], "dtype": "float32"}, {"value": "value_6", "shape": [768, 768], "dtype": "float16"}]
Native sites: 10: extern "C" __global__ void __launch_bounds__(256) main_kernel(float* __restrict__ X, half_t* __restrict__ Y) {
Dataflow: {"edges": [], "bindings": [{"call": 0, "params": [{"value": "input_1", "shape": [2304, 768], "dtype": "float32"}, {"value": "value_5", "shape": [2304, 768], "dtype": "float16"}]}, {"call": 1, "params": [{"value": "input_3", "shape": [768, 768], "dtype": "float32"}, {"value": "value_6", "shape": [768, 768], "dtype": "float16"}]}, {"call": 2, "params": [{"value": "input_0", "shape": [128, 512, 768], "dtype": "float32"}, {"value": "value_5", "shape": [2304, 768], "dtype": "float16"}, {"value": "input_2", "shape": [2304], "dtype": "float32"}, {"value": "value_7", "shape": [128, 512, 2304], "dtype": "float16"}]}, {"call": 3, "params": [{"value": "value_7", "shape": [128, 512, 2304], "dtype": "float16"}, {"value": "value_8", "shape": [128, 512, 768], "dtype": "float16"}]}, {"call": 4, "params": [{"value": "value_8", "shape": [128, 512, 768], "dtype": "float16"}, {"value": "value_6", "shape": [768, 768], "dtype": "float16"}, {"value": "input_4", "shape": [768], "dtype": "float32"}, {"value": "value_9", "shape": [128, 512, 768], "dtype": "float32"}]}], "interpretation": "Observed tensor return/argument dependencies where available; allocation/ABI bindings are not proof of memory read/write or synchronization."}
Counter status: available | candidate association: unknown | missing: []
Device kernel: {"symbol": "void at::native::(anonymous namespace)::distribution_elementwise_grid_stride_kernel<float, 4, at::native::templates::cuda::normal_and_transform<float, float, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1}>(at::TensorIteratorBase&, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1})::{lambda(hiprandStatePhilox4_32_10*)#2}, at::native::(anonymous namespace)::distribution_nullary_kernel<float, float, HIP_vector_type<float, 4u>, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_and_transform<float, float, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1}>(at::TensorIteratorBase&, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1})::{lambda(hiprandStatePhilox4_32_10*)#2}, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1}>(at::TensorIteratorBase&, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_and_transform<float, float, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1}>(at::TensorIteratorBase&, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1})::{lambda(hiprandStatePhilox4_32_10*)#2} const&, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1})::{lambda(int, float)#1}>(long, at::PhiloxCudaState, at::native::templates::cuda::normal_and_transform<float, float, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1}>(at::TensorIteratorBase&, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1})::{lambda(hiprandStatePhilox4_32_10*)#2}, at::native::(anonymous namespace)::distribution_nullary_kernel<float, float, HIP_vector_type<float, 4u>, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_and_transform<float, float, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1}>(at::TensorIteratorBase&, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1})::{lambda(hiprandStatePhilox4_32_10*)#2}, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1}>(at::TensorIteratorBase&, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_and_transform<float, float, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1}>(at::TensorIteratorBase&, at::CUDAGeneratorImpl*, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1})::{lambda(hiprandStatePhilox4_32_10*)#2} const&, at::native::templates::cuda::normal_kernel<at::CUDAGeneratorImpl*>(at::TensorBase const&, double, double, at::CUDAGeneratorImpl*)::{lambda()#1}::operator()() const::{lambda()#2}::operator()() const::{lambda(float)#1})::{lambda(int, float)#1})", "kernel_id": "5", "dispatches": 1, "metrics": {"GRBM_GUI_ACTIVE": {"median": 6559529.0, "min": 6559529.0, "max": 6559529.0, "samples": 1}, "SQ_INSTS_SALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 1}, "SQ_INSTS_VALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 1}, "SQ_WAVES": {"median": 2048.0, "min": 2048.0, "max": 2048.0, "samples": 1}}, "trace_samples": 1, "trace_median_ms": 2.317412, "candidate_calls": [], "association": "unmapped_or_setup"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "37", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 4572407.5, "min": 4567882.0, "max": 5426983.0, "samples": 4}, "SQ_INSTS_SALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 4}, "SQ_INSTS_VALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 4}, "SQ_WAVES": {"median": 36864.0, "min": 36864.0, "max": 37376.0, "samples": 4}}, "trace_samples": 4, "trace_median_ms": 1.818829, "candidate_calls": [0, 1, 2, 3, 4], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "38", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 3485692.0, "min": 3474272.0, "max": 3982729.0, "samples": 4}, "SQ_INSTS_SALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 4}, "SQ_INSTS_VALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 4}, "SQ_WAVES": {"median": 65536.0, "min": 65536.0, "max": 65536.0, "samples": 4}}, "trace_samples": 4, "trace_median_ms": 1.2825665, "candidate_calls": [0, 1, 2, 3, 4], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "39", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 1619491.0, "min": 1616589.0, "max": 1635509.0, "samples": 4}, "SQ_INSTS_SALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 4}, "SQ_INSTS_VALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 4}, "SQ_WAVES": {"median": 24576.0, "min": 24576.0, "max": 24576.0, "samples": 4}}, "trace_samples": 4, "trace_median_ms": 0.611023, "candidate_calls": [0, 1, 2, 3, 4], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "35", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 72783.0, "min": 69055.0, "max": 74766.0, "samples": 4}, "SQ_INSTS_SALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 4}, "SQ_INSTS_VALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 4}, "SQ_WAVES": {"median": 13824.0, "min": 13824.0, "max": 13824.0, "samples": 4}}, "trace_samples": 4, "trace_median_ms": 0.018799999999999997, "candidate_calls": [0, 1, 2, 3, 4], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "36", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 40518.0, "min": 38840.0, "max": 41352.0, "samples": 4}, "SQ_INSTS_SALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 4}, "SQ_INSTS_VALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 4}, "SQ_WAVES": {"median": 4608.0, "min": 4608.0, "max": 4608.0, "samples": 4}}, "trace_samples": 4, "trace_median_ms": 0.0074, "candidate_calls": [0, 1, 2, 3, 4], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "__amd_rocclr_copyBuffer", "kernel_id": "25", "dispatches": 2, "metrics": {"GRBM_GUI_ACTIVE": {"median": 44190.0, "min": 42097.0, "max": 46283.0, "samples": 2}, "SQ_INSTS_SALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 2}, "SQ_INSTS_VALU": {"median": 0.0, "min": 0.0, "max": 0.0, "samples": 2}, "SQ_WAVES": {"median": 24.0, "min": 16.0, "max": 32.0, "samples": 2}}, "trace_samples": 2, "trace_median_ms": 0.0038599999999999997, "candidate_calls": [], "association": "unmapped_or_setup"}
Full report: /home/xiaoxinyu/dsl-native-opt/runs/kb-ours-rocm-mingpt-combinations-astra-low-20260926-082611/sessions/1/.cache/diagnosis/report.json
