# Diagnostic report

2026-09-26 10:05:13  [TileLang:tilelang.cache.kernel_cache:WARNING] (kernel_cache.py:380): Found kernel 'main' in memory cache. For better performance, consider using `@tilelang.jit` instead of direct kernel caching.
2026-09-26 10:05:14  [TileLang:tilelang.cache.kernel_cache:WARNING] (kernel_cache.py:380): Found kernel 'main' in memory cache. For better performance, consider using `@tilelang.jit` instead of direct kernel caching.
2026-09-26 10:05:14  [TileLang:tilelang.cache.kernel_cache:WARNING] (kernel_cache.py:380): Found kernel 'main' in memory cache. For better performance, consider using `@tilelang.jit` instead of direct kernel caching.
2026-09-26 10:05:14  [TileLang:tilelang.cache.kernel_cache:WARNING] (kernel_cache.py:380): Found kernel 'main' in memory cache. For better performance, consider using `@tilelang.jit` instead of direct kernel caching.
2026-09-26 10:05:14  [TileLang:tilelang.cache.kernel_cache:WARNING] (kernel_cache.py:380): Found kernel 'main' in memory cache. For better performance, consider using `@tilelang.jit` instead of direct kernel caching.
2026-09-26 10:05:14  [TileLang:tilelang.cache.kernel_cache:WARNING] (kernel_cache.py:380): Found kernel 'main' in memory cache. For better performance, consider using `@tilelang.jit` instead of direct kernel caching.
Level: High | ranking measurement: {"path": "trajectory/20260926_100232_iter-8/output.txt", "valid": true, "reference_and_protocol_verified": true, "COMPILED": "True", "CORRECT": "True", "RUNTIME": "5.681180953979492", "REF_RUNTIME": "44.953800201416016", "SPEEDUP": "7.912756267678336"}
HIP call-group costs are diagnostic, not per-device-kernel ranking or hardware floors.

