Kernels:
Trusted publisher
Uploaded using `kernel-builder`.
Browse files- build/torch-xpu/_ops.py +1 -1
- build/torch-xpu/metadata.json +11 -10
- build/torch-xpu/metadata.json.sigstore +1 -1
- build/torch-xpu/ops/common/backends/intracard.py +1 -1
- build/torch-xpu/ops/common/backends/tilelang/__init__.py +3 -3
- build/torch-xpu/ops/common/backends/tilelang/parallel_attn_bwd.py +1 -1
- build/torch-xpu/ops/deltaformer/parallel.py +1 -1
- build/torch-xpu/ops/kda/backends/flashkda.py +1 -1
- build/torch-xpu/ops/kda/backends/tilelang/__init__.py +1 -1
- build/torch-xpu/utils/layers.py +332 -0
build/torch-xpu/_ops.py
CHANGED
|
@@ -28,7 +28,7 @@ def get_backend() -> str:
|
|
| 28 |
|
| 29 |
def _find_ops_name() -> str:
|
| 30 |
kernel_name = "fla"
|
| 31 |
-
unique_id = "
|
| 32 |
backend = get_backend()
|
| 33 |
return f"_{kernel_name}_{backend}_{unique_id}"
|
| 34 |
|
|
|
|
| 28 |
|
| 29 |
def _find_ops_name() -> str:
|
| 30 |
kernel_name = "fla"
|
| 31 |
+
unique_id = "486b34c"
|
| 32 |
backend = get_backend()
|
| 33 |
return f"_{kernel_name}_{backend}_{unique_id}"
|
| 34 |
|
build/torch-xpu/metadata.json
CHANGED
|
@@ -1,6 +1,6 @@
|
|
| 1 |
{
|
| 2 |
"name": "fla",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"kernels-minver": "0.14.0",
|
| 6 |
"license": "MIT",
|
|
@@ -15,7 +15,7 @@
|
|
| 15 |
"algorithm": "sha256",
|
| 16 |
"files": {
|
| 17 |
"__init__.py": "bVRNYUhThvULeUaovvQFZCJ9IZZsC+0kU/HbnXdYT1A=",
|
| 18 |
-
"_ops.py": "
|
| 19 |
"layers.py": "/39XVQW4d4GHeNROqT67ujKbItdteLBtm0PT6zq41YQ=",
|
| 20 |
"modules/__init__.py": "6ZjjRMt6Rf+2OegXDC8w58RD1IW3SBK+dxBoZa4Uohc=",
|
| 21 |
"modules/activations.py": "Lp7xLUf7r785nt93NFhqoSQM+G7AK2e6kviE2/0TRJs=",
|
|
@@ -82,10 +82,10 @@
|
|
| 82 |
"ops/comba/wy_fast.py": "EsJfyG3HDGY9rgJ97cn2SAq1DeqnXdgj5mzdOlWtz2I=",
|
| 83 |
"ops/common/__init__.py": "47DEQpj8HBSa+/TImW+5JCeuQeRkm5NMpJWZG3hSuFU=",
|
| 84 |
"ops/common/backends/__init__.py": "aEhY8S2W0I6qbjedaCYjqLt4VaNumWM+bjKtPEJSAN0=",
|
| 85 |
-
"ops/common/backends/intracard.py": "
|
| 86 |
-
"ops/common/backends/tilelang/__init__.py": "
|
| 87 |
"ops/common/backends/tilelang/chunk_bwd.py": "5zXyBzqIXRgoZdZ8zjLNEIz3jBksHpaVxnqUk66ckjQ=",
|
| 88 |
-
"ops/common/backends/tilelang/parallel_attn_bwd.py": "
|
| 89 |
"ops/common/backends/tilelang/parallel_attn_fwd.py": "7INMyGs1d3H8BKcjRnMvu/NHMm9VeNFTETR5z7ZO15c=",
|
| 90 |
"ops/common/chunk_delta_h.py": "QmkSsbbyOHai9cVaMB8O/8ep4NxiM6ieUVKmja6USw8=",
|
| 91 |
"ops/common/chunk_h.py": "imyz71p9AD6nOtStpN+zzgE0yVbGnq4DO3FitrO3hvc=",
|
|
@@ -111,7 +111,7 @@
|
|
| 111 |
"ops/deltaformer/__init__.py": "51ZmNFoHy3P0lVn/2EHKu1QS2BMC6h+hnAwSIr9szYg=",
|
| 112 |
"ops/deltaformer/invcum.py": "ren181AxlHBv2MoZIHmSGMmJ0sW6HDb3t8fqG7gDELU=",
|
| 113 |
"ops/deltaformer/naive.py": "8b88rviKkDZpW+F9GwgEkVkwuSioECwgRJX//CVksnI=",
|
| 114 |
-
"ops/deltaformer/parallel.py": "
|
| 115 |
"ops/dsa/__init__.py": "ZNyws889LPBiY8gwB5zTs2vQwJYV8bX5jg6Cldkkq9s=",
|
| 116 |
"ops/dsa/naive.py": "lKgrdKOoE/qCdF4DtrI0uEkRdB28kUYbtn8hO63fzyo=",
|
| 117 |
"ops/forgetting_attn/__init__.py": "A6ozg9zdEsz4lgUAXqEIUY8CeNp+5IZqYKOpChpj/A4=",
|
|
@@ -179,8 +179,8 @@
|
|
| 179 |
"ops/hgrn/naive.py": "XKo+5zJ+G4elkgK5ARsoXm5WYkG1QidTKNIuPT4k588=",
|
| 180 |
"ops/kda/__init__.py": "JFZPAQH4eiYFbpBh7Xccrw2PbZoAtNve1wHqdSW0Wss=",
|
| 181 |
"ops/kda/backends/__init__.py": "pk1xLku49AblKggFmQjmGlr1oNqUgdzxggzzgx8yeZk=",
|
| 182 |
-
"ops/kda/backends/flashkda.py": "
|
| 183 |
-
"ops/kda/backends/tilelang/__init__.py": "
|
| 184 |
"ops/kda/backends/tilelang/chunk_bwd_dqkg.py": "Fk1eT3M/KPFvfVRrJQ6C0158h5GaR72YK3tgSxkKoUo=",
|
| 185 |
"ops/kda/chunk.py": "LnoIH4MI8gHR8613fOKLqJGVNQkbtzjEJ0R7qX8/kQE=",
|
| 186 |
"ops/kda/chunk_bwd.py": "AahkP3AepxWsoyGCZHlwUGc7rGIHJh68Im5r/Zs3MK0=",
|
|
@@ -302,7 +302,8 @@
|
|
| 302 |
"utils/_decorators.py": "8OxDii2b+D+deTCQLPK1sdFx7McS/Wi9+AO67mlLrYg=",
|
| 303 |
"utils/_device.py": "yXcM5ZkBLCOvHl/Dh25SveTgbSemR6PkFAl/UabLYWE=",
|
| 304 |
"utils/_testing.py": "eFchc/74QJDXWv5T1YmIZFzRRpFWh+3gaYpQ1doztzs=",
|
| 305 |
-
"utils/ascend_ub_manager.py": "TcjVL6L7YiNwsa7YZ4+9KrIPQ2RmG7PJCfmXYc24uEc="
|
|
|
|
| 306 |
}
|
| 307 |
},
|
| 308 |
"provenance": {
|
|
@@ -312,7 +313,7 @@
|
|
| 312 |
"dirty": false
|
| 313 |
},
|
| 314 |
"kernel": {
|
| 315 |
-
"commit": "
|
| 316 |
"dirty": false
|
| 317 |
}
|
| 318 |
}
|
|
|
|
| 1 |
{
|
| 2 |
"name": "fla",
|
| 3 |
+
"id": "_fla_xpu_486b34c",
|
| 4 |
"version": 1,
|
| 5 |
"kernels-minver": "0.14.0",
|
| 6 |
"license": "MIT",
|
|
|
|
| 15 |
"algorithm": "sha256",
|
| 16 |
"files": {
|
| 17 |
"__init__.py": "bVRNYUhThvULeUaovvQFZCJ9IZZsC+0kU/HbnXdYT1A=",
|
| 18 |
+
"_ops.py": "RqdWVrqfjdvXo+8ZZWZw/MgkOJ0IdLQ6Zy8LRljdYC8=",
|
| 19 |
"layers.py": "/39XVQW4d4GHeNROqT67ujKbItdteLBtm0PT6zq41YQ=",
|
| 20 |
"modules/__init__.py": "6ZjjRMt6Rf+2OegXDC8w58RD1IW3SBK+dxBoZa4Uohc=",
|
| 21 |
"modules/activations.py": "Lp7xLUf7r785nt93NFhqoSQM+G7AK2e6kviE2/0TRJs=",
|
|
|
|
| 82 |
"ops/comba/wy_fast.py": "EsJfyG3HDGY9rgJ97cn2SAq1DeqnXdgj5mzdOlWtz2I=",
|
| 83 |
"ops/common/__init__.py": "47DEQpj8HBSa+/TImW+5JCeuQeRkm5NMpJWZG3hSuFU=",
|
| 84 |
"ops/common/backends/__init__.py": "aEhY8S2W0I6qbjedaCYjqLt4VaNumWM+bjKtPEJSAN0=",
|
| 85 |
+
"ops/common/backends/intracard.py": "yXACesLmdhG1PwAFJuquz4F+hqSNtFvXwZ51YeEN8FM=",
|
| 86 |
+
"ops/common/backends/tilelang/__init__.py": "2Zz/LkJfy8oOvqUcen6iQDMcf2prznZpYhHAc5QYW3E=",
|
| 87 |
"ops/common/backends/tilelang/chunk_bwd.py": "5zXyBzqIXRgoZdZ8zjLNEIz3jBksHpaVxnqUk66ckjQ=",
|
| 88 |
+
"ops/common/backends/tilelang/parallel_attn_bwd.py": "w+KjVu0voRjcFhmwWfIEYaSwmlk17oIu9EKRvjNIAro=",
|
| 89 |
"ops/common/backends/tilelang/parallel_attn_fwd.py": "7INMyGs1d3H8BKcjRnMvu/NHMm9VeNFTETR5z7ZO15c=",
|
| 90 |
"ops/common/chunk_delta_h.py": "QmkSsbbyOHai9cVaMB8O/8ep4NxiM6ieUVKmja6USw8=",
|
| 91 |
"ops/common/chunk_h.py": "imyz71p9AD6nOtStpN+zzgE0yVbGnq4DO3FitrO3hvc=",
|
|
|
|
| 111 |
"ops/deltaformer/__init__.py": "51ZmNFoHy3P0lVn/2EHKu1QS2BMC6h+hnAwSIr9szYg=",
|
| 112 |
"ops/deltaformer/invcum.py": "ren181AxlHBv2MoZIHmSGMmJ0sW6HDb3t8fqG7gDELU=",
|
| 113 |
"ops/deltaformer/naive.py": "8b88rviKkDZpW+F9GwgEkVkwuSioECwgRJX//CVksnI=",
|
| 114 |
+
"ops/deltaformer/parallel.py": "sDHrnr2fM4+0T2cyAJoI2VP9Lan3NmaaDxk9JLtyLWk=",
|
| 115 |
"ops/dsa/__init__.py": "ZNyws889LPBiY8gwB5zTs2vQwJYV8bX5jg6Cldkkq9s=",
|
| 116 |
"ops/dsa/naive.py": "lKgrdKOoE/qCdF4DtrI0uEkRdB28kUYbtn8hO63fzyo=",
|
| 117 |
"ops/forgetting_attn/__init__.py": "A6ozg9zdEsz4lgUAXqEIUY8CeNp+5IZqYKOpChpj/A4=",
|
|
|
|
| 179 |
"ops/hgrn/naive.py": "XKo+5zJ+G4elkgK5ARsoXm5WYkG1QidTKNIuPT4k588=",
|
| 180 |
"ops/kda/__init__.py": "JFZPAQH4eiYFbpBh7Xccrw2PbZoAtNve1wHqdSW0Wss=",
|
| 181 |
"ops/kda/backends/__init__.py": "pk1xLku49AblKggFmQjmGlr1oNqUgdzxggzzgx8yeZk=",
|
| 182 |
+
"ops/kda/backends/flashkda.py": "g63Eq+K84JInsin5NVb428RIRPhL3wH+6Y4a+OHSFqk=",
|
| 183 |
+
"ops/kda/backends/tilelang/__init__.py": "Y+1//U4pzuPg2mbHoCymDfs186j482cI85ZOPxi88E0=",
|
| 184 |
"ops/kda/backends/tilelang/chunk_bwd_dqkg.py": "Fk1eT3M/KPFvfVRrJQ6C0158h5GaR72YK3tgSxkKoUo=",
|
| 185 |
"ops/kda/chunk.py": "LnoIH4MI8gHR8613fOKLqJGVNQkbtzjEJ0R7qX8/kQE=",
|
| 186 |
"ops/kda/chunk_bwd.py": "AahkP3AepxWsoyGCZHlwUGc7rGIHJh68Im5r/Zs3MK0=",
|
|
|
|
| 302 |
"utils/_decorators.py": "8OxDii2b+D+deTCQLPK1sdFx7McS/Wi9+AO67mlLrYg=",
|
| 303 |
"utils/_device.py": "yXcM5ZkBLCOvHl/Dh25SveTgbSemR6PkFAl/UabLYWE=",
|
| 304 |
"utils/_testing.py": "eFchc/74QJDXWv5T1YmIZFzRRpFWh+3gaYpQ1doztzs=",
|
| 305 |
+
"utils/ascend_ub_manager.py": "TcjVL6L7YiNwsa7YZ4+9KrIPQ2RmG7PJCfmXYc24uEc=",
|
| 306 |
+
"utils/layers.py": "afB10gGysmmVfVRlPnduPpuvcdAeDAstmqPC0sYJbRw="
|
| 307 |
}
|
| 308 |
},
|
| 309 |
"provenance": {
|
|
|
|
| 313 |
"dirty": false
|
| 314 |
},
|
| 315 |
"kernel": {
|
| 316 |
+
"commit": "486b34c8bc724f1905fd3fdc04b599106f1b7cfd",
|
| 317 |
"dirty": false
|
| 318 |
}
|
| 319 |
}
|
build/torch-xpu/metadata.json.sigstore
CHANGED
|
@@ -1 +1 @@
|
|
| 1 |
-
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSzCCBtCgAwIBAgIUcLUwS8ZGOZpTFckuBYad/EcRfowwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwOTEwMTE0OTMxWhcNMjYwOTEwMTE1OTMxWjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEatgOlVutRSeN/IvVs4uAv++8mKGRC4AvEpPC27SsWFOaRSB/RqATHWZ7ZthokfIW+A2o7BC0EnY7xcmhXsAg+KOCBe8wggXrMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQU8SE+bjIqYvNZOVIkHZw7RWAGsvAwHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoNTE2MjI1MWY3MjhhNmRhMWUyNDg2NTRiYjgyODdhZTczYTg2NmU5YjATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoNTE2MjI1MWY3MjhhNmRhMWUyNDg2NTRiYjgyODdhZTczYTg2NmU5YjAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoNTE2MjI1MWY3MjhhNmRhMWUyNDg2NTRiYjgyODdhZTczYTg2NmU5YjAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKDUxNjIyNTFmNzI4YTZkYTFlMjQ4NjU0YmI4Mjg3YWU3M2E4NjZlOWIwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzQ0NzI5Njc1NjAvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBiQYKKwYBBAHWeQIEAgR7BHkAdwB1AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoIsmzWwAAAQDAEYwRAIgEmyE1Irj+3mbcSfTbPl/+qoSfK75GnTJcV/cERL8vpsCIHMLYPZH4bj3hFMwDfn+hYctcj2lYMm7ALw0D6eOyJm8MAoGCCqGSM49BAMDA2kAMGYCMQDZSGDVfjjoaSRprSCwCGdjhmBLb27GO0wY/A6qZ4Xm6oSZBz7X7mPqQYjzy7Zl+PsCMQCDsgRonv+WuX4vR9WuIAIt3tlKSJglO4drA/L036cdxNGKfCizz4cVL/jTuU7okYo="}, "tlogEntries":[{"logIndex":"2782559873", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1789040971", "inclusionPromise":{"signedEntryTimestamp":"MEUCICAggujvWxO0+gtCJG6mum03QRHXwUEn4sDsJZEbMC8MAiEA/6YkxTjfiRZOHDzXHNVAV2/3XvlNpjz6sKm3Fyi+6qU="}, "inclusionProof":{"logIndex":"2660655611", "rootHash":"PYChoaAu0DlBu3LnoGGl2kwkLognz1MM+FweD0iRDq0=", "treeSize":"2660655623", "hashes":["36NDVNJJRb/TxTltgEWZFT51la19no2AmyHlOGKLGNk=", "PkqxwnlTzZE+S1ceB7cyGKDAOHaerkoS2CShqaPjjNs=", "jZzvno980PbNjQ6xB3TPp+jlnhhTsViVpZmrPBJPga4=", "fTIWdfNxGi6gaECbGk6CAlGdQsTtj6IsgG23aTfqabM=", "8jVR/L0rN/PtZq/Gg9byg06dgRh9flcMnXQhxiEsLiI=", "GJMx1Rev3xjPVzbjowVfCTjHOLM0j0Kz3dv/Bu532YA=", "ccyi8YQyZPNkAajV5ZFx5ZjNM/8yCv/iz0ENI8BGxFk=", "bV6QggOT5F+OezAYVS9tSqw3SVvF5y0b+SlDt+A0GWU=", "w0Ag7xIBEGCrsN0LkTnYa0IhTceEzYH45XGewd0Q1ZA=", "Ar3PQ2cYb0oAvZ4R9ScV6n1Qe9lr7qTDncX6LbR/Q3Q=", "goGYcVKIESdUGbpgYDNg0Cqp9hv6+UtnPmDDW6IBBG0=", "zmD+IBktdnbjs1Rw8qpVHtVgEgovTUgye5qyufoyF2M=", "TKG+Q1VKaHO0wYQXG0ay0oIleqTjGQY9pObaWFWMKoo=", "tUKyb7C480RhDVXaNDboFd+KJmVTMK/FWW47rRdB15c=", "wa3SxR/hCfRbs8pzwM4D8wuxTOnbBYK/CGQe1GqX7HA=", "RWKTC8lK6QH9wY0VD43p3P3N0jCodw0jaVv1bOki3J4=", "4SexKjDhe33uQFmbzX5SUsB0SyhTm14ObSvNGPpKDFk=", "7tD2K8mAqGTMENrP2YXWz71BeTWQPCiYd3CESAgQLz8=", "eXSLrLqKSxRTixSP/k4YimriB2lufao6trqTrsxxlyg=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2660655623\nPYChoaAu0DlBu3LnoGGl2kwkLognz1MM+FweD0iRDq0=\n\n— rekor.sigstore.dev wNI9ajBFAiEA86V21lwoQm7xDgf5MXWy84xKtMmIuc5uRs5rA87BfusCIET5Q99AqNpait7bntYY7JRRe27shTRHHvn4Gwm4H4Ee\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiI4ZjBlNWUzNWMwYTRmZWVhYWEwNzQwZDdlYjFiODBmNGI1ODE1NTQxZGI1MmExZTE2ZTAyN2Y1MTMzYzUyZThkIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJUUMzYmRJaE90eHdJVFhWd2MyM3crWlcra212SjBFdmgwUy9yQk5JTFlEek1nSWdYb0dnK2lPbXlMT1Z3ZFAxU0dTL0lhb2crelNjd1lQUHQ3L1JEK3RXYklzPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRla05EUW5SRFowRjNTVUpCWjBsVlkweFZkMU00V2tkUFduQlVSbU5yZFVKWllXUXZSV05TWm05M2QwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlVUlhkTlZFVXdUMVJOZUZkb1kwNU5hbGwzVDFSRmQwMVVSVEZQVkUxNFYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZoZEdkUGJGWjFkRkpUWlU0dlNYWldjelIxUVhZckt6aHRTMGRTUXpSQmRrVndVRU1LTWpkVGMxZEdUMkZTVTBJdlVuRkJWRWhYV2pkYWRHaHZhMlpKVnl0Qk1tODNRa013Ulc1Wk4zaGpiV2hZYzBGbkswdFBRMEpsT0hkbloxaHlUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlU0VTBVckNtSnFTWEZaZGs1YVQxWkphMGhhZHpkU1YwRkhjM1pCZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOU9WRVV5VFdwSk1VMVhXVE5OYW1ob1RtMVNhRTFYVlhsT1JHY3lDazVVVW1sWmFtZDVUMFJrYUZwVVkzcFpWR2N5VG0xVk5WbHFRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMDVVUlRKTmFra3hUVmRaTTAxcWFHaE9iVkpvVFZkVmVVNUVaekpPVkZKcFdXcG5lVTlFWkdnS1dsUmplbGxVWnpKT2JWVTFXV3BCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMDVVUlRKTmFra3hUVmRaTTAxcWFHaE9iVkpvVFZkVmVVNUVaeklLVGxSU2FWbHFaM2xQUkdSb1dsUmplbGxVWnpKT2JWVTFXV3BCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwUlZlRTVxU1hrS1RsUkdiVTU2U1RSWlZGcHJXVlJHYkUxcVVUUk9hbFV3V1cxSk5FMXFaek5aVjFVelRUSkZORTVxV214UFYwbDNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWxFd1RucEpOVTVxWXpGT2FrRjJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwVVZsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTTjBKSWEwRUtaSGRDTVVGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUpjMjE2VjNkQlFVRlJSQXBCUlZsM1VrRkpaMFZ0ZVVVeFNYSnFLek50WW1OVFpsUmlVR3d2SzNGdlUyWkxOelZIYmxSS1kxWXZZMFZTVERoMmNITkRTVWhOVEZsUVdrZzBZbW96Q21oR1RYZEVabTRyYUZsamRHTnFNbXhaVFcwM1FVeDNNRVEyWlU5NVNtMDRUVUZ2UjBORGNVZFRUVFE1UWtGTlJFRXlhMEZOUjFsRFRWRkVXbE5IUkZZS1ptcHFiMkZUVW5CeVUwTjNRMGRrYW1odFFreGlNamRIVHpCM1dTOUJObkZhTkZodE5tOVRXa0o2TjFnM2JWQnhVVmxxZW5rM1dtd3JVSE5EVFZGRFJBcHpaMUp2Ym5ZclYzVllOSFpTT1ZkMVNVRkpkRE4wYkV0VFNtZHNUelJrY2tFdlREQXpObU5rZUU1SFMyWkRhWHA2TkdOV1RDOXFWSFZWTjI5cldXODlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyTADAgEAMIICwAYJKoZIhvcNAQcCoIICsTCCAq0CAQMxDTALBglghkgBZQMEAgEwgbgGCyqGSIb3DQEJEAEEoIGoBIGlMIGiAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQg8HrbePstO33we4pBO9MQ1YxSBb4uPfIwwUDxO5YvCWYCFQDXgQs9s3Imma8MQ0RvPPGRndKtxxgPMjAyNjA5MTAxMTQ5MzFaMAMCAQGgMqQwMC4xFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEVMBMGA1UEAxMMc2lnc3RvcmUtdHNhoAAxggHaMIIB1gIBATBRMDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAsGCWCGSAFlAwQCAaCB/DAaBgkqhkiG9w0BCQMxDQYLKoZIhvcNAQkQAQQwHAYJKoZIhvcNAQkFMQ8XDTI2MDkxMDExNDkzMVowLwYJKoZIhvcNAQkEMSIEIBfB3puZkKtIJ5hht9N2huUtcwDnZUnPrFo1MrOso2FKMIGOBgsqhkiG9w0BCRACLzF/MH0wezB5BCCF+Se8B6tiysO0Q1bBDvyBssaIP9p6uebYcNnROs0FtzBVMD2kOzA5MRUwEwYDVQQKEwxzaWdzdG9yZS5kZXYxIDAeBgNVBAMTF3NpZ3N0b3JlLXRzYS1zZWxmc2lnbmVkAhQ6E1QvDJBh7rzBQy/Lio6LKiOLDDAKBggqhkjOPQQDAgRmMGQCMAySFaX0alMcOpofNMtDltijC+JaUFyXcmZsAENwoXNWaPB3gIHrwjEizcm84VtrIgIwbf7wJumvSBESom90MvknXKCJyntPzEK+0ZTMLRuQe7k+ykpnhJqsHL5dkszrjTKl"}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"jw5eNcCk/uqqB0DX6xuA9LWBVUHbUqHhbgJ/UTPFLo0="}, "signature":"MEUCIQC3bdIhOtxwITXVwc23w+ZW+kmvJ0Evh0S/rBNILYDzMgIgXoGg+iOmyLOVwdP1SGS/Iaog+zScwYPPt7/RD+tWbIs="}}
|
|
|
|
| 1 |
+
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHTDCCBtKgAwIBAgIUP6oOKnWdWXzuNikaR3pS10hAhX8wCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwOTExMDkzNDI3WhcNMjYwOTExMDk0NDI3WjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAE/I/niYH7aq8Aiu59OOSSHh3ffZvQyfzJMIvha8VQzuHubNblOx3KgR5RQnzJQRwoX8CamYKnfMnxVDwkKsnt4aOCBfEwggXtMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUDX14lWno9y4R7tqwCDEOgLPgXk8wHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoNDg2YjM0YzhiYzcyNGYxOTA1ZmQzZmRjMDRiNTk5MTA2ZjFiN2NmZDATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoNDg2YjM0YzhiYzcyNGYxOTA1ZmQzZmRjMDRiNTk5MTA2ZjFiN2NmZDAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoNDg2YjM0YzhiYzcyNGYxOTA1ZmQzZmRjMDRiNTk5MTA2ZjFiN2NmZDAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKDQ4NmIzNGM4YmM3MjRmMTkwNWZkM2ZkYzA0YjU5OTEwNmYxYjdjZmQwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzQ1ODQ1MzcxNTEvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBiwYKKwYBBAHWeQIEAgR9BHsAeQB3AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoI/RgtoAAAQDAEgwRgIhAJCtkLOT44GW874XrNU/Zjf9VGoFgxGH6RpBoIQldd2ZAiEA5ZGLc8CwVhteUFSPbBwtZqLVFT4z6LU/8Yyor4L4ZkMwCgYIKoZIzj0EAwMDaAAwZQIwPJa49KcsvrsxcIFr7h1654usOnFK1oIkdxEKKUrkrTNR1bjitOPr0aT+gbnWyrD3AjEAo/TOW4pMsrhaPgOeyRuBFSyt7+4xbgWuUk8flv4rJibX6MHk/ezY3kJwlHRZw5fM"}, "tlogEntries":[{"logIndex":"2792115965", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1789119267", "inclusionPromise":{"signedEntryTimestamp":"MEUCICRwKxdXgQvuwJKQ0pwoax3eal52G4brQqr6z/SBxXjxAiEA8hE4w0dNt2fauW9D4NTL1RWGfyXMC/uGLz51NlIT/Es="}, "inclusionProof":{"logIndex":"2670211703", "rootHash":"y0w1WsHl2NY/58/9VElZaSb66x0Rf7K38UbSHRu++iY=", "treeSize":"2670211729", "hashes":["9XOS9/3dyWltcLRvJBBE436Q0t7eSgTKl3zqmgSt5A8=", "ibPeYw+FXtq5FdU8i0GhFj2VJ3zJuo5EmvaPE6MElcg=", "+tuYqJAHaURKCOqXTiQ2lRB+bU5sXGSVMnXiMyXeavI=", "AULHa1uxZHtdR+f2mP0ry+VXHHPaEOPNRR9r+qTHRbY=", "VNWfa+WI7PHGXZfUgDVKKuqQRKIfszEyvovMnsEqwbY=", "EhmEo3Q20R/DLPBhrveFlgtupxXTetFjchPEucwy/0c=", "a6kEQGoVZCWNSX3xudh46fveV+bXCZypb1VkDs7wDao=", "ku5E92I8CXUxOqmyKKc9czv7AseU5MJsY7PW8viBTL0=", "RlHHXFrRKT9kTpQoNaVKXdXxZG5nRNzACv43hrsDJ24=", "c+/cPCxMARuFdCZhcFZD7TxfnzGnBopBeOg3TDr4c68=", "1aFX/gxNiOXrxEXMHUxDEaPw7osJHuZlNTsx2S1nWcE=", "CQNgRg90pO6+IqGHeFwZV3GOx/yBXjDoB8Xld1+4prA=", "/mjQogAQefxrs6hDWFie0uOlNqZc4omr6I+XxshnZ1Y=", "Ari+W6zWoRpeR1Lw5lnM64DXE7cw7awR/QTfo7IO3PM=", "4SexKjDhe33uQFmbzX5SUsB0SyhTm14ObSvNGPpKDFk=", "7tD2K8mAqGTMENrP2YXWz71BeTWQPCiYd3CESAgQLz8=", "eXSLrLqKSxRTixSP/k4YimriB2lufao6trqTrsxxlyg=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2670211729\ny0w1WsHl2NY/58/9VElZaSb66x0Rf7K38UbSHRu++iY=\n\n— rekor.sigstore.dev wNI9ajBFAiEAxh8gdxiAF0Khd6mF+w9VD+QKpvee0+V2X1OyJzidejoCICuyOJHgn1HGVThB+SUtwV33hYYI/CTn1KVnbxgZyP3g\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiI3OTc5YjkyNjNkYzhjOTVhOTZjZTI0ZmMzNzRkZDNiODYzZmUwYjRhOWVjODUyZjJiMzBhNmY5MmY3MzZmZDc1In19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJUUREY0E3KzhmSWR5bzZEeCsveXExYVpGU092cWZuMStqaitwUnpnd05PRDNBSWdiTERLZC9GVFRCRE5KZ2VFRWVBbERPVEZCL2NYSURUT2xQVFlMUVhwV0pzPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFVSRU5EUW5STFowRjNTVUpCWjBsVlVEWnZUMHR1VjJSWFdIcDFUbWxyWVZJemNGTXhNR2hCYUZnNGQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlVUlhoTlJHdDZUa1JKTTFkb1kwNU5hbGwzVDFSRmVFMUVhekJPUkVrelYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVV2U1M5dWFWbElOMkZ4T0VGcGRUVTVUMDlUVTBob00yWm1XblpSZVdaNlNrMUpkbWdLWVRoV1VYcDFTSFZpVG1Kc1QzZ3pTMmRTTlZKUmJucEtVVkozYjFnNFEyRnRXVXR1WmsxdWVGWkVkMnRMYzI1ME5HRlBRMEptUlhkbloxaDBUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZFV0RFMENteFhibTg1ZVRSU04zUnhkME5FUlU5blRGQm5XR3M0ZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOU9SR2N5V1dwTk1GbDZhR2xaZW1ONVRrZFplRTlVUVRGYWJWRjZDbHB0VW1wTlJGSnBUbFJyTlUxVVFUSmFha1pwVGpKT2JWcEVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMDVFWnpKWmFrMHdXWHBvYVZsNlkzbE9SMWw0VDFSQk1WcHRVWHBhYlZKcVRVUlNhVTVVYXpVS1RWUkJNbHBxUm1sT01rNXRXa1JCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMDVFWnpKWmFrMHdXWHBvYVZsNlkzbE9SMWw0VDFSQk1WcHRVWG9LV20xU2FrMUVVbWxPVkdzMVRWUkJNbHBxUm1sT01rNXRXa1JCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwUlJORTV0U1hvS1RrZE5ORmx0VFROTmFsSnRUVlJyZDA1WFdtdE5NbHByV1hwQk1GbHFWVFZQVkVWM1RtMVplRmxxWkdwYWJWRjNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWxFeFQwUlJNVTE2WTNoT1ZFVjJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwZDFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT1VKSWMwRUtaVkZDTTBGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUpMMUpuZEc5QlFVRlJSQXBCUldkM1VtZEphRUZLUTNSclRFOVVORFJIVnpnM05GaHlUbFV2V21wbU9WWkhiMFpuZUVkSU5sSndRbTlKVVd4a1pESmFRV2xGUVRWYVIweGpPRU4zQ2xab2RHVlZSbE5RWWtKM2RGcHhURlpHVkRSNk5reFZMemhaZVc5eU5FdzBXbXROZDBObldVbExiMXBKZW1vd1JVRjNUVVJoUVVGM1dsRkpkMUJLWVRRS09VdGpjM1p5YzNoalNVWnlOMmd4TmpVMGRYTlBia1pMTVc5SmEyUjRSVXRMVlhKcmNsUk9VakZpYW1sMFQxQnlNR0ZVSzJkaWJsZDVja1F6UVdwRlFRcHZMMVJQVnpSd1RYTnlhR0ZRWjA5bGVWSjFRa1pUZVhRM0t6UjRZbWRYZFZWck9HWnNkalJ5U21saVdEWk5TR3N2WlhwWk0ydEtkMnhJVWxwM05XWk5DaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyjADAgEAMIICwQYJKoZIhvcNAQcCoIICsjCCAq4CAQMxDTALBglghkgBZQMEAgEwgbcGCyqGSIb3DQEJEAEEoIGnBIGkMIGhAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgFAvgVDvYoGUz36ZsrrR+oTzdDojGBOfZOY2FcqgzTvsCFCwuwm6L5A8bOhvhfw5x2WvFF9AJGA8yMDI2MDkxMTA5MzQyN1owAwIBAaAypDAwLjEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MRUwEwYDVQQDEwxzaWdzdG9yZS10c2GgADGCAdwwggHYAgEBMFEwOTEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MSAwHgYDVQQDExdzaWdzdG9yZS10c2Etc2VsZnNpZ25lZAIUOhNULwyQYe68wUMvy4qOiyojiwwwCwYJYIZIAWUDBAIBoIH8MBoGCSqGSIb3DQEJAzENBgsqhkiG9w0BCRABBDAcBgkqhkiG9w0BCQUxDxcNMjYwOTExMDkzNDI3WjAvBgkqhkiG9w0BCQQxIgQgORls50DJBhFi3Jv/9cjXP4GHNdWywmOi1SUtTieXWmYwgY4GCyqGSIb3DQEJEAIvMX8wfTB7MHkEIIX5J7wHq2LKw7RDVsEO/IGyxog/2nq55thw2dE6zQW3MFUwPaQ7MDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAoGCCqGSM49BAMCBGgwZgIxANEJy5RSrvF/5wQd5+bZ4C63AhPV0uPTOG/OnwIAMi4HgZ3WTXDWrJQYL8qXstB7jwIxAIBhi64lObeTLrzKsgANqpXj5NFR4PU86ptb/pun/9DepAFOWDs85y8e9qlJmfgYjQ=="}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"eXm5Jj3IyVqWziT8N03TuGP+C0qeyFLyswpvkvc2/XU="}, "signature":"MEUCIQDDcA7+8fIdyo6Dx+/yq1aZFSOvqfn1+jj+pRzgwNOD3AIgbLDKd/FTTBDNJgeEEeAlDOTFB/cXIDTOlPTYLQXpWJs="}}
|
build/torch-xpu/ops/common/backends/intracard.py
CHANGED
|
@@ -82,7 +82,7 @@ class IntraCardCPBackend(BaseBackend):
|
|
| 82 |
chunk_indices: torch.LongTensor | None = None,
|
| 83 |
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor | None]:
|
| 84 |
"""Intra-card CP implementation of chunk_gated_delta_rule_fwd_h."""
|
| 85 |
-
from ...ops.common.intracard_cp import intracard_fwd_h
|
| 86 |
|
| 87 |
return intracard_fwd_h(
|
| 88 |
k=k, w=w, u=u, g=g, gk=gk,
|
|
|
|
| 82 |
chunk_indices: torch.LongTensor | None = None,
|
| 83 |
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor | None]:
|
| 84 |
"""Intra-card CP implementation of chunk_gated_delta_rule_fwd_h."""
|
| 85 |
+
from ....ops.common.intracard_cp import intracard_fwd_h
|
| 86 |
|
| 87 |
return intracard_fwd_h(
|
| 88 |
k=k, w=w, u=u, g=g, gk=gk,
|
build/torch-xpu/ops/common/backends/tilelang/__init__.py
CHANGED
|
@@ -96,7 +96,7 @@ class TileLangBackend(BaseBackend):
|
|
| 96 |
chunk_size: int = 64,
|
| 97 |
chunk_indices: torch.LongTensor | None = None,
|
| 98 |
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor | None, torch.Tensor | None]:
|
| 99 |
-
from ...ops.common.backends.tilelang.chunk_bwd import (
|
| 100 |
chunk_bwd_dqkwg_tilelang,
|
| 101 |
)
|
| 102 |
return chunk_bwd_dqkwg_tilelang(
|
|
@@ -145,7 +145,7 @@ class TileLangBackend(BaseBackend):
|
|
| 145 |
cu_seqlens: torch.LongTensor | None = None,
|
| 146 |
chunk_indices: torch.LongTensor | None = None,
|
| 147 |
) -> tuple[torch.Tensor, torch.Tensor]:
|
| 148 |
-
from ...ops.common.backends.tilelang.parallel_attn_fwd import (
|
| 149 |
parallel_attn_fwd_tilelang,
|
| 150 |
)
|
| 151 |
return parallel_attn_fwd_tilelang(
|
|
@@ -190,7 +190,7 @@ class TileLangBackend(BaseBackend):
|
|
| 190 |
cu_seqlens: torch.LongTensor | None = None,
|
| 191 |
chunk_indices: torch.LongTensor | None = None,
|
| 192 |
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor | None, torch.Tensor | None]:
|
| 193 |
-
from ...ops.common.backends.tilelang.parallel_attn_bwd import (
|
| 194 |
parallel_attn_bwd_tilelang,
|
| 195 |
)
|
| 196 |
return parallel_attn_bwd_tilelang(
|
|
|
|
| 96 |
chunk_size: int = 64,
|
| 97 |
chunk_indices: torch.LongTensor | None = None,
|
| 98 |
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor | None, torch.Tensor | None]:
|
| 99 |
+
from .....ops.common.backends.tilelang.chunk_bwd import (
|
| 100 |
chunk_bwd_dqkwg_tilelang,
|
| 101 |
)
|
| 102 |
return chunk_bwd_dqkwg_tilelang(
|
|
|
|
| 145 |
cu_seqlens: torch.LongTensor | None = None,
|
| 146 |
chunk_indices: torch.LongTensor | None = None,
|
| 147 |
) -> tuple[torch.Tensor, torch.Tensor]:
|
| 148 |
+
from .....ops.common.backends.tilelang.parallel_attn_fwd import (
|
| 149 |
parallel_attn_fwd_tilelang,
|
| 150 |
)
|
| 151 |
return parallel_attn_fwd_tilelang(
|
|
|
|
| 190 |
cu_seqlens: torch.LongTensor | None = None,
|
| 191 |
chunk_indices: torch.LongTensor | None = None,
|
| 192 |
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor | None, torch.Tensor | None]:
|
| 193 |
+
from .....ops.common.backends.tilelang.parallel_attn_bwd import (
|
| 194 |
parallel_attn_bwd_tilelang,
|
| 195 |
)
|
| 196 |
return parallel_attn_bwd_tilelang(
|
build/torch-xpu/ops/common/backends/tilelang/parallel_attn_bwd.py
CHANGED
|
@@ -254,7 +254,7 @@ def parallel_attn_bwd_tilelang(
|
|
| 254 |
chunk_indices = prepare_chunk_indices(cu_seqlens, BT)
|
| 255 |
|
| 256 |
# reuse Triton preprocess for delta (works with unpadded tensors)
|
| 257 |
-
from ...ops.attn.parallel import parallel_attn_bwd_preprocess
|
| 258 |
delta = parallel_attn_bwd_preprocess(o, do)
|
| 259 |
|
| 260 |
# No padding on the last dim: TileLang T.copy auto-zero-fills OOB columns
|
|
|
|
| 254 |
chunk_indices = prepare_chunk_indices(cu_seqlens, BT)
|
| 255 |
|
| 256 |
# reuse Triton preprocess for delta (works with unpadded tensors)
|
| 257 |
+
from .....ops.attn.parallel import parallel_attn_bwd_preprocess
|
| 258 |
delta = parallel_attn_bwd_preprocess(o, do)
|
| 259 |
|
| 260 |
# No padding on the last dim: TileLang T.copy auto-zero-fills OOB columns
|
build/torch-xpu/ops/deltaformer/parallel.py
CHANGED
|
@@ -23,7 +23,7 @@ except ImportError:
|
|
| 23 |
)
|
| 24 |
flash_attn_func = None
|
| 25 |
|
| 26 |
-
from ...
|
| 27 |
|
| 28 |
BLOCK_SIZE_C = 512
|
| 29 |
|
|
|
|
| 23 |
)
|
| 24 |
flash_attn_func = None
|
| 25 |
|
| 26 |
+
from ...utils.layers import pad_input, unpad_input
|
| 27 |
|
| 28 |
BLOCK_SIZE_C = 512
|
| 29 |
|
build/torch-xpu/ops/kda/backends/flashkda.py
CHANGED
|
@@ -18,7 +18,7 @@ import torch
|
|
| 18 |
from ....ops.backends import BaseBackend
|
| 19 |
|
| 20 |
if TYPE_CHECKING:
|
| 21 |
-
from ...ops.cp import FLACPContext
|
| 22 |
|
| 23 |
|
| 24 |
class FlashKDABackend(BaseBackend):
|
|
|
|
| 18 |
from ....ops.backends import BaseBackend
|
| 19 |
|
| 20 |
if TYPE_CHECKING:
|
| 21 |
+
from ....ops.cp import FLACPContext
|
| 22 |
|
| 23 |
|
| 24 |
class FlashKDABackend(BaseBackend):
|
build/torch-xpu/ops/kda/backends/tilelang/__init__.py
CHANGED
|
@@ -72,7 +72,7 @@ class KDATileLangBackend(BaseBackend):
|
|
| 72 |
chunk_size: int = 64,
|
| 73 |
chunk_indices: torch.LongTensor | None = None,
|
| 74 |
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
|
| 75 |
-
from ...ops.kda.backends.tilelang.chunk_bwd_dqkg import (
|
| 76 |
chunk_kda_bwd_wy_dqkg_fused_tilelang,
|
| 77 |
)
|
| 78 |
return chunk_kda_bwd_wy_dqkg_fused_tilelang(
|
|
|
|
| 72 |
chunk_size: int = 64,
|
| 73 |
chunk_indices: torch.LongTensor | None = None,
|
| 74 |
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
|
| 75 |
+
from .....ops.kda.backends.tilelang.chunk_bwd_dqkg import (
|
| 76 |
chunk_kda_bwd_wy_dqkg_fused_tilelang,
|
| 77 |
)
|
| 78 |
return chunk_kda_bwd_wy_dqkg_fused_tilelang(
|
build/torch-xpu/utils/layers.py
ADDED
|
@@ -0,0 +1,332 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright (c) 2023-2026, Songlin Yang, Yu Zhang, Zhiyuan Li
|
| 2 |
+
#
|
| 3 |
+
# This source code is licensed under the MIT license found in the
|
| 4 |
+
# LICENSE file in the root directory of this source tree.
|
| 5 |
+
# For a list of all contributors, visit:
|
| 6 |
+
# https://github.com/fla-org/flash-linear-attention/graphs/contributors
|
| 7 |
+
|
| 8 |
+
# Code is adapted from flash-attn.bert_padding.py
|
| 9 |
+
|
| 10 |
+
# NOTE: This is a copy of `fla.layers.utils` but we do not sync the FLA layers in hf-kernels
|
| 11 |
+
|
| 12 |
+
import torch
|
| 13 |
+
from einops import rearrange, repeat
|
| 14 |
+
|
| 15 |
+
from ..ops.utils.index import (
|
| 16 |
+
prepare_cu_seqlens_from_lens,
|
| 17 |
+
prepare_cu_seqlens_from_mask,
|
| 18 |
+
prepare_lens_from_mask,
|
| 19 |
+
)
|
| 20 |
+
from . import tensor_cache
|
| 21 |
+
|
| 22 |
+
_LAYER_IDX_REQUIRED_MSG = "{cls} requires `layer_idx` when `past_key_values` is provided."
|
| 23 |
+
|
| 24 |
+
|
| 25 |
+
class IndexFirstAxis(torch.autograd.Function):
|
| 26 |
+
|
| 27 |
+
@staticmethod
|
| 28 |
+
def forward(ctx, x, indices):
|
| 29 |
+
ctx.save_for_backward(indices)
|
| 30 |
+
assert x.ndim >= 2
|
| 31 |
+
ctx.first_axis_dim, other_shape = x.shape[0], x.shape[1:]
|
| 32 |
+
second_dim = other_shape.numel()
|
| 33 |
+
# TD [2022-03-04] For some reason torch.gather is a bit faster than indexing.
|
| 34 |
+
# return x[indices]
|
| 35 |
+
return torch.gather(
|
| 36 |
+
rearrange(x, "b ... -> b (...)"), 0, repeat(indices, "z -> z d", d=second_dim),
|
| 37 |
+
).reshape(-1, *other_shape)
|
| 38 |
+
|
| 39 |
+
@staticmethod
|
| 40 |
+
def backward(ctx, do):
|
| 41 |
+
(indices,) = ctx.saved_tensors
|
| 42 |
+
assert do.ndim >= 2
|
| 43 |
+
other_shape = do.shape[1:]
|
| 44 |
+
do = rearrange(do, "b ... -> b (...)")
|
| 45 |
+
dx = torch.zeros(
|
| 46 |
+
[ctx.first_axis_dim, do.shape[1]],
|
| 47 |
+
device=do.device,
|
| 48 |
+
dtype=do.dtype,
|
| 49 |
+
)
|
| 50 |
+
# TD [2022-03-04] For some reason torch.scatter is a bit faster than indexing.
|
| 51 |
+
# dx[indices] = do
|
| 52 |
+
dx.scatter_(0, repeat(indices, "z -> z d", d=do.shape[1]), do)
|
| 53 |
+
return dx.reshape(ctx.first_axis_dim, *other_shape), None
|
| 54 |
+
|
| 55 |
+
|
| 56 |
+
index_first_axis = IndexFirstAxis.apply
|
| 57 |
+
|
| 58 |
+
|
| 59 |
+
class IndexPutFirstAxis(torch.autograd.Function):
|
| 60 |
+
|
| 61 |
+
@staticmethod
|
| 62 |
+
def forward(ctx, x, indices, first_axis_dim):
|
| 63 |
+
ctx.save_for_backward(indices)
|
| 64 |
+
assert indices.ndim == 1
|
| 65 |
+
assert x.ndim >= 2
|
| 66 |
+
y = torch.zeros(first_axis_dim, *x.shape[1:], device=x.device, dtype=x.dtype)
|
| 67 |
+
# TODO [2022-03-04] For some reason torch.scatter is a bit faster than indexing.
|
| 68 |
+
y[indices] = x
|
| 69 |
+
# y.scatter_(0, repeat(indices, 'z -> z d', d=x.shape[1]), x)
|
| 70 |
+
return y
|
| 71 |
+
|
| 72 |
+
@staticmethod
|
| 73 |
+
def backward(ctx, do):
|
| 74 |
+
(indices,) = ctx.saved_tensors
|
| 75 |
+
# TODO [2022-03-04] For some reason torch.gather is a bit faster than indexing.
|
| 76 |
+
dx = do[indices]
|
| 77 |
+
# dx = torch.gather(do, 0, repeat(indices, 'z -> z d', d=do.shape[1]))
|
| 78 |
+
return dx, None, None
|
| 79 |
+
|
| 80 |
+
|
| 81 |
+
index_put_first_axis = IndexPutFirstAxis.apply
|
| 82 |
+
|
| 83 |
+
|
| 84 |
+
@tensor_cache
|
| 85 |
+
def _get_unpad_indices_and_cu(
|
| 86 |
+
attention_mask: torch.Tensor,
|
| 87 |
+
) -> tuple[torch.Tensor, torch.Tensor]:
|
| 88 |
+
lens = prepare_lens_from_mask(attention_mask)
|
| 89 |
+
indices = torch.nonzero(attention_mask.flatten(), as_tuple=False).flatten()
|
| 90 |
+
cu_seqlens = prepare_cu_seqlens_from_lens(lens)
|
| 91 |
+
return indices, cu_seqlens
|
| 92 |
+
|
| 93 |
+
|
| 94 |
+
@tensor_cache
|
| 95 |
+
def _get_last_token_unpad_indices_and_cu(
|
| 96 |
+
attention_mask: torch.Tensor,
|
| 97 |
+
) -> tuple[torch.Tensor, torch.Tensor]:
|
| 98 |
+
return _get_unpad_indices_and_cu(attention_mask[:, -1:])
|
| 99 |
+
|
| 100 |
+
|
| 101 |
+
def get_unpad_indices_and_cu(
|
| 102 |
+
attention_mask: torch.Tensor,
|
| 103 |
+
q_len: int | None = None,
|
| 104 |
+
) -> tuple[torch.Tensor, torch.Tensor]:
|
| 105 |
+
"""Return unpadding indices and cumulative sequence lengths."""
|
| 106 |
+
if q_len is None or attention_mask.shape[-1] == q_len:
|
| 107 |
+
return _get_unpad_indices_and_cu(attention_mask)
|
| 108 |
+
if q_len == 1:
|
| 109 |
+
return _get_last_token_unpad_indices_and_cu(attention_mask)
|
| 110 |
+
return _get_unpad_indices_and_cu(attention_mask[:, -q_len:])
|
| 111 |
+
|
| 112 |
+
|
| 113 |
+
@tensor_cache
|
| 114 |
+
def get_unpad_data(
|
| 115 |
+
attention_mask: torch.Tensor,
|
| 116 |
+
) -> tuple[torch.Tensor, torch.Tensor, int]:
|
| 117 |
+
"""
|
| 118 |
+
Retrieves indexing data required to repad unpadded (ragged) tensors.
|
| 119 |
+
|
| 120 |
+
Args:
|
| 121 |
+
attention_mask (`torch.Tensor`):
|
| 122 |
+
Boolean or int tensor of shape (batch_size, sequence_length), 1 means valid and 0 means not valid.
|
| 123 |
+
|
| 124 |
+
Return:
|
| 125 |
+
indices (`torch.Tensor`):
|
| 126 |
+
The indices of non-masked tokens from the flattened input sequence.
|
| 127 |
+
cu_seqlens (`torch.Tensor`):
|
| 128 |
+
The cumulative sequence lengths, used to index into ragged (unpadded) tensors.
|
| 129 |
+
`cu_seqlens` shape is [batch_size + 1].
|
| 130 |
+
max_seqlen_in_batch (`int`):
|
| 131 |
+
Maximum sequence length in batch.
|
| 132 |
+
"""
|
| 133 |
+
lens = prepare_lens_from_mask(attention_mask)
|
| 134 |
+
indices = torch.nonzero(attention_mask.flatten(), as_tuple=False).flatten()
|
| 135 |
+
max_seqlen_in_batch = lens.max().item()
|
| 136 |
+
cu_seqlens = prepare_cu_seqlens_from_mask(attention_mask)
|
| 137 |
+
return indices, cu_seqlens, max_seqlen_in_batch
|
| 138 |
+
|
| 139 |
+
|
| 140 |
+
def unpad_hidden_states(
|
| 141 |
+
hidden_states: torch.Tensor,
|
| 142 |
+
cu_seqlens: torch.LongTensor | None,
|
| 143 |
+
attention_mask: torch.Tensor | None,
|
| 144 |
+
q_len: int,
|
| 145 |
+
) -> tuple[torch.Tensor, torch.Tensor | None, torch.LongTensor | None]:
|
| 146 |
+
"""Unpad hidden states via attention_mask, unless explicit cu_seqlens are given.
|
| 147 |
+
|
| 148 |
+
Explicit `cu_seqlens` take precedence: when they are provided, the input is
|
| 149 |
+
returned unchanged and no unpadding happens. Only when `cu_seqlens` is None
|
| 150 |
+
and `attention_mask` is present are the hidden states unpadded via the mask.
|
| 151 |
+
|
| 152 |
+
Arguments:
|
| 153 |
+
hidden_states (`torch.Tensor`):
|
| 154 |
+
Input hidden states with padding. Shape: [batch_size, q_len, ...].
|
| 155 |
+
cu_seqlens (`torch.LongTensor | None`):
|
| 156 |
+
Explicit cumulative sequence lengths provided by the caller.
|
| 157 |
+
attention_mask (`torch.Tensor | None`):
|
| 158 |
+
Boolean or int tensor of shape (batch_size, sequence_length), 1 means valid and 0 means not valid.
|
| 159 |
+
q_len (`int`):
|
| 160 |
+
Target length.
|
| 161 |
+
|
| 162 |
+
Return:
|
| 163 |
+
hidden_states (`torch.Tensor`):
|
| 164 |
+
The unpadded hidden states of shape [1, total_tokens, ...],
|
| 165 |
+
or the input unchanged when unpadding is skipped.
|
| 166 |
+
indices (`torch.Tensor | None`):
|
| 167 |
+
The indices of non-masked tokens, or None when unpadding is skipped.
|
| 168 |
+
cu_seqlens (`torch.LongTensor | None`):
|
| 169 |
+
The cumulative sequence lengths derived from the mask,
|
| 170 |
+
or the explicit value passed in when unpadding is skipped.
|
| 171 |
+
"""
|
| 172 |
+
if cu_seqlens is None and attention_mask is not None:
|
| 173 |
+
indices, cu_seqlens = get_unpad_indices_and_cu(attention_mask, q_len)
|
| 174 |
+
hidden_states = index_first_axis(rearrange(hidden_states, "b s ... -> (b s) ..."), indices).unsqueeze(0)
|
| 175 |
+
return hidden_states, indices, cu_seqlens
|
| 176 |
+
return hidden_states, None, cu_seqlens
|
| 177 |
+
|
| 178 |
+
|
| 179 |
+
def repad_hidden_states(
|
| 180 |
+
hidden_states: torch.Tensor,
|
| 181 |
+
indices: torch.Tensor | None,
|
| 182 |
+
batch_size: int,
|
| 183 |
+
q_len: int,
|
| 184 |
+
) -> torch.Tensor:
|
| 185 |
+
"""Inverse of `unpad_hidden_states`; a no-op when `indices` is None.
|
| 186 |
+
|
| 187 |
+
Arguments:
|
| 188 |
+
hidden_states (`torch.Tensor`):
|
| 189 |
+
Hidden states of shape [1, total_tokens, ...].
|
| 190 |
+
indices (`torch.Tensor | None`):
|
| 191 |
+
The indices returned by `unpad_hidden_states`, or None.
|
| 192 |
+
batch_size (`int`):
|
| 193 |
+
Batch size for the padded sequence.
|
| 194 |
+
q_len (`int`):
|
| 195 |
+
Target length.
|
| 196 |
+
|
| 197 |
+
Return:
|
| 198 |
+
hidden_states of shape [batch_size, q_len, ...],
|
| 199 |
+
or the input unchanged when `indices` is None.
|
| 200 |
+
"""
|
| 201 |
+
if indices is None:
|
| 202 |
+
return hidden_states
|
| 203 |
+
return pad_input(hidden_states.squeeze(0), indices, batch_size, q_len)
|
| 204 |
+
|
| 205 |
+
|
| 206 |
+
def unpad_input(
|
| 207 |
+
q: torch.Tensor | tuple[torch.Tensor, ...],
|
| 208 |
+
states: tuple[torch.Tensor],
|
| 209 |
+
attention_mask: torch.Tensor,
|
| 210 |
+
q_len: int,
|
| 211 |
+
keepdim: bool = False,
|
| 212 |
+
):
|
| 213 |
+
"""
|
| 214 |
+
Unpads query, key, and values tensors, using a single dimension for all tokens
|
| 215 |
+
even though they belong to different batches.
|
| 216 |
+
|
| 217 |
+
|
| 218 |
+
Arguments:
|
| 219 |
+
q (`torch.Tensor` or `Tuple[torch.Tensor]`):
|
| 220 |
+
Query state with padding. Shape: [batch_size, q_len, ...].
|
| 221 |
+
When it is a tuple, do unpadding for each tensor in the tuple.
|
| 222 |
+
states (`Tuple[torch.Tensor]`):
|
| 223 |
+
Attention state with padding. Shape: [batch_size, seq_len, ...].
|
| 224 |
+
attention_mask (`torch.Tensor`):
|
| 225 |
+
Boolean or int tensor of shape [batch_size, sequence_length], 1 means valid and 0 means not valid.
|
| 226 |
+
q_len (`int`):
|
| 227 |
+
Target length.
|
| 228 |
+
keepdim (`bool`):
|
| 229 |
+
Whether to keep the batch dimension. Default: `False`.
|
| 230 |
+
|
| 231 |
+
Return:
|
| 232 |
+
q (`torch.Tensor` or `Tuple[torch.Tensor]`):
|
| 233 |
+
Query state without padding.
|
| 234 |
+
Shape: [1, total_target_length, ...] if `keepdim=True` else [total_target_length, ...].
|
| 235 |
+
When the `q` passed in is a tuple, return a tuple of such unpadded tensors.
|
| 236 |
+
states (`Tuple[torch.Tensor]`):
|
| 237 |
+
Attention state without padding.
|
| 238 |
+
Shape: [1, total_source_length, ...] if `keepdim=True` else [total_source_length, ...].
|
| 239 |
+
indices_q (`torch.Tensor`):
|
| 240 |
+
The indices of non-masked tokens from the flattened input target sequence.
|
| 241 |
+
(cu_seqlens_q, cu_seqlens_k) (`Tuple[torch.LongTensor, torch.LongTensor]`):
|
| 242 |
+
The cumulative sequence lengths for the target (query) and source (key, value),
|
| 243 |
+
used to index into ragged (unpadded) tensors.
|
| 244 |
+
`cu_seqlens` shape is [batch_size + 1].
|
| 245 |
+
(max_seqlen_in_batch_q, max_seqlen_in_batch_k) (`Tuple[int, int]`):
|
| 246 |
+
Maximum sequence length in batch (`max_seqlen_in_batch_q` for the target sequence
|
| 247 |
+
i.e. query, `max_seqlen_in_batch_k` for the source sequence i.e. key/value).
|
| 248 |
+
"""
|
| 249 |
+
indices_k, cu_seqlens_k, max_seqlen_in_batch_k = get_unpad_data(attention_mask)
|
| 250 |
+
batch_size, seq_len, *_ = states[0].shape
|
| 251 |
+
|
| 252 |
+
state = tuple(
|
| 253 |
+
index_first_axis(rearrange(s, "b s ... -> (b s) ..."), indices_k)
|
| 254 |
+
for s in states
|
| 255 |
+
)
|
| 256 |
+
if isinstance(q, torch.Tensor):
|
| 257 |
+
q = (q,)
|
| 258 |
+
cast_tuple = True
|
| 259 |
+
else:
|
| 260 |
+
cast_tuple = False
|
| 261 |
+
|
| 262 |
+
if q_len == seq_len:
|
| 263 |
+
q = tuple(index_first_axis(rearrange(q_, "b s ... -> (b s) ..."), indices_k) for q_ in q)
|
| 264 |
+
cu_seqlens_q = cu_seqlens_k
|
| 265 |
+
max_seqlen_in_batch_q = max_seqlen_in_batch_k
|
| 266 |
+
indices_q = indices_k
|
| 267 |
+
elif q_len == 1:
|
| 268 |
+
max_seqlen_in_batch_q = 1
|
| 269 |
+
cu_seqlens_q = torch.arange(batch_size + 1, dtype=torch.int32, device=q[0].device)
|
| 270 |
+
indices_q = cu_seqlens_q[:-1]
|
| 271 |
+
q = tuple(q_.squeeze(1) for q_ in q)
|
| 272 |
+
else:
|
| 273 |
+
raise NotImplementedError("We only support either q_len == k_len (prefilling) or q_len == 1 (decoding)")
|
| 274 |
+
|
| 275 |
+
if keepdim:
|
| 276 |
+
q = tuple(q_.unsqueeze(0) for q_ in q)
|
| 277 |
+
state = tuple(s.unsqueeze(0) for s in state)
|
| 278 |
+
if cast_tuple:
|
| 279 |
+
q = q[0]
|
| 280 |
+
|
| 281 |
+
return (
|
| 282 |
+
q,
|
| 283 |
+
state,
|
| 284 |
+
indices_q,
|
| 285 |
+
(cu_seqlens_q, cu_seqlens_k),
|
| 286 |
+
(max_seqlen_in_batch_q, max_seqlen_in_batch_k),
|
| 287 |
+
)
|
| 288 |
+
|
| 289 |
+
|
| 290 |
+
def pad_input(
|
| 291 |
+
hidden_states: torch.Tensor,
|
| 292 |
+
indices: torch.LongTensor,
|
| 293 |
+
batch_size: int,
|
| 294 |
+
seq_len: int,
|
| 295 |
+
) -> torch.Tensor:
|
| 296 |
+
"""
|
| 297 |
+
Args:
|
| 298 |
+
hidden_states ([total_tokens, ...]):
|
| 299 |
+
where total_tokens denotes the number of tokens in selected in attention_mask.
|
| 300 |
+
indices ([total_tokens]):
|
| 301 |
+
the indices that represent the non-masked tokens of the original padded input sequence.
|
| 302 |
+
batch_size (int):
|
| 303 |
+
batch_size size for the padded sequence.
|
| 304 |
+
seq_len (int):
|
| 305 |
+
maximum sequence length for the padded sequence.
|
| 306 |
+
|
| 307 |
+
Return:
|
| 308 |
+
hidden_states of shape [batch_size, seq_len, ...]
|
| 309 |
+
"""
|
| 310 |
+
output = index_put_first_axis(hidden_states, indices, batch_size * seq_len)
|
| 311 |
+
return rearrange(output, "(b s) ... -> b s ...", b=batch_size)
|
| 312 |
+
|
| 313 |
+
|
| 314 |
+
def require_cache_layer_idx(module, past_key_values):
|
| 315 |
+
layer_idx = getattr(module, "layer_idx", None)
|
| 316 |
+
if past_key_values is not None and layer_idx is None:
|
| 317 |
+
raise ValueError(_LAYER_IDX_REQUIRED_MSG.format(cls=module.__class__.__name__))
|
| 318 |
+
return layer_idx
|
| 319 |
+
|
| 320 |
+
|
| 321 |
+
def get_layer_cache(module, past_key_values):
|
| 322 |
+
layer_idx = require_cache_layer_idx(module, past_key_values)
|
| 323 |
+
if past_key_values is not None and len(past_key_values) > layer_idx:
|
| 324 |
+
return past_key_values[layer_idx]
|
| 325 |
+
return None
|
| 326 |
+
|
| 327 |
+
|
| 328 |
+
def update_layer_cache(module, past_key_values, **kwargs):
|
| 329 |
+
layer_idx = require_cache_layer_idx(module, past_key_values)
|
| 330 |
+
if past_key_values is not None:
|
| 331 |
+
return past_key_values.update(layer_idx=layer_idx, **kwargs)
|
| 332 |
+
return None
|