The LinearAttention output of tilelang are inconsistent with the FLA baseline

[cailun01]

Hi, TileLang team!

I run examples/linear_attention/example_linear_attn_fwd.py and examples/linear_attention/example_linear_attn_bwd.py and got:

Failed some tests!❌
Triton latency: 3.718 ms
TileLang latency: 2.905 ms
Speedup: 1.280x

Failed some tests!❌
Triton latency: 11.189 ms
TileLang latency: 8.127 ms
Speedup: 1.377x

torch.allclose(o, o_ref) in example_linear_attn_fwd.py and torch.allclose(dq, q.grad), torch.allclose(dk, k.grad), torch.allclose(dv, v.grad) in example_linear_attn_bwd.py return False.

Could you please share some troubleshooting ideas?

Testbed setup

I built tilelang from source by method 3

tilelang version: 8fe3540
flash-linear-attention version: 195b74db140bada1ebf5a7bac8b11af0faf4ca59

PyTorch version: 2.8.0a0+5228986c39.nv25.06
Is debug build: False
CUDA used to build PyTorch: 12.9
ROCM used to build PyTorch: N/A

OS: Ubuntu 24.04.2 LTS (aarch64)
GCC version: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0
Clang version: Could not collect
CMake version: version 4.0.3
Libc version: glibc-2.39

Python version: 3.12.3 (main, Feb  4 2025, 14:48:35) [GCC 13.3.0] (64-bit runtime)
Python platform: Linux-6.6.47-002.ant8.aarch64-aarch64-with-glibc2.39
Is CUDA available: True
CUDA runtime version: 12.9.86
CUDA_MODULE_LOADING set to: LAZY
GPU models and configuration: 
GPU 0: NVIDIA GB200
GPU 1: NVIDIA GB200
GPU 2: NVIDIA GB200
GPU 3: NVIDIA GB200

Nvidia driver version: 580.65.06
cuDNN version: Probably one of the following:
/usr/lib/aarch64-linux-gnu/libcudnn.so.9.10.2
/usr/lib/aarch64-linux-gnu/libcudnn_adv.so.9.10.2
/usr/lib/aarch64-linux-gnu/libcudnn_cnn.so.9.10.2
/usr/lib/aarch64-linux-gnu/libcudnn_engines_precompiled.so.9.10.2
/usr/lib/aarch64-linux-gnu/libcudnn_engines_runtime_compiled.so.9.10.2
/usr/lib/aarch64-linux-gnu/libcudnn_graph.so.9.10.2
/usr/lib/aarch64-linux-gnu/libcudnn_heuristic.so.9.10.2
/usr/lib/aarch64-linux-gnu/libcudnn_ops.so.9.10.2
HIP runtime version: N/A
MIOpen runtime version: N/A
Is XNNPACK available: True

CPU:
Architecture:                         aarch64
CPU op-mode(s):                       64-bit
Byte Order:                           Little Endian
CPU(s):                               144
On-line CPU(s) list:                  0-143
Vendor ID:                            ARM
Model name:                           Neoverse-V2
Model:                                0
Thread(s) per core:                   1
Core(s) per cluster:                  72
Socket(s):                            -
Cluster(s):                           2
Stepping:                             r0p0
Frequency boost:                      disabled
CPU(s) scaling MHz:                   95%
CPU max MHz:                          3411.0000
CPU min MHz:                          81.0000
BogoMIPS:                             2000.00
Flags:                                fp asimd evtstrm aes pmull sha1 sha2 crc32 atomics fphp asimdhp cpuid asimdrdm jscvt fcma lrcpc dcpop sha3 sm3 sm4 asimddp sha512 sve asimdfhm dit uscat ilrcpc flagm sb dcpodp sve2 sveaes svepmull svebitperm svesha3 svesm4 flagm2 frint svei8mm svebf16 i8mm bf16 dgh
L1d cache:                            9 MiB (144 instances)
L1i cache:                            9 MiB (144 instances)
L2 cache:                             144 MiB (144 instances)
L3 cache:                             228 MiB (2 instances)
NUMA node(s):                         34
NUMA node0 CPU(s):                    0-71
NUMA node1 CPU(s):                    72-143
NUMA node2 CPU(s):                    
NUMA node3 CPU(s):                    
NUMA node4 CPU(s):                    
NUMA node5 CPU(s):                    
NUMA node6 CPU(s):                    
NUMA node7 CPU(s):                    
NUMA node8 CPU(s):                    
NUMA node9 CPU(s):                    
NUMA node10 CPU(s):                   
NUMA node11 CPU(s):                   
NUMA node12 CPU(s):                   
NUMA node13 CPU(s):                   
NUMA node14 CPU(s):                   
NUMA node15 CPU(s):                   
NUMA node16 CPU(s):                   
NUMA node17 CPU(s):                   
NUMA node18 CPU(s):                   
NUMA node19 CPU(s):                   
NUMA node20 CPU(s):                   
NUMA node21 CPU(s):                   
NUMA node22 CPU(s):                   
NUMA node23 CPU(s):                   
NUMA node24 CPU(s):                   
NUMA node25 CPU(s):                   
NUMA node26 CPU(s):                   
NUMA node27 CPU(s):                   
NUMA node28 CPU(s):                   
NUMA node29 CPU(s):                   
NUMA node30 CPU(s):                   
NUMA node31 CPU(s):                   
NUMA node32 CPU(s):                   
NUMA node33 CPU(s):                   
Vulnerability Gather data sampling:   Not affected
Vulnerability Itlb multihit:          Not affected
Vulnerability L1tf:                   Not affected
Vulnerability Mds:                    Not affected
Vulnerability Meltdown:               Not affected
Vulnerability Mmio stale data:        Not affected
Vulnerability Reg file data sampling: Not affected
Vulnerability Retbleed:               Not affected
Vulnerability Spec rstack overflow:   Not affected
Vulnerability Spec store bypass:      Not affected
Vulnerability Spectre v1:             Mitigation; __user pointer sanitization
Vulnerability Spectre v2:             Not affected
Vulnerability Srbds:                  Not affected
Vulnerability Tsx async abort:        Not affected