Call 4: 2.786430 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 95, "function": "forward"} | Native: {"file": "kernel_4.cpp", "line": 1, "end_line": 210}
Tensor bindings: [{"value": "value_8", "shape": [128, 512, 2304], "dtype": "float16"}, {"value": "value_9", "shape": [128, 512, 768], "dtype": "float16"}]
High source: 93:         ph=self.cp(projection_weight)
94:         q=self.qkv(xh,wh,qkv_bias)
95:         y=self.attn(q)
96:         return self.proj(y,ph,projection_bias)
97:     __call__=forward
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[];
46:     __syncthreads();
60:       __syncthreads();
70:       *((tl_v8f32*)s + j_1) = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(
99:       __syncthreads();
143:       __syncthreads();
175:     __syncthreads();
183:       __syncthreads();
195:       *((tl_v8f32*)o + j_3) = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(

Call 3: 1.758769 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 94, "function": "forward"} | Native: {"file": "kernel_3.cpp", "line": 1, "end_line": 63}
Tensor bindings: [{"value": "value_5", "shape": [128, 512, 768], "dtype": "float16"}, {"value": "value_6", "shape": [2304, 768], "dtype": "float16"}, {"value": "input_2", "shape": [2304], "dtype": "float32"}, {"value": "value_8", "shape": [128, 512, 2304], "dtype": "float16"}]
High source: 92:         wh=self.cw(qkv_weight)
93:         ph=self.cp(projection_weight)
94:         q=self.qkv(xh,wh,qkv_bias)
95:         y=self.attn(q)
96:         return self.proj(y,ph,projection_bias)
Native sites: 10: extern "C" __global__ void __launch_bounds__(256) main_kernel(float* __restrict__ B, half_t* __restrict__ W, half_t* __restrict__ X, half_t* __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 5: 0.624367 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 96, "function": "forward"} | Native: {"file": "kernel_5.cpp", "line": 1, "end_line": 63}
Tensor bindings: [{"value": "value_9", "shape": [128, 512, 768], "dtype": "float16"}, {"value": "value_7", "shape": [768, 768], "dtype": "float16"}, {"value": "input_4", "shape": [768], "dtype": "float32"}, {"value": "value_10", "shape": [128, 512, 768], "dtype": "float32"}]
High source: 94:         q=self.qkv(xh,wh,qkv_bias)
95:         y=self.attn(q)
96:         return self.proj(y,ph,projection_bias)
97:     __call__=forward
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.543246 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 91, "function": "forward"} | Native: {"file": "kernel_0.cpp", "line": 1, "end_line": 26}
Tensor bindings: [{"value": "input_0", "shape": [128, 512, 768], "dtype": "float32"}, {"value": "value_5", "shape": [128, 512, 768], "dtype": "float16"}]
High source: 89:     def to(self,*args,**kwargs): return self
90:     def forward(self,x,qkv_weight,qkv_bias,projection_weight,projection_bias):
91:         xh=self.cx(x)
92:         wh=self.cw(qkv_weight)
93:         ph=self.cp(projection_weight)
Native sites: 10: extern "C" __global__ void __launch_bounds__(256) main_kernel(float* __restrict__ X, half_t* __restrict__ Y) {

Call 2: 0.039100 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 93, "function": "forward"} | Native: {"file": "kernel_2.cpp", "line": 1, "end_line": 26}
Tensor bindings: [{"value": "input_3", "shape": [768, 768], "dtype": "float32"}, {"value": "value_7", "shape": [768, 768], "dtype": "float16"}]
High source: 91:         xh=self.cx(x)
92:         wh=self.cw(qkv_weight)
93:         ph=self.cp(projection_weight)
94:         q=self.qkv(xh,wh,qkv_bias)
95:         y=self.attn(q)
Native sites: 10: extern "C" __global__ void __launch_bounds__(256) main_kernel(float* __restrict__ X, half_t* __restrict__ Y) {

Call 1: 0.022840 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 92, "function": "forward"} | Native: {"file": "kernel_1.cpp", "line": 1, "end_line": 26}
Tensor bindings: [{"value": "input_1", "shape": [2304, 768], "dtype": "float32"}, {"value": "value_6", "shape": [2304, 768], "dtype": "float16"}]
High source: 90:     def forward(self,x,qkv_weight,qkv_bias,projection_weight,projection_bias):
91:         xh=self.cx(x)
92:         wh=self.cw(qkv_weight)
93:         ph=self.cp(projection_weight)
94:         q=self.qkv(xh,wh,qkv_bias)
Native sites: 10: extern "C" __global__ void __launch_bounds__(256) main_kernel(float* __restrict__ X, half_t* __restrict__ Y) {
Dataflow: {"edges": [{"value": "value_5", "producer": 0, "consumer": 3}, {"value": "value_6", "producer": 1, "consumer": 3}, {"value": "value_8", "producer": 3, "consumer": 4}, {"value": "value_9", "producer": 4, "consumer": 5}, {"value": "value_7", "producer": 2, "consumer": 5}], "bindings": [{"call": 0, "params": [{"value": "input_0", "shape": [128, 512, 768], "dtype": "float32"}, {"value": "value_5", "shape": [128, 512, 768], "dtype": "float16"}]}, {"call": 1, "params": [{"value": "input_1", "shape": [2304, 768], "dtype": "float32"}, {"value": "value_6", "shape": [2304, 768], "dtype": "float16"}]}, {"call": 2, "params": [{"value": "input_3", "shape": [768, 768], "dtype": "float32"}, {"value": "value_7", "shape": [768, 768], "dtype": "float16"}]}, {"call": 3, "params": [{"value": "value_5", "shape": [128, 512, 768], "dtype": "float16"}, {"value": "value_6", "shape": [2304, 768], "dtype": "float16"}, {"value": "input_2", "shape": [2304], "dtype": "float32"}, {"value": "value_8", "shape": [128, 512, 2304], "dtype": "float16"}]}, {"call": 4, "params": [{"value": "value_8", "shape": [128, 512, 2304], "dtype": "float16"}, {"value": "value_9", "shape": [128, 512, 768], "dtype": "float16"}]}, {"call": 5, "params": [{"value": "value_9", "shape": [128, 512, 768], "dtype": "float16"}, {"value": "value_7", "shape": [768, 768], "dtype": "float16"}, {"value": "input_4", "shape": [768], "dtype": "float32"}, {"value": "value_10", "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": "main_kernel", "kernel_id": "39", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 8091358.0, "min": 8034585.0, "max": 8121671.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": 2.81965, "candidate_calls": [0, 1, 2, 3, 4, 5], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "38", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 4649661.5, "min": 4644120.0, "max": 4665699.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": 73728.0, "min": 73728.0, "max": 74240.0, "samples": 4}}, "trace_samples": 4, "trace_median_ms": 1.804399, "candidate_calls": [0, 1, 2, 3, 4, 5], "association": "ambiguous_symbol"}
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": "4", "dispatches": 1, "metrics": {"GRBM_GUI_ACTIVE": {"median": 4334802.0, "min": 4334802.0, "max": 4334802.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": 1.530456, "candidate_calls": [], "association": "unmapped_or_setup"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "40", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 1617338.5, "min": 1601250.0, "max": 1623234.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.6009260000000001, "candidate_calls": [0, 1, 2, 3, 4, 5], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "35", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 1540571.0, "min": 1532063.0, "max": 2387041.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": 393216.0, "min": 393216.0, "max": 393472.0, "samples": 4}}, "trace_samples": 4, "trace_median_ms": 0.5301855, "candidate_calls": [0, 1, 2, 3, 4, 5], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "36", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 61713.0, "min": 61339.0, "max": 84065.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.01512, "candidate_calls": [0, 1, 2, 3, 4, 5], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "37", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 35171.0, "min": 34982.0, "max": 36972.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.006, "candidate_calls": [0, 1, 2, 3, 4, 5], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "__amd_rocclr_copyBuffer", "kernel_id": "21", "dispatches": 2, "metrics": {"GRBM_GUI_ACTIVE": {"median": 43856.5, "min": 41070.0, "max": 46643.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.00364, "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
