# Diagnostic report

2026-09-26 10:41:55  [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:41:55  [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:41:55  [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:41:55  [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:41:55  [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_104102_iter-43/output.txt", "valid": true, "reference_and_protocol_verified": true, "COMPILED": "True", "CORRECT": "True", "RUNTIME": "5.529796123504639", "REF_RUNTIME": "44.953800201416016", "SPEEDUP": "8.129377502786754"}
HIP call-group costs are diagnostic, not per-device-kernel ranking or hardware floors.

Call 3: 2.893782 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 98, "function": "forward"} | Native: {"file": "kernel_3.cpp", "line": 1, "end_line": 210}
Tensor bindings: [{"value": "value_7", "shape": [3, 128, 8, 512, 96], "dtype": "float16"}, {"value": "value_8", "shape": [128, 512, 768], "dtype": "float16"}]
High source: 96:         ph=self.cp(projection_weight)
97:         q=self.qkv(x,wh,qkv_bias)
98:         y=self.attn(q)
99:         return self.proj(y,ph,projection_bias)
100:     __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 2: 1.845769 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 97, "function": "forward"} | Native: {"file": "kernel_2.cpp", "line": 1, "end_line": 73}
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": [3, 128, 8, 512, 96], "dtype": "float16"}]
High source: 95:         wh=self.cw(qkv_weight)
96:         ph=self.cp(projection_weight)
97:         q=self.qkv(x,wh,qkv_bias)
98:         y=self.attn(q)
99:         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, 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 4: 0.626533 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 99, "function": "forward"} | Native: {"file": "kernel_4.cpp", "line": 1, "end_line": 63}
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"}]
High source: 97:         q=self.qkv(x,wh,qkv_bias)
98:         y=self.attn(q)
99:         return self.proj(y,ph,projection_bias)
100:     __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 1: 0.067081 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 96, "function": "forward"} | Native: {"file": "kernel_1.cpp", "line": 1, "end_line": 26}
Tensor bindings: [{"value": "input_3", "shape": [768, 768], "dtype": "float32"}, {"value": "value_6", "shape": [768, 768], "dtype": "float16"}]
High source: 94:     def forward(self,x,qkv_weight,qkv_bias,projection_weight,projection_bias):
95:         wh=self.cw(qkv_weight)
96:         ph=self.cp(projection_weight)
97:         q=self.qkv(x,wh,qkv_bias)
98:         y=self.attn(q)
Native sites: 10: extern "C" __global__ void __launch_bounds__(256) main_kernel(float* __restrict__ X, half_t* __restrict__ Y) {

Call 0: 0.027400 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 95, "function": "forward"} | Native: {"file": "kernel_0.cpp", "line": 1, "end_line": 26}
Tensor bindings: [{"value": "input_1", "shape": [2304, 768], "dtype": "float32"}, {"value": "value_5", "shape": [2304, 768], "dtype": "float16"}]
High source: 93:     def to(self,*args,**kwargs): return self
94:     def forward(self,x,qkv_weight,qkv_bias,projection_weight,projection_bias):
95:         wh=self.cw(qkv_weight)
96:         ph=self.cp(projection_weight)
97:         q=self.qkv(x,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": 2}, {"value": "value_7", "producer": 2, "consumer": 3}, {"value": "value_8", "producer": 3, "consumer": 4}, {"value": "value_6", "producer": 1, "consumer": 4}], "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": [3, 128, 8, 512, 96], "dtype": "float16"}]}, {"call": 3, "params": [{"value": "value_7", "shape": [3, 128, 8, 512, 96], "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": "main_kernel", "kernel_id": "38", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 8410761.5, "min": 8344961.0, "max": 8933370.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.943923, "candidate_calls": [0, 1, 2, 3, 4], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "37", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 4747939.5, "min": 4738085.0, "max": 5608954.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.8761800000000002, "candidate_calls": [0, 1, 2, 3, 4], "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": "7", "dispatches": 1, "metrics": {"GRBM_GUI_ACTIVE": {"median": 5143815.0, "min": 5143815.0, "max": 5143815.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.816559, "candidate_calls": [], "association": "unmapped_or_setup"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "39", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 1618637.5, "min": 1616002.0, "max": 1625071.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.605512, "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": 71432.0, "min": 67613.0, "max": 71976.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.018520000000000002, "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": 39215.0, "min": 36789.0, "max": 39999.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.0074405, "candidate_calls": [0, 1, 2, 3, 4], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "__amd_rocclr_copyBuffer", "kernel_id": "21", "dispatches": 2, "metrics": {"GRBM_GUI_ACTIVE": {"median": 44134.0, "min": 40851.0, "max": 47417.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.00366, "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
