# Diagnostic report

2026-09-26 09:58:18  [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 09:58:18  [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 09:58:18  [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_095748_iter-2/output.txt", "valid": true, "reference_and_protocol_verified": true, "COMPILED": "True", "CORRECT": "True", "RUNTIME": "21.111492156982422", "REF_RUNTIME": "44.953800201416016", "SPEEDUP": "2.1293521020278985"}
HIP call-group costs are diagnostic, not per-device-kernel ranking or hardware floors.

Call 0: 13.119339 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 128, "function": "forward"} | Native: {"file": "kernel_0.cpp", "line": 1, "end_line": 71}
Tensor bindings: [{"value": "input_0", "shape": [128, 512, 768], "dtype": "float32"}, {"value": "input_1", "shape": [2304, 768], "dtype": "float32"}, {"value": "input_2", "shape": [2304], "dtype": "float32"}, {"value": "value_5", "shape": [128, 512, 2304], "dtype": "float16"}]
High source: 126:     def to(self,*args,**kwargs): return self
127:     def forward(self,x,qkv_weight,qkv_bias,projection_weight,projection_bias):
128:         q=self.qkv(x,qkv_weight,qkv_bias)
129:         y=self.attn(q)
130:         return self.proj(y,projection_weight,projection_bias)
Native sites: 10: extern "C" __global__ void __launch_bounds__(128) main_kernel(float* __restrict__ B, float* __restrict__ W, float* __restrict__ X, half_t* __restrict__ Y) {
11:   extern __shared__ __align__(1024) uchar buf_dyn_shmem[];
21:     __syncthreads();
42:       __syncthreads();
55:       *((tl_v8f32*)c + ((i_3 * 2) + j_1)) = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(

Call 2: 4.797291 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 130, "function": "forward"} | Native: {"file": "kernel_2.cpp", "line": 1, "end_line": 71}
Tensor bindings: [{"value": "value_6", "shape": [128, 512, 768], "dtype": "float32"}, {"value": "input_3", "shape": [768, 768], "dtype": "float32"}, {"value": "input_4", "shape": [768], "dtype": "float32"}, {"value": "value_7", "shape": [128, 512, 768], "dtype": "float32"}]
High source: 128:         q=self.qkv(x,qkv_weight,qkv_bias)
129:         y=self.attn(q)
130:         return self.proj(y,projection_weight,projection_bias)
131:     __call__=forward
Native sites: 10: extern "C" __global__ void __launch_bounds__(128) main_kernel(float* __restrict__ B, float* __restrict__ W, float* __restrict__ X, float* __restrict__ Y) {
11:   extern __shared__ __align__(1024) uchar buf_dyn_shmem[];
21:     __syncthreads();
42:       __syncthreads();
55:       *((tl_v8f32*)c + ((i_3 * 2) + j_1)) = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(

Call 1: 3.181594 ms | compiler_and_execution_matched
High: {"file": "solution/ModelNew.py", "line": 129, "function": "forward"} | Native: {"file": "kernel_1.cpp", "line": 1, "end_line": 210}
Tensor bindings: [{"value": "value_5", "shape": [128, 512, 2304], "dtype": "float16"}, {"value": "value_6", "shape": [128, 512, 768], "dtype": "float32"}]
High source: 127:     def forward(self,x,qkv_weight,qkv_bias,projection_weight,projection_bias):
128:         q=self.qkv(x,qkv_weight,qkv_bias)
129:         y=self.attn(q)
130:         return self.proj(y,projection_weight,projection_bias)
131:     __call__=forward
Native sites: 10: extern "C" __global__ void __launch_bounds__(128) main_kernel(half_t* __restrict__ Q, float* __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(
Dataflow: {"edges": [{"value": "value_5", "producer": 0, "consumer": 1}, {"value": "value_6", "producer": 1, "consumer": 2}], "bindings": [{"call": 0, "params": [{"value": "input_0", "shape": [128, 512, 768], "dtype": "float32"}, {"value": "input_1", "shape": [2304, 768], "dtype": "float32"}, {"value": "input_2", "shape": [2304], "dtype": "float32"}, {"value": "value_5", "shape": [128, 512, 2304], "dtype": "float16"}]}, {"call": 1, "params": [{"value": "value_5", "shape": [128, 512, 2304], "dtype": "float16"}, {"value": "value_6", "shape": [128, 512, 768], "dtype": "float32"}]}, {"call": 2, "params": [{"value": "value_6", "shape": [128, 512, 768], "dtype": "float32"}, {"value": "input_3", "shape": [768, 768], "dtype": "float32"}, {"value": "input_4", "shape": [768], "dtype": "float32"}, {"value": "value_7", "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": "35", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 41279124.0, "min": 39067523.0, "max": 42966863.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": 147456.0, "min": 147456.0, "max": 147968.0, "samples": 4}}, "trace_samples": 4, "trace_median_ms": 13.0697775, "candidate_calls": [0, 1, 2], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "37", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 15736668.0, "min": 13633749.0, "max": 15760840.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": 49152.0, "min": 49152.0, "max": 49152.0, "samples": 4}}, "trace_samples": 4, "trace_median_ms": 4.7845505, "candidate_calls": [0, 1, 2], "association": "ambiguous_symbol"}
Device kernel: {"symbol": "main_kernel", "kernel_id": "36", "dispatches": 4, "metrics": {"GRBM_GUI_ACTIVE": {"median": 10426505.5, "min": 10357565.0, "max": 10440848.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": 3.1745140000000003, "candidate_calls": [0, 1, 2], "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": "1", "dispatches": 1, "metrics": {"GRBM_GUI_ACTIVE": {"median": 4844488.0, "min": 4844488.0, "max": 4844488.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.711218, "candidate_calls": [], "association": "unmapped_or_setup"}
Device kernel: {"symbol": "__amd_rocclr_copyBuffer", "kernel_id": "21", "dispatches": 2, "metrics": {"GRBM_GUI_ACTIVE": {"median": 43508.0, "min": 41031.0, "max": 45985.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.00362, "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