Versions of relevant libraries:
[pip3] flake8==7.3.0
[pip3] mypy_extensions==1.1.0
[pip3] numpy==1.26.4
[pip3] nvidia-cublas-cu12==12.9.1.4
[pip3] nvidia-cuda-cupti-cu12==12.9.79
[pip3] nvidia-cuda-nvrtc-cu12==12.9.86
[pip3] nvidia-cuda-runtime-cu12==12.9.79
[pip3] nvidia-cudnn-cu12==9.11.0.98
[pip3] nvidia-cudnn-frontend==1.12.0
[pip3] nvidia-cufft-cu12==11.4.1.4
[pip3] nvidia-curand-cu12==10.3.10.19
[pip3] nvidia-cusolver-cu12==11.7.5.82
[pip3] nvidia-cusparse-cu12==12.5.10.65
[pip3] nvidia-cusparselt-cu12==0.7.1
[pip3] nvidia-nccl-cu12==2.27.6
[pip3] nvidia-nvjitlink-cu12==12.9.86
[pip3] nvidia-nvtx-cu12==12.9.79
[pip3] nvtx==0.2.11
[pip3] onnx==1.17.0
[pip3] onnx-ir==0.1.4
[pip3] onnxscript==0.3.1
[pip3] optree==0.16.0
[pip3] pynvjitlink==0.3.0
[pip3] pytorch-triton==3.3.0+git96316ce52.nvinternal
[pip3] torch==2.8.0a0+5228986c39.nv25.6
[pip3] torch_tensorrt==2.8.0a0
[pip3] torchao==0.11.0+git
[pip3] torchprofile==0.0.4
[pip3] torchvision==0.22.0a0+95f10a4e

[LeiWang1999]

cc @tzj-fxz @Rachmanino

[Rachmanino]

Hi @cailun01, thanks for your report. I'll take a look.

[cailun01]

Here is the generated code:

// modified by tilelang_callback_cuda_postproc
#include <tl_templates/cuda/gemm.h>
#include <tl_templates/cuda/copy.h>
#include <tl_templates/cuda/reduce.h>
#include <tl_templates/cuda/ldsm.h>
#include <tl_templates/cuda/threadblock_swizzle.h>
#include <tl_templates/cuda/debug.h>
#ifdef ENABLE_BF16
#include <tl_templates/cuda/cuda_bf16_fallbacks.cuh>
#endif

