Kernels:
Trusted publisher
Uploaded using `kernel-builder`.
Browse files- build/torch-cuda/_ops.py +1 -1
- build/torch-cuda/metadata.json +11 -10
- build/torch-cuda/metadata.json.sigstore +1 -1
- build/torch-cuda/ops/common/backends/intracard.py +1 -1
- build/torch-cuda/ops/common/backends/tilelang/__init__.py +3 -3
- build/torch-cuda/ops/common/backends/tilelang/parallel_attn_bwd.py +1 -1
- build/torch-cuda/ops/deltaformer/parallel.py +1 -1
- build/torch-cuda/ops/kda/backends/flashkda.py +1 -1
- build/torch-cuda/ops/kda/backends/tilelang/__init__.py +1 -1
- build/torch-cuda/utils/layers.py +332 -0
build/torch-cuda/_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-cuda/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_cuda_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-cuda/metadata.json.sigstore
CHANGED
|
@@ -1 +1 @@
|
|
| 1 |
-
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSjCCBtGgAwIBAgIUAfIxdBmJAkWnIir4TqaIj8X6i6EwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwOTEwMTE0OTM0WhcNMjYwOTEwMTE1OTM0WjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEA1th0kFcSYl5CidY3vZ+bglLs1ETBvN8auCsizhTX4GhrqfpLQP4PvzNyHzoXO7xv5Zf6PrKxEtBq0j3mFrPOqOCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUTfpG7pEKsHtcB7m6CMpLBhMK2w0wHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoNTE2MjI1MWY3MjhhNmRhMWUyNDg2NTRiYjgyODdhZTczYTg2NmU5YjATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoNTE2MjI1MWY3MjhhNmRhMWUyNDg2NTRiYjgyODdhZTczYTg2NmU5YjAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoNTE2MjI1MWY3MjhhNmRhMWUyNDg2NTRiYjgyODdhZTczYTg2NmU5YjAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKDUxNjIyNTFmNzI4YTZkYTFlMjQ4NjU0YmI4Mjg3YWU3M2E4NjZlOWIwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzQ0NzI5Njc1NjAvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoIsm2/kAAAQDAEcwRQIgJDUKR1/KP5BbF+m+pNr2aaNggZwN1+iN4Pn+bHZ81a0CIQDTol7sGGY/YlDttLCuYY1BjElz5fcHArW4vGvOZTm9JTAKBggqhkjOPQQDAwNnADBkAjA2Thr0Xxh5WkPGNlKGgn+qI82QxVftzFVdm13wpRiuyGi51vhxoVITlhjy9wZgK+ICMGrtuKuFGiyEzj0wS/Ak7qupFRsMdbh0qpVdHzYXkeuQMEzhQHV60QlkuvOMpyh6TA=="}, "tlogEntries":[{"logIndex":"2782561056", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1789040977", "inclusionPromise":{"signedEntryTimestamp":"MEUCIQCT3Pugf26BhtG4Tt/O6moKE4cU/xGUOqArDDqtPzL1PQIgVnWg/8GkCEfrmpr7ca5dfgcX+Qz/2XXKs2DeNa+2/Zo="}, "inclusionProof":{"logIndex":"2660656794", "rootHash":"N2psO4ZFh6veACXIA+JwMmBHRvveg93UUr6viiHgCOg=", "treeSize":"2660656951", "hashes":["ah8zFCy6HcTzmbroSlrwyjpfDETR4euwEq9Zas5NBe8=", "fOL58qM6jQF+l7N0zoQS/SOZKWAt6JFJPXIC4F4Yud0=", "hb15UG+koSTR1S+FZX+Nrm9nOApr9FIG5dCnq06zhzk=", "bZrve1jTrW7MwiilNYiy1YDIfZOtn60RCMpXezIqePU=", "jCFBd2R/VjH6mPAYmXWxdrdjQVtFoL6tYs/c/Hdx7x8=", "XfHe6DYP5WfkB7CFt84v7uzSxr6SsHhZRXGicPPVYIQ=", "xx0KSiOIdErhKcwSAdP/imfnkXcQqfU6/M3/oyEnmTY=", "XAQ+oXEjJSnVoIz6BAmmDr8AerALQkqb9VVHi7qzj/A=", "p6QDR3X5qkqLe0ZpXdKk9TKw1Bkx3xE3X3BRIPqIVRA=", "TSXKS1ABF1IwZnHcWB2NpY67qUpCZ/agLECgdpGf8JQ=", "Fp1ZUtwb5aOGNNySeAWbtDDxZttr2iJ0WJr/xg06vwg=", "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\n2660656951\nN2psO4ZFh6veACXIA+JwMmBHRvveg93UUr6viiHgCOg=\n\n— rekor.sigstore.dev wNI9ajBEAiBH24jujNgCiU872D0kP+OGFx36Bi3z/XYNXcsM8BzdPAIgYe1Eaf0u0ramwCbW4BxswdRbJT6oUMpyjRQhMJLwBAY=\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiI1OWVlYmMxMGVkYjhjYmE5NmNmMWM0NjVjNTM4YmEwOWM5ZWU1NGU1MWYzNmUwNzljYzMxMzhkNmE1OTdjZThkIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJUUQ0Y1pJYU4rWitwSkd0dUlwSXdTNmZnRmV4OVdvTi9GTGl3UStUTjFCRnd3SWdad1lsK3d3cmFaQ1E2SWNLMnNEWWZaeUh5RWdCL3AzTVFoMHhvTXQzSklzPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRha05EUW5SSFowRjNTVUpCWjBsVlFXWkplR1JDYlVwQmExZHVTV2x5TkZSeFlVbHFPRmcyYVRaRmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlVUlhkTlZFVXdUMVJOTUZkb1kwNU5hbGwzVDFSRmQwMVVSVEZQVkUwd1YycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZCTVhSb01HdEdZMU5aYkRWRGFXUlpNM1phSzJKbmJFeHpNVVZVUW5aT09HRjFRM01LYVhwb1ZGZzBSMmh5Y1dad1RGRlFORkIyZWs1NVNIcHZXRTgzZUhZMVdtWTJVSEpMZUVWMFFuRXdhak50Um5KUVQzRlBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZVWm5CSENqZHdSVXR6U0hSalFqZHROa05OY0V4Q2FFMUxNbmN3ZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOU9WRVV5VFdwSk1VMVhXVE5OYW1ob1RtMVNhRTFYVlhsT1JHY3lDazVVVW1sWmFtZDVUMFJrYUZwVVkzcFpWR2N5VG0xVk5WbHFRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMDVVUlRKTmFra3hUVmRaTTAxcWFHaE9iVkpvVFZkVmVVNUVaekpPVkZKcFdXcG5lVTlFWkdnS1dsUmplbGxVWnpKT2JWVTFXV3BCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMDVVUlRKTmFra3hUVmRaTTAxcWFHaE9iVkpvVFZkVmVVNUVaeklLVGxSU2FWbHFaM2xQUkdSb1dsUmplbGxVWnpKT2JWVTFXV3BCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwUlZlRTVxU1hrS1RsUkdiVTU2U1RSWlZGcHJXVlJHYkUxcVVUUk9hbFV3V1cxSk5FMXFaek5aVjFVelRUSkZORTVxV214UFYwbDNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWxFd1RucEpOVTVxWXpGT2FrRjJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUpjMjB5TDJ0QlFVRlJSQXBCUldOM1VsRkpaMHBFVlV0U01TOUxVRFZDWWtZcmJTdHdUbkl5WVdGT1oyZGFkMDR4SzJsT05GQnVLMkpJV2pneFlUQkRTVkZFVkc5c04zTkhSMWt2Q2xsc1JIUjBURU4xV1ZreFFtcEZiSG8xWm1OSVFYSlhOSFpIZGs5YVZHMDVTbFJCUzBKblozRm9hMnBQVUZGUlJFRjNUbTVCUkVKclFXcEJNbFJvY2pBS1dIaG9OVmRyVUVkT2JFdEhaMjRyY1VrNE1sRjRWbVowZWtaV1pHMHhNM2R3VW1sMWVVZHBOVEYyYUhodlZrbFViR2hxZVRsM1dtZExLMGxEVFVkeWRBcDFTM1ZHUjJsNVJYcHFNSGRUTDBGck4zRjFjRVpTYzAxa1ltZ3djWEJXWkVoNldWaHJaWFZSVFVWNmFGRklWall3VVd4cmRYWlBUWEI1YURaVVFUMDlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyjADAgEAMIICwQYJKoZIhvcNAQcCoIICsjCCAq4CAQMxDTALBglghkgBZQMEAgEwgbgGCyqGSIb3DQEJEAEEoIGoBIGlMIGiAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQg0mIPHE4XV+9v6qQUiZ3WsI4EjZ2rd/reFEdLEVJIuwwCFQDtblrtRHqxeW4FW6I2/Fd5g4VrARgPMjAyNjA5MTAxMTQ5MzRaMAMCAQGgMqQwMC4xFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEVMBMGA1UEAxMMc2lnc3RvcmUtdHNhoAAxggHbMIIB1wIBATBRMDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAsGCWCGSAFlAwQCAaCB/DAaBgkqhkiG9w0BCQMxDQYLKoZIhvcNAQkQAQQwHAYJKoZIhvcNAQkFMQ8XDTI2MDkxMDExNDkzNFowLwYJKoZIhvcNAQkEMSIEIEAZvSCkXGiy9zCJeFkpctHEK2nIY8tHyRzsl4o85EEOMIGOBgsqhkiG9w0BCRACLzF/MH0wezB5BCCF+Se8B6tiysO0Q1bBDvyBssaIP9p6uebYcNnROs0FtzBVMD2kOzA5MRUwEwYDVQQKEwxzaWdzdG9yZS5kZXYxIDAeBgNVBAMTF3NpZ3N0b3JlLXRzYS1zZWxmc2lnbmVkAhQ6E1QvDJBh7rzBQy/Lio6LKiOLDDAKBggqhkjOPQQDAgRnMGUCMB3Ma6dCowk+B2d51t3UsrOZXNBHglVool+tFSFxyI+bE067XHyY8JkVYb40gWRZgQIxAJ9Q4NQu+zYBgYlWRORk0SyoaVz/YAjceFFNm8uVNknx+SeNLgVq/d0pjci7KKH5SQ=="}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"We68EO24y6ls8cRlxTi6CcnuVOUfNuB5zDE41qWXzo0="}, "signature":"MEUCIQD4cZIaN+Z+pJGtuIpIwS6fgFex9WoN/FLiwQ+TN1BFwwIgZwYl+wwraZCQ6IcK2sDYfZyHyEgB/p3MQh0xoMt3JIs="}}
|
|
|
|
| 1 |
+
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json","verificationMaterial":{"certificate":{"rawBytes":"MIIHSzCCBtCgAwIBAgIUARbiR8eKelzOl01rJcf7WTqALI0wCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwOTExMDkzMjU1WhcNMjYwOTExMDk0MjU1WjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEc3QyjxYTM01NW6MwtrWEfdPikX87mbJ7mujpTlq3eapfgE35IlUzwWFpiE95onxMDQkCCPoWYWn2nwme/JciWKOCBe8wggXrMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUqtVZYZOx0gjLEDqiZ4IXvAKfLAIwHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoNDg2YjM0YzhiYzcyNGYxOTA1ZmQzZmRjMDRiNTk5MTA2ZjFiN2NmZDATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoNDg2YjM0YzhiYzcyNGYxOTA1ZmQzZmRjMDRiNTk5MTA2ZjFiN2NmZDAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoNDg2YjM0YzhiYzcyNGYxOTA1ZmQzZmRjMDRiNTk5MTA2ZjFiN2NmZDAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKDQ4NmIzNGM4YmM3MjRmMTkwNWZkM2ZkYzA0YjU5OTEwNmYxYjdjZmQwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzQ1ODQ1MzcxNTEvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBiQYKKwYBBAHWeQIEAgR7BHkAdwB1AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoI/QHSkAAAQDAEYwRAIgOJrsq+e9Dv9anbKkAPRyTJ0yL5LoCYvgAAQJ7zCrNJsCIGAsUve9QjmqGfhVKjALqgA13IyCQkQIwUBxiw8hfpPEMAoGCCqGSM49BAMDA2kAMGYCMQDMsCfcnURFrEa/8IsfmNg6/H+QgQu6ob/mA6gBlhEmUpf/TlhIi8hqHiuGpfBvgVkCMQDAuvFWJgJxfo1ZuakJs/SZECk8tO4rB0ViDs3BqRb1T3UZlcGYWdOSZoXbJRZB1Xw="},"tlogEntries":[{"logIndex":"2792100816","logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="},"kindVersion":{"kind":"hashedrekord","version":"0.0.1"},"integratedTime":"1789119176","inclusionPromise":{"signedEntryTimestamp":"MEYCIQC8FOWDeGq+VQcuUmPiuT7xCY0U+QWT21D8Fn2SXOneGgIhAJyZsQnnynWrvYsVIevMTDMofL1fP4vrDsQxGWiFnxCG"},"inclusionProof":{"logIndex":"2670196554","rootHash":"PjrNq6jqtgwqukCerHsa2zZFJiyrpIkJ6aEzfpb9vQM=","treeSize":"2670196586","hashes":["vCnQisKzLaBS9QYrhoNqVkAQezkoMolsswM7XA/Bv2o=","+pawtnsgA/q6l1bHF7RfzuEinR/2yhbcm4SIiKHEF64=","KNu1vgEUzlloYfUFooFyaliGUxosI6UnRETB2vFUWOc=","1VRwHxQtoH3R1cSEqzOdf52ifAljt8g+b3ONdu1oPuk=","EW78gA6/dv/NnOoG975yzTJXnO35RNuCZdx1+lAV/Vs=","h/qseIZpiebPhddIc4F+1JGPxqZ/mrqGRWrsFfgSBuo=","qVaW3UK2fdsaCcn8Ev7hNOi9OW6o9LlsO9VqZusjbg4=","NEjH87ZiMRR+YIysVI82F+HU3io9kutWbQvN8b+e7Gc=","KsTbYWTnQrSEU/0L9BxW3EKi1K1OBZAWQlFwDifhXAE=","9MaTDeuzozFzDElxaj6vwWw2epKAH9ZeT1njno4wAt8=","fPi+si31P6UOUJTHazN7wzii6QSvGEztgJKgVMN0fKg=","q3NVVYG2lC8kuX1fXpyJP2hEuonNdxb7dSDcud/2ho8=","oAS8o4MBzmTN1LVnviaIdsOVrMY+Cg2MTqyze73BGnY=","15Vzt0nU/9+QQYX0Pzg5ZLUfglcPqm8aCy9EsAjangc=","IG+pl3sU6YoRSqWPHtiSQbRJPCFcf4JeGc2fHPdqSYI=","4dTtD+QBzvTOC8xDaENhkUTbPKJNt3aLFv5v9LBySks=","lqiz5Og+wRcVZ/oMWDixqal9HNmtLYvg6HLjQs/hRzY=","/mjQogAQefxrs6hDWFie0uOlNqZc4omr6I+XxshnZ1Y=","Ari+W6zWoRpeR1Lw5lnM64DXE7cw7awR/QTfo7IO3PM=","4SexKjDhe33uQFmbzX5SUsB0SyhTm14ObSvNGPpKDFk=","7tD2K8mAqGTMENrP2YXWz71BeTWQPCiYd3CESAgQLz8=","eXSLrLqKSxRTixSP/k4YimriB2lufao6trqTrsxxlyg=","SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=","xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="],"checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2670196586\nPjrNq6jqtgwqukCerHsa2zZFJiyrpIkJ6aEzfpb9vQM=\n\n— rekor.sigstore.dev wNI9ajBGAiEA3gsGwSguZlBOKZ2FgO+rCMv9TyoGUs/cGrOXKn6R9x8CIQCKNc8J9NocPTolWh7nhnD/bq8QnVOQWR0D7vzQhw+LpA==\n"}},"canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiJmNjQxMzZmMTMxN2IxZDBjZjJlYzU5Nzc5OTNhMTFlYmE2OTk1ZGQ3ZDU3ODNiMjQ5ODRlMWUyYTBkYTk2NDRiIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJUUNaamQ1bldBMW5mZUY4b3I1VWRpbUdaN1kyL3M5S00weFpnK2tMaUxlWEd3SWdDTFk3MlorUFdpeVpVOHNxU3Jnc0g3bTl0d3dCZjUwQ0V2UW1jR2Y3ZzA0PSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRla05EUW5SRFowRjNTVUpCWjBsVlFWSmlhVkk0WlV0bGJIcFBiREF4Y2twalpqZFhWSEZCVEVrd2QwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlVUlhoTlJHdDZUV3BWTVZkb1kwNU5hbGwzVDFSRmVFMUVhekJOYWxVeFYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZqTTFGNWFuaFpWRTB3TVU1WE5rMTNkSEpYUldaa1VHbHJXRGczYldKS04yMTFhbkFLVkd4eE0yVmhjR1puUlRNMVNXeFZlbmRYUm5CcFJUazFiMjU0VFVSUmEwTkRVRzlYV1ZkdU1tNTNiV1V2U21OcFYwdFBRMEpsT0hkbloxaHlUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZ4ZEZaYUNsbGFUM2d3WjJwTVJVUnhhVm8wU1ZoMlFVdG1URUZKZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOU9SR2N5V1dwTk1GbDZhR2xaZW1ONVRrZFplRTlVUVRGYWJWRjZDbHB0VW1wTlJGSnBUbFJyTlUxVVFUSmFha1pwVGpKT2JWcEVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMDVFWnpKWmFrMHdXWHBvYVZsNlkzbE9SMWw0VDFSQk1WcHRVWHBhYlZKcVRVUlNhVTVVYXpVS1RWUkJNbHBxUm1sT01rNXRXa1JCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMDVFWnpKWmFrMHdXWHBvYVZsNlkzbE9SMWw0VDFSQk1WcHRVWG9LV20xU2FrMUVVbWxPVkdzMVRWUkJNbHBxUm1sT01rNXRXa1JCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwUlJORTV0U1hvS1RrZE5ORmx0VFROTmFsSnRUVlJyZDA1WFdtdE5NbHByV1hwQk1GbHFWVFZQVkVWM1RtMVplRmxxWkdwYWJWRjNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWxFeFQwUlJNVTE2WTNoT1ZFVjJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwVVZsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTTjBKSWEwRUtaSGRDTVVGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUpMMUZJVTJ0QlFVRlJSQXBCUlZsM1VrRkpaMDlLY25OeEsyVTVSSFk1WVc1aVMydEJVRko1VkVvd2VVdzFURzlEV1hablFVRlJTamQ2UTNKT1NuTkRTVWRCYzFWMlpUbFJhbTF4Q2tkbWFGWkxha0ZNY1dkQk1UTkplVU5SYTFGSmQxVkNlR2wzT0dobWNGQkZUVUZ2UjBORGNVZFRUVFE1UWtGTlJFRXlhMEZOUjFsRFRWRkVUWE5EWm1NS2JsVlNSbkpGWVM4NFNYTm1iVTVuTmk5SUsxRm5VWFUyYjJJdmJVRTJaMEpzYUVWdFZYQm1MMVJzYUVscE9HaHhTR2wxUjNCbVFuWm5WbXREVFZGRVFRcDFka1pYU21kS2VHWnZNVnAxWVd0S2N5OVRXa1ZEYXpoMFR6UnlRakJXYVVSek0wSnhVbUl4VkROVldteGpSMWxYWkU5VFdtOVlZa3BTV2tJeFdIYzlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}],"timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyTADAgEAMIICwAYJKoZIhvcNAQcCoIICsTCCAq0CAQMxDTALBglghkgBZQMEAgEwgbcGCyqGSIb3DQEJEAEEoIGnBIGkMIGhAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgv7nQ6Y1Bw93wHqffBORVCJ3tFvwhvQTdUvhRv7fMDpkCFGTFypLH4RjYScTiaarigsHHYsB7GA8yMDI2MDkxMTA5MzI1NlowAwIBAaAypDAwLjEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MRUwEwYDVQQDEwxzaWdzdG9yZS10c2GgADGCAdswggHXAgEBMFEwOTEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MSAwHgYDVQQDExdzaWdzdG9yZS10c2Etc2VsZnNpZ25lZAIUOhNULwyQYe68wUMvy4qOiyojiwwwCwYJYIZIAWUDBAIBoIH8MBoGCSqGSIb3DQEJAzENBgsqhkiG9w0BCRABBDAcBgkqhkiG9w0BCQUxDxcNMjYwOTExMDkzMjU2WjAvBgkqhkiG9w0BCQQxIgQgZ+IwqL+SDLuP+9GK2VnADhdMh0rt0YkOpk2BL/2o/QQwgY4GCyqGSIb3DQEJEAIvMX8wfTB7MHkEIIX5J7wHq2LKw7RDVsEO/IGyxog/2nq55thw2dE6zQW3MFUwPaQ7MDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAoGCCqGSM49BAMCBGcwZQIwHPV2E7lEFqomqOU1zz6lg8PiaWdq7J3EJ44GgV8l1t6okdrOXiBC/Yt2zMXdj7cjAjEAy0Pn+x32HHaQkQdSVDhXdlbTdrZ5ZYk1eDyjWb+ABLJVjCMsoAvXq83kE78hCQay"}]}},"messageSignature":{"messageDigest":{"algorithm":"SHA2_256","digest":"9kE28TF7HQzy7Fl3mToR66aZXdfVeDskmE4eKg2pZEs="},"signature":"MEUCIQCZjd5nWA1nfeF8or5UdimGZ7Y2/s9KM0xZg+kLiLeXGwIgCLY72Z+PWiyZU8sqSrgsH7m9twwBf50CEvQmcGf7g04="}}
|
build/torch-cuda/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-cuda/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-cuda/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-cuda/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-cuda/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-cuda/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-cuda/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
|