extern "C" __global__ void chunk_linear_attn_fwd_kernel(half_t* __restrict__ K, half_t* __restrict__ O, half_t* __restrict__ Q, half_t* __restrict__ V, float* __restrict__ final_state);
extern "C" __global__ void __launch_bounds__(128, 1) chunk_linear_attn_fwd_kernel(half_t* __restrict__ K, half_t* __restrict__ O, half_t* __restrict__ Q, half_t* __restrict__ V, float* __restrict__ final_state) {
  extern __shared__ __align__(1024) uchar buf_dyn_shmem[];
  float h[32];
  float s[32];
  float o[32];
  #pragma unroll
  for (int i = 0; i < 16; ++i) {
    *(float2*)(h + (i * 2)) = make_float2(0x0p+0f/*0.000000e+00*/, 0x0p+0f/*0.000000e+00*/);
  }
  const dim3 blockIdx = tl::rasterization2DRow<10>();
  #pragma unroll
  for (int i_1 = 0; i_1 < 4; ++i_1) {
    tl::cp_async_gs<16>(buf_dyn_shmem+(((((i_1 * 2048) + ((((int)threadIdx.x) >> 3) * 128)) + (((((((int)threadIdx.x) & 63) >> 5) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 64)) + (((((((int)threadIdx.x) & 31) >> 4) + ((((int)threadIdx.x) & 3) >> 1)) & 1) * 32)) + (((((((int)threadIdx.x) & 15) >> 3) + (((int)threadIdx.x) & 1)) & 1) * 16)), K+(((((((((int)blockIdx.z) >> 5) * 33554432) + (i_1 * 131072)) + ((((int)threadIdx.x) >> 3) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.y) * 64)) + ((((int)threadIdx.x) & 7) * 8)));
  }
  #pragma unroll
  for (int i_2 = 0; i_2 < 4; ++i_2) {
    tl::cp_async_gs<16>(buf_dyn_shmem+((((((i_2 * 2048) + ((((int)threadIdx.x) >> 3) * 128)) + (((((((int)threadIdx.x) & 63) >> 5) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 64)) + (((((((int)threadIdx.x) & 31) >> 4) + ((((int)threadIdx.x) & 3) >> 1)) & 1) * 32)) + (((((((int)threadIdx.x) & 15) >> 3) + (((int)threadIdx.x) & 1)) & 1) * 16)) + 16384), V+(((((((((int)blockIdx.z) >> 5) * 33554432) + (i_2 * 131072)) + ((((int)threadIdx.x) >> 3) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.x) * 64)) + ((((int)threadIdx.x) & 7) * 8)));
  }
  tl::cp_async_commit();
  #pragma unroll
  for (int i_3 = 0; i_3 < 4; ++i_3) {
    uint4 __1;
    ulonglong4 __2;
      ulonglong4 __3;
      uint4 v_ = *(uint4*)(Q + (((((((((int)blockIdx.z) >> 5) * 33554432) + (i_3 * 131072)) + ((((int)threadIdx.x) >> 3) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.y) * 64)) + ((((int)threadIdx.x) & 7) * 8)));
      ((float2*)(&(__3.x)))->x = (float)(((half2*)(&(v_.x)))->x);
      ((float2*)(&(__3.x)))->y = (float)(((half2*)(&(v_.x)))->y);
      ((float2*)(&(__3.y)))->x = (float)(((half2*)(&(v_.y)))->x);
      ((float2*)(&(__3.y)))->y = (float)(((half2*)(&(v_.y)))->y);
      ((float2*)(&(__3.z)))->x = (float)(((half2*)(&(v_.z)))->x);
      ((float2*)(&(__3.z)))->y = (float)(((half2*)(&(v_.z)))->y);
      ((float2*)(&(__3.w)))->x = (float)(((half2*)(&(v_.w)))->x);
      ((float2*)(&(__3.w)))->y = (float)(((half2*)(&(v_.w)))->y);
      ulonglong4 v__1 = make_ulonglong4(*(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/), *(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/), *(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/), *(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/));
      ((float2*)(&(__2.x)))->x = (((float2*)(&(__3.x)))->x*((float2*)(&(v__1.x)))->x);
      ((float2*)(&(__2.x)))->y = (((float2*)(&(__3.x)))->y*((float2*)(&(v__1.x)))->y);
      ((float2*)(&(__2.y)))->x = (((float2*)(&(__3.y)))->x*((float2*)(&(v__1.y)))->x);
      ((float2*)(&(__2.y)))->y = (((float2*)(&(__3.y)))->y*((float2*)(&(v__1.y)))->y);
      ((float2*)(&(__2.z)))->x = (((float2*)(&(__3.z)))->x*((float2*)(&(v__1.z)))->x);
      ((float2*)(&(__2.z)))->y = (((float2*)(&(__3.z)))->y*((float2*)(&(v__1.z)))->y);
      ((float2*)(&(__2.w)))->x = (((float2*)(&(__3.w)))->x*((float2*)(&(v__1.w)))->x);
      ((float2*)(&(__2.w)))->y = (((float2*)(&(__3.w)))->y*((float2*)(&(v__1.w)))->y);
    ((half2*)(&(__1.x)))->x = (half_t)(((float2*)(&(__2.x)))->x);
    ((half2*)(&(__1.x)))->y = (half_t)(((float2*)(&(__2.x)))->y);
    ((half2*)(&(__1.y)))->x = (half_t)(((float2*)(&(__2.y)))->x);
    ((half2*)(&(__1.y)))->y = (half_t)(((float2*)(&(__2.y)))->y);
    ((half2*)(&(__1.z)))->x = (half_t)(((float2*)(&(__2.z)))->x);
    ((half2*)(&(__1.z)))->y = (half_t)(((float2*)(&(__2.z)))->y);
    ((half2*)(&(__1.w)))->x = (half_t)(((float2*)(&(__2.w)))->x);
    ((half2*)(&(__1.w)))->y = (half_t)(((float2*)(&(__2.w)))->y);
    *(uint4*)(((half_t*)buf_dyn_shmem) + ((((((i_3 * 1024) + ((((int)threadIdx.x) >> 3) * 64)) + (((((((int)threadIdx.x) & 63) >> 5) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 32)) + (((((((int)threadIdx.x) & 31) >> 4) + ((((int)threadIdx.x) & 3) >> 1)) & 1) * 16)) + (((((((int)threadIdx.x) & 15) >> 3) + (((int)threadIdx.x) & 1)) & 1) * 8)) + 16384)) = __1;
  }
  tl::cp_async_commit();
  #pragma unroll
  for (int i_4 = 0; i_4 < 4; ++i_4) {
    tl::cp_async_gs<16>(buf_dyn_shmem+((((((i_4 * 2048) + ((((int)threadIdx.x) >> 3) * 128)) + (((((((int)threadIdx.x) & 63) >> 5) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 64)) + (((((((int)threadIdx.x) & 31) >> 4) + ((((int)threadIdx.x) & 3) >> 1)) & 1) * 32)) + (((((((int)threadIdx.x) & 15) >> 3) + (((int)threadIdx.x) & 1)) & 1) * 16)) + 8192), K+((((((((((int)blockIdx.z) >> 5) * 33554432) + (i_4 * 131072)) + ((((int)threadIdx.x) >> 3) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.y) * 64)) + ((((int)threadIdx.x) & 7) * 8)) + 524288));
  }
  #pragma unroll
  for (int i_5 = 0; i_5 < 4; ++i_5) {
    tl::cp_async_gs<16>(buf_dyn_shmem+((((((i_5 * 2048) + ((((int)threadIdx.x) >> 3) * 128)) + (((((((int)threadIdx.x) & 63) >> 5) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 64)) + (((((((int)threadIdx.x) & 31) >> 4) + ((((int)threadIdx.x) & 3) >> 1)) & 1) * 32)) + (((((((int)threadIdx.x) & 15) >> 3) + (((int)threadIdx.x) & 1)) & 1) * 16)) + 24576), V+((((((((((int)blockIdx.z) >> 5) * 33554432) + (i_5 * 131072)) + ((((int)threadIdx.x) >> 3) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.x) * 64)) + ((((int)threadIdx.x) & 7) * 8)) + 524288));
  }
  tl::cp_async_commit();
  #pragma unroll
  for (int i_6 = 0; i_6 < 4; ++i_6) {
    uint4 __4;
    ulonglong4 __5;
      ulonglong4 __6;
      uint4 v__2 = *(uint4*)(Q + ((((((((((int)blockIdx.z) >> 5) * 33554432) + (i_6 * 131072)) + ((((int)threadIdx.x) >> 3) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.y) * 64)) + ((((int)threadIdx.x) & 7) * 8)) + 524288));
      ((float2*)(&(__6.x)))->x = (float)(((half2*)(&(v__2.x)))->x);
      ((float2*)(&(__6.x)))->y = (float)(((half2*)(&(v__2.x)))->y);
      ((float2*)(&(__6.y)))->x = (float)(((half2*)(&(v__2.y)))->x);
      ((float2*)(&(__6.y)))->y = (float)(((half2*)(&(v__2.y)))->y);
      ((float2*)(&(__6.z)))->x = (float)(((half2*)(&(v__2.z)))->x);
      ((float2*)(&(__6.z)))->y = (float)(((half2*)(&(v__2.z)))->y);
      ((float2*)(&(__6.w)))->x = (float)(((half2*)(&(v__2.w)))->x);
      ((float2*)(&(__6.w)))->y = (float)(((half2*)(&(v__2.w)))->y);
      ulonglong4 v__3 = make_ulonglong4(*(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/), *(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/), *(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/), *(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/));
      ((float2*)(&(__5.x)))->x = (((float2*)(&(__6.x)))->x*((float2*)(&(v__3.x)))->x);
      ((float2*)(&(__5.x)))->y = (((float2*)(&(__6.x)))->y*((float2*)(&(v__3.x)))->y);
      ((float2*)(&(__5.y)))->x = (((float2*)(&(__6.y)))->x*((float2*)(&(v__3.y)))->x);
      ((float2*)(&(__5.y)))->y = (((float2*)(&(__6.y)))->y*((float2*)(&(v__3.y)))->y);
      ((float2*)(&(__5.z)))->x = (((float2*)(&(__6.z)))->x*((float2*)(&(v__3.z)))->x);
      ((float2*)(&(__5.z)))->y = (((float2*)(&(__6.z)))->y*((float2*)(&(v__3.z)))->y);
      ((float2*)(&(__5.w)))->x = (((float2*)(&(__6.w)))->x*((float2*)(&(v__3.w)))->x);
      ((float2*)(&(__5.w)))->y = (((float2*)(&(__6.w)))->y*((float2*)(&(v__3.w)))->y);
    ((half2*)(&(__4.x)))->x = (half_t)(((float2*)(&(__5.x)))->x);
    ((half2*)(&(__4.x)))->y = (half_t)(((float2*)(&(__5.x)))->y);
    ((half2*)(&(__4.y)))->x = (half_t)(((float2*)(&(__5.y)))->x);
    ((half2*)(&(__4.y)))->y = (half_t)(((float2*)(&(__5.y)))->y);
    ((half2*)(&(__4.z)))->x = (half_t)(((float2*)(&(__5.z)))->x);
    ((half2*)(&(__4.z)))->y = (half_t)(((float2*)(&(__5.z)))->y);
    ((half2*)(&(__4.w)))->x = (half_t)(((float2*)(&(__5.w)))->x);
    ((half2*)(&(__4.w)))->y = (half_t)(((float2*)(&(__5.w)))->y);
    *(uint4*)(((half_t*)buf_dyn_shmem) + ((((((i_6 * 1024) + ((((int)threadIdx.x) >> 3) * 64)) + (((((((int)threadIdx.x) & 63) >> 5) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 32)) + (((((((int)threadIdx.x) & 31) >> 4) + ((((int)threadIdx.x) & 3) >> 1)) & 1) * 16)) + (((((((int)threadIdx.x) & 15) >> 3) + (((int)threadIdx.x) & 1)) & 1) * 8)) + 20480)) = __4;
  }
  tl::cp_async_commit();
  for (int i_7 = 0; i_7 < 62; ++i_7) {
    tl::cp_async_wait<2>();
    __syncthreads();
    tl::gemm_ss<64, 64, 64, 2, 2, 0, 1, 1, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[(((i_7 & 1) * 4096) + 16384)])), (&(((half_t*)buf_dyn_shmem)[((i_7 & 1) * 4096)])), (&(s[0])));
    __syncthreads();
    #pragma unroll
    for (int i_8 = 0; i_8 < 16; ++i_8) {
      for (int vec_s = 0; vec_s < 2; ++vec_s) {
        float condval;
        if (((((((i_8 >> 2) * 16) + ((((int)threadIdx.x) >> 6) * 8)) + ((((int)threadIdx.x) & 3) * 2)) + vec_s) <= ((((((i_8 & 3) >> 1) * 32) + (((((int)threadIdx.x) & 63) >> 5) * 16)) + ((i_8 & 1) * 8)) + ((((int)threadIdx.x) & 31) >> 2)))) {
          condval = s[((i_8 * 2) + vec_s)];
        } else {
          condval = 0x0p+0f/*0.000000e+00*/;
        }
        ((half_t*)buf_dyn_shmem)[((((((((((((i_8 & 3) >> 1) * 2048) + (((((int)threadIdx.x) & 63) >> 5) * 1024)) + ((i_8 & 1) * 512)) + (((((int)threadIdx.x) & 31) >> 2) * 64)) + (((((((((i_8 >> 2) * 16) + ((((int)threadIdx.x) >> 6) * 8)) + ((((int)threadIdx.x) & 3) * 2)) + vec_s) >> 5) + ((((int)threadIdx.x) & 31) >> 4)) & 1) * 32)) + (((((((int)threadIdx.x) & 15) >> 3) + ((i_8 & 7) >> 2)) & 1) * 16)) + ((((((int)threadIdx.x) >> 6) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 8)) + ((((int)threadIdx.x) & 3) * 2)) + vec_s) + 24576)] = ((half_t)condval);
      }
    }
    tl::fence_proxy_async();
    tl::cp_async_wait<3>();
    __syncthreads();
    tl::gemm_ss<64, 64, 64, 2, 2, 0, 0, 1, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[24576])), (&(((half_t*)buf_dyn_shmem)[(((i_7 & 1) * 4096) + 8192)])), (&(o[0])));
    __syncthreads();
    #pragma unroll
    for (int i_9 = 0; i_9 < 4; ++i_9) {
      tl::ptx_stmatrix_x4((&(((half_t*)buf_dyn_shmem)[(((((((((((int)threadIdx.x) & 31) >> 4) * 2048) + (((((int)threadIdx.x) & 63) >> 5) * 1024)) + ((((int)threadIdx.x) & 15) * 64)) + (((((((int)threadIdx.x) & 7) >> 2) + (i_9 >> 1)) & 1) * 32)) + (((((((int)threadIdx.x) & 3) >> 1) + (i_9 & 1)) & 1) * 16)) + ((((((int)threadIdx.x) >> 6) + (((int)threadIdx.x) & 1)) & 1) * 8)) + 28672)])), __pack_half2(((half_t)h[(i_9 * 8)]), ((half_t)h[((i_9 * 8) + 1)])), __pack_half2(((half_t)h[((i_9 * 8) + 2)]), ((half_t)h[((i_9 * 8) + 3)])), __pack_half2(((half_t)h[((i_9 * 8) + 4)]), ((half_t)h[((i_9 * 8) + 5)])), __pack_half2(((half_t)h[((i_9 * 8) + 6)]), ((half_t)h[((i_9 * 8) + 7)])));
    }
    tl::fence_proxy_async();
    tl::cp_async_wait<3>();
    __syncthreads();
    tl::gemm_ss<64, 64, 64, 2, 2, 1, 0, 0, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[((i_7 & 1) * 4096)])), (&(((half_t*)buf_dyn_shmem)[(((i_7 & 1) * 4096) + 8192)])), (&(h[0])));
    __syncthreads();
    #pragma unroll
    for (int i_10 = 0; i_10 < 4; ++i_10) {
      tl::cp_async_gs<16>(buf_dyn_shmem+(((((((i_7 & 1) * 8192) + (i_10 * 2048)) + ((((int)threadIdx.x) >> 3) * 128)) + (((((((int)threadIdx.x) & 63) >> 5) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 64)) + (((((((int)threadIdx.x) & 31) >> 4) + ((((int)threadIdx.x) & 3) >> 1)) & 1) * 32)) + (((((((int)threadIdx.x) & 15) >> 3) + (((int)threadIdx.x) & 1)) & 1) * 16)), K+(((((((((((int)blockIdx.z) >> 5) * 33554432) + (i_7 * 524288)) + (i_10 * 131072)) + ((((int)threadIdx.x) >> 3) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.y) * 64)) + ((((int)threadIdx.x) & 7) * 8)) + 1048576));
    }
    #pragma unroll
    for (int i_11 = 0; i_11 < 4; ++i_11) {
      tl::cp_async_gs<16>(buf_dyn_shmem+((((((((i_7 & 1) * 8192) + (i_11 * 2048)) + ((((int)threadIdx.x) >> 3) * 128)) + (((((((int)threadIdx.x) & 63) >> 5) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 64)) + (((((((int)threadIdx.x) & 31) >> 4) + ((((int)threadIdx.x) & 3) >> 1)) & 1) * 32)) + (((((((int)threadIdx.x) & 15) >> 3) + (((int)threadIdx.x) & 1)) & 1) * 16)) + 16384), V+(((((((((((int)blockIdx.z) >> 5) * 33554432) + (i_7 * 524288)) + (i_11 * 131072)) + ((((int)threadIdx.x) >> 3) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.x) * 64)) + ((((int)threadIdx.x) & 7) * 8)) + 1048576));
    }
    tl::cp_async_commit();
    tl::fence_proxy_async();
    tl::cp_async_wait<3>();
    __syncthreads();
    tl::gemm_ss<64, 64, 64, 2, 2, 0, 0, 0, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[(((i_7 & 1) * 4096) + 16384)])), (&(((half_t*)buf_dyn_shmem)[28672])), (&(o[0])));
    __syncthreads();
    #pragma unroll
    for (int i_12 = 0; i_12 < 4; ++i_12) {
      uint4 __7;
      ulonglong4 __8;
        ulonglong4 __9;
        uint4 v__4 = *(uint4*)(Q + (((((((((((int)blockIdx.z) >> 5) * 33554432) + (i_7 * 524288)) + (i_12 * 131072)) + ((((int)threadIdx.x) >> 3) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.y) * 64)) + ((((int)threadIdx.x) & 7) * 8)) + 1048576));
        ((float2*)(&(__9.x)))->x = (float)(((half2*)(&(v__4.x)))->x);
        ((float2*)(&(__9.x)))->y = (float)(((half2*)(&(v__4.x)))->y);
        ((float2*)(&(__9.y)))->x = (float)(((half2*)(&(v__4.y)))->x);
        ((float2*)(&(__9.y)))->y = (float)(((half2*)(&(v__4.y)))->y);
        ((float2*)(&(__9.z)))->x = (float)(((half2*)(&(v__4.z)))->x);
        ((float2*)(&(__9.z)))->y = (float)(((half2*)(&(v__4.z)))->y);
        ((float2*)(&(__9.w)))->x = (float)(((half2*)(&(v__4.w)))->x);
        ((float2*)(&(__9.w)))->y = (float)(((half2*)(&(v__4.w)))->y);
        ulonglong4 v__5 = make_ulonglong4(*(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/), *(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/), *(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/), *(unsigned long long*)&make_float2(0x1p-4f/*6.250000e-02*/, 0x1p-4f/*6.250000e-02*/));
        ((float2*)(&(__8.x)))->x = (((float2*)(&(__9.x)))->x*((float2*)(&(v__5.x)))->x);
        ((float2*)(&(__8.x)))->y = (((float2*)(&(__9.x)))->y*((float2*)(&(v__5.x)))->y);
        ((float2*)(&(__8.y)))->x = (((float2*)(&(__9.y)))->x*((float2*)(&(v__5.y)))->x);
        ((float2*)(&(__8.y)))->y = (((float2*)(&(__9.y)))->y*((float2*)(&(v__5.y)))->y);
        ((float2*)(&(__8.z)))->x = (((float2*)(&(__9.z)))->x*((float2*)(&(v__5.z)))->x);
        ((float2*)(&(__8.z)))->y = (((float2*)(&(__9.z)))->y*((float2*)(&(v__5.z)))->y);
        ((float2*)(&(__8.w)))->x = (((float2*)(&(__9.w)))->x*((float2*)(&(v__5.w)))->x);
        ((float2*)(&(__8.w)))->y = (((float2*)(&(__9.w)))->y*((float2*)(&(v__5.w)))->y);
      ((half2*)(&(__7.x)))->x = (half_t)(((float2*)(&(__8.x)))->x);
      ((half2*)(&(__7.x)))->y = (half_t)(((float2*)(&(__8.x)))->y);
      ((half2*)(&(__7.y)))->x = (half_t)(((float2*)(&(__8.y)))->x);
      ((half2*)(&(__7.y)))->y = (half_t)(((float2*)(&(__8.y)))->y);
      ((half2*)(&(__7.z)))->x = (half_t)(((float2*)(&(__8.z)))->x);
      ((half2*)(&(__7.z)))->y = (half_t)(((float2*)(&(__8.z)))->y);
      ((half2*)(&(__7.w)))->x = (half_t)(((float2*)(&(__8.w)))->x);
      ((half2*)(&(__7.w)))->y = (half_t)(((float2*)(&(__8.w)))->y);
      *(uint4*)(((half_t*)buf_dyn_shmem) + ((((((((i_7 & 1) * 4096) + (i_12 * 1024)) + ((((int)threadIdx.x) >> 3) * 64)) + (((((((int)threadIdx.x) & 63) >> 5) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 32)) + (((((((int)threadIdx.x) & 31) >> 4) + ((((int)threadIdx.x) & 3) >> 1)) & 1) * 16)) + (((((((int)threadIdx.x) & 15) >> 3) + (((int)threadIdx.x) & 1)) & 1) * 8)) + 16384)) = __7;
    }
    tl::cp_async_commit();
    #pragma unroll
    for (int i_13 = 0; i_13 < 16; ++i_13) {
      uint1 __10;
      float2 v__6 = *(float2*)(o + (i_13 * 2));
      ((half2*)(&(__10.x)))->x = (half_t)(v__6.x);
      ((half2*)(&(__10.x)))->y = (half_t)(v__6.y);
      *(uint1*)(O + ((((((((((((((int)blockIdx.y) * 268435456) + ((((int)blockIdx.z) >> 5) * 33554432)) + (i_7 * 524288)) + (((i_13 & 3) >> 1) * 262144)) + (((((int)threadIdx.x) & 63) >> 5) * 131072)) + ((i_13 & 1) * 65536)) + (((((int)threadIdx.x) & 31) >> 2) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.x) * 64)) + ((i_13 >> 2) * 16)) + ((((int)threadIdx.x) >> 6) * 8)) + ((((int)threadIdx.x) & 3) * 2))) = __10;
    }
  }
  tl::cp_async_wait<2>();
  __syncthreads();
  tl::gemm_ss<64, 64, 64, 2, 2, 0, 1, 1, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[16384])), (&(((half_t*)buf_dyn_shmem)[0])), (&(s[0])));
  __syncthreads();
  #pragma unroll
  for (int i_14 = 0; i_14 < 16; ++i_14) {
    for (int vec_s_1 = 0; vec_s_1 < 2; ++vec_s_1) {
      float condval_1;
      if (((((((i_14 >> 2) * 16) + ((((int)threadIdx.x) >> 6) * 8)) + ((((int)threadIdx.x) & 3) * 2)) + vec_s_1) <= ((((((i_14 & 3) >> 1) * 32) + (((((int)threadIdx.x) & 63) >> 5) * 16)) + ((i_14 & 1) * 8)) + ((((int)threadIdx.x) & 31) >> 2)))) {
        condval_1 = s[((i_14 * 2) + vec_s_1)];
      } else {
        condval_1 = 0x0p+0f/*0.000000e+00*/;
      }
      ((half_t*)buf_dyn_shmem)[((((((((((((i_14 & 3) >> 1) * 2048) + (((((int)threadIdx.x) & 63) >> 5) * 1024)) + ((i_14 & 1) * 512)) + (((((int)threadIdx.x) & 31) >> 2) * 64)) + (((((((((i_14 >> 2) * 16) + ((((int)threadIdx.x) >> 6) * 8)) + ((((int)threadIdx.x) & 3) * 2)) + vec_s_1) >> 5) + ((((int)threadIdx.x) & 31) >> 4)) & 1) * 32)) + (((((((int)threadIdx.x) & 15) >> 3) + ((i_14 & 7) >> 2)) & 1) * 16)) + ((((((int)threadIdx.x) >> 6) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 8)) + ((((int)threadIdx.x) & 3) * 2)) + vec_s_1) + 24576)] = ((half_t)condval_1);
    }
  }
  tl::fence_proxy_async();
  tl::cp_async_wait<3>();
  __syncthreads();
  tl::gemm_ss<64, 64, 64, 2, 2, 0, 0, 1, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[24576])), (&(((half_t*)buf_dyn_shmem)[8192])), (&(o[0])));
  __syncthreads();
  #pragma unroll
  for (int i_15 = 0; i_15 < 4; ++i_15) {
    tl::ptx_stmatrix_x4((&(((half_t*)buf_dyn_shmem)[(((((((((((int)threadIdx.x) & 31) >> 4) * 2048) + (((((int)threadIdx.x) & 63) >> 5) * 1024)) + ((((int)threadIdx.x) & 15) * 64)) + (((((((int)threadIdx.x) & 7) >> 2) + (i_15 >> 1)) & 1) * 32)) + (((((((int)threadIdx.x) & 3) >> 1) + (i_15 & 1)) & 1) * 16)) + ((((((int)threadIdx.x) >> 6) + (((int)threadIdx.x) & 1)) & 1) * 8)) + 28672)])), __pack_half2(((half_t)h[(i_15 * 8)]), ((half_t)h[((i_15 * 8) + 1)])), __pack_half2(((half_t)h[((i_15 * 8) + 2)]), ((half_t)h[((i_15 * 8) + 3)])), __pack_half2(((half_t)h[((i_15 * 8) + 4)]), ((half_t)h[((i_15 * 8) + 5)])), __pack_half2(((half_t)h[((i_15 * 8) + 6)]), ((half_t)h[((i_15 * 8) + 7)])));
  }
  tl::fence_proxy_async();
  tl::cp_async_wait<3>();
  __syncthreads();
  tl::gemm_ss<64, 64, 64, 2, 2, 1, 0, 0, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[0])), (&(((half_t*)buf_dyn_shmem)[8192])), (&(h[0])));
  tl::cp_async_wait<2>();
  __syncthreads();
  tl::gemm_ss<64, 64, 64, 2, 2, 0, 0, 0, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[16384])), (&(((half_t*)buf_dyn_shmem)[28672])), (&(o[0])));
  #pragma unroll
  for (int i_16 = 0; i_16 < 16; ++i_16) {
    uint1 __11;
    float2 v__7 = *(float2*)(o + (i_16 * 2));
    ((half2*)(&(__11.x)))->x = (half_t)(v__7.x);
    ((half2*)(&(__11.x)))->y = (half_t)(v__7.y);
    *(uint1*)(O + ((((((((((((((int)blockIdx.y) * 268435456) + ((((int)blockIdx.z) >> 5) * 33554432)) + (((i_16 & 3) >> 1) * 262144)) + (((((int)threadIdx.x) & 63) >> 5) * 131072)) + ((i_16 & 1) * 65536)) + (((((int)threadIdx.x) & 31) >> 2) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.x) * 64)) + ((i_16 >> 2) * 16)) + ((((int)threadIdx.x) >> 6) * 8)) + ((((int)threadIdx.x) & 3) * 2)) + 32505856)) = __11;
  }
  tl::cp_async_wait<0>();
  __syncthreads();
  tl::gemm_ss<64, 64, 64, 2, 2, 0, 1, 1, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[20480])), (&(((half_t*)buf_dyn_shmem)[4096])), (&(s[0])));
  __syncthreads();
  #pragma unroll
  for (int i_17 = 0; i_17 < 16; ++i_17) {
    for (int vec_s_2 = 0; vec_s_2 < 2; ++vec_s_2) {
      float condval_2;
      if (((((((i_17 >> 2) * 16) + ((((int)threadIdx.x) >> 6) * 8)) + ((((int)threadIdx.x) & 3) * 2)) + vec_s_2) <= ((((((i_17 & 3) >> 1) * 32) + (((((int)threadIdx.x) & 63) >> 5) * 16)) + ((i_17 & 1) * 8)) + ((((int)threadIdx.x) & 31) >> 2)))) {
        condval_2 = s[((i_17 * 2) + vec_s_2)];
      } else {
        condval_2 = 0x0p+0f/*0.000000e+00*/;
      }
      ((half_t*)buf_dyn_shmem)[((((((((((((i_17 & 3) >> 1) * 2048) + (((((int)threadIdx.x) & 63) >> 5) * 1024)) + ((i_17 & 1) * 512)) + (((((int)threadIdx.x) & 31) >> 2) * 64)) + (((((((((i_17 >> 2) * 16) + ((((int)threadIdx.x) >> 6) * 8)) + ((((int)threadIdx.x) & 3) * 2)) + vec_s_2) >> 5) + ((((int)threadIdx.x) & 31) >> 4)) & 1) * 32)) + (((((((int)threadIdx.x) & 15) >> 3) + ((i_17 & 7) >> 2)) & 1) * 16)) + ((((((int)threadIdx.x) >> 6) + ((((int)threadIdx.x) & 7) >> 2)) & 1) * 8)) + ((((int)threadIdx.x) & 3) * 2)) + vec_s_2) + 24576)] = ((half_t)condval_2);
    }
  }
  tl::fence_proxy_async();
  tl::cp_async_wait<1>();
  __syncthreads();
  tl::gemm_ss<64, 64, 64, 2, 2, 0, 0, 1, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[24576])), (&(((half_t*)buf_dyn_shmem)[12288])), (&(o[0])));
  __syncthreads();
  #pragma unroll
  for (int i_18 = 0; i_18 < 4; ++i_18) {
    tl::ptx_stmatrix_x4((&(((half_t*)buf_dyn_shmem)[(((((((((((int)threadIdx.x) & 31) >> 4) * 2048) + (((((int)threadIdx.x) & 63) >> 5) * 1024)) + ((((int)threadIdx.x) & 15) * 64)) + (((((((int)threadIdx.x) & 7) >> 2) + (i_18 >> 1)) & 1) * 32)) + (((((((int)threadIdx.x) & 3) >> 1) + (i_18 & 1)) & 1) * 16)) + ((((((int)threadIdx.x) >> 6) + (((int)threadIdx.x) & 1)) & 1) * 8)) + 28672)])), __pack_half2(((half_t)h[(i_18 * 8)]), ((half_t)h[((i_18 * 8) + 1)])), __pack_half2(((half_t)h[((i_18 * 8) + 2)]), ((half_t)h[((i_18 * 8) + 3)])), __pack_half2(((half_t)h[((i_18 * 8) + 4)]), ((half_t)h[((i_18 * 8) + 5)])), __pack_half2(((half_t)h[((i_18 * 8) + 6)]), ((half_t)h[((i_18 * 8) + 7)])));
  }
  tl::fence_proxy_async();
  tl::cp_async_wait<1>();
  __syncthreads();
  tl::gemm_ss<64, 64, 64, 2, 2, 1, 0, 0, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[4096])), (&(((half_t*)buf_dyn_shmem)[12288])), (&(h[0])));
  tl::cp_async_wait<0>();
  __syncthreads();
  tl::gemm_ss<64, 64, 64, 2, 2, 0, 0, 0, 64, 64, 0, 0>((&(((half_t*)buf_dyn_shmem)[20480])), (&(((half_t*)buf_dyn_shmem)[28672])), (&(o[0])));
  #pragma unroll
  for (int i_19 = 0; i_19 < 16; ++i_19) {
    uint1 __12;
    float2 v__8 = *(float2*)(o + (i_19 * 2));
    ((half2*)(&(__12.x)))->x = (half_t)(v__8.x);
    ((half2*)(&(__12.x)))->y = (half_t)(v__8.y);
    *(uint1*)(O + ((((((((((((((int)blockIdx.y) * 268435456) + ((((int)blockIdx.z) >> 5) * 33554432)) + (((i_19 & 3) >> 1) * 262144)) + (((((int)threadIdx.x) & 63) >> 5) * 131072)) + ((i_19 & 1) * 65536)) + (((((int)threadIdx.x) & 31) >> 2) * 8192)) + ((((int)blockIdx.z) & 31) * 256)) + (((int)blockIdx.x) * 64)) + ((i_19 >> 2) * 16)) + ((((int)threadIdx.x) >> 6) * 8)) + ((((int)threadIdx.x) & 3) * 2)) + 33030144)) = __12;
  }
  #pragma unroll
  for (int i_20 = 0; i_20 < 16; ++i_20) {
    *(float2*)(final_state + ((((((((((((int)blockIdx.z) * 65536) + (((int)blockIdx.y) * 16384)) + (((i_20 & 3) >> 1) * 8192)) + (((((int)threadIdx.x) & 63) >> 5) * 4096)) + ((i_20 & 1) * 2048)) + (((((int)threadIdx.x) & 31) >> 2) * 256)) + (((int)blockIdx.x) * 64)) + ((i_20 >> 2) * 16)) + ((((int)threadIdx.x) >> 6) * 8)) + ((((int)threadIdx.x) & 3) * 2))) = *(float2*)(h + (i_20 * 2));
  }
}


#define ERROR_BUF_SIZE 1024
static char error_buf[ERROR_BUF_SIZE];

extern "C" const char* get_last_error() {
    return error_buf;
}

extern "C" int init() {
    error_buf[0] = '\0';
    
    cudaError_t result_chunk_linear_attn_fwd_kernel = cudaFuncSetAttribute(chunk_linear_attn_fwd_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 65536);
    if (result_chunk_linear_attn_fwd_kernel != CUDA_SUCCESS) {
        snprintf(error_buf, ERROR_BUF_SIZE, "Failed to set the allowed dynamic shared memory size to %d with error: %s", 65536, cudaGetErrorString(result_chunk_linear_attn_fwd_kernel));
        return -1;
    }

    return 0;
}

extern "C" int call(half_t* __restrict__ Q, half_t* __restrict__ K, half_t* __restrict__ V, half_t* __restrict__ O, float* __restrict__ final_state, cudaStream_t stream=cudaStreamDefault) {
        chunk_linear_attn_fwd_kernel<<<dim3(4, 4, 256), dim3(128, 1, 1), 65536, stream>>>(K, O, Q, V, final_state);
        TILELANG_CHECK_LAST_ERROR("chunk_linear_attn_fwd_kernel");

        return 0;
}