Kernels:
Trusted publisher
Uploaded using `kernel-builder`.
Browse filesThis view is limited to 50 files because it contains too many changes. See raw diff
- build/torch212-cxx11-cu126-x86_64-linux/{_megablocks_cuda_c4d0dc5.abi3.so → _megablocks_cuda_addf474.abi3.so} +2 -2
- build/torch212-cxx11-cu126-x86_64-linux/_ops.py +3 -3
- build/torch212-cxx11-cu126-x86_64-linux/megablocks/__init__.py +0 -26
- build/torch212-cxx11-cu126-x86_64-linux/metadata.json +7 -7
- build/torch212-cxx11-cu126-x86_64-linux/metadata.json.sigstore +1 -1
- build/{torch212-cxx11-cu132-x86_64-linux/_megablocks_cuda_c4d0dc5.abi3.so → torch212-cxx11-cu130-x86_64-linux/_megablocks_cuda_addf474.abi3.so} +2 -2
- build/torch212-cxx11-cu130-x86_64-linux/_ops.py +3 -3
- build/torch212-cxx11-cu130-x86_64-linux/megablocks/__init__.py +0 -26
- build/torch212-cxx11-cu130-x86_64-linux/metadata.json +9 -6
- build/torch212-cxx11-cu130-x86_64-linux/metadata.json.sigstore +1 -1
- build/{torch212-cxx11-cu130-x86_64-linux/_megablocks_cuda_c4d0dc5.abi3.so → torch212-cxx11-cu132-x86_64-linux/_megablocks_cuda_addf474.abi3.so} +2 -2
- build/torch212-cxx11-cu132-x86_64-linux/_ops.py +3 -3
- build/torch212-cxx11-cu132-x86_64-linux/megablocks/__init__.py +0 -26
- build/torch212-cxx11-cu132-x86_64-linux/metadata.json +9 -6
- build/torch212-cxx11-cu132-x86_64-linux/metadata.json.sigstore +1 -1
- build/torch213-cxx11-cu126-x86_64-linux/{_megablocks_cuda_c4d0dc5.abi3.so → _megablocks_cuda_addf474.abi3.so} +2 -2
- build/torch213-cxx11-cu126-x86_64-linux/_ops.py +3 -3
- build/torch213-cxx11-cu126-x86_64-linux/megablocks/__init__.py +0 -26
- build/torch213-cxx11-cu126-x86_64-linux/metadata.json +7 -7
- build/torch213-cxx11-cu126-x86_64-linux/metadata.json.sigstore +1 -1
- build/torch213-cxx11-cu130-x86_64-linux/_megablocks_cuda_addf474.abi3.so +3 -0
- build/torch213-cxx11-cu130-x86_64-linux/_megablocks_cuda_c4d0dc5.abi3.so +0 -3
- build/torch213-cxx11-cu130-x86_64-linux/_ops.py +3 -3
- build/torch213-cxx11-cu130-x86_64-linux/megablocks/__init__.py +0 -26
- build/torch213-cxx11-cu130-x86_64-linux/metadata.json +9 -6
- build/torch213-cxx11-cu130-x86_64-linux/metadata.json.sigstore +1 -1
- build/torch213-cxx11-cu132-x86_64-linux/_megablocks_cuda_addf474.abi3.so +3 -0
- build/torch213-cxx11-cu132-x86_64-linux/_megablocks_cuda_c4d0dc5.abi3.so +0 -3
- build/torch213-cxx11-cu132-x86_64-linux/_ops.py +3 -3
- build/torch213-cxx11-cu132-x86_64-linux/megablocks/__init__.py +0 -26
- build/torch213-cxx11-cu132-x86_64-linux/metadata.json +9 -6
- build/torch213-cxx11-cu132-x86_64-linux/metadata.json.sigstore +1 -1
- build/torch214-cxx11-cu126-x86_64-linux/__init__.py +205 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/__init__.py +0 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/_triton_kernels/__init__.py +0 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/_triton_kernels/gmm.py +574 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/adapter.py +53 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/configs.py +5 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/gmm.py +567 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/__init__.py +0 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/_triton/__init__.py +0 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/_triton/arch_info.py +46 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/_triton/pid_preprocessing.py +100 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/gmm_common.py +752 -0
- build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/logger.py +47 -0
- build/torch214-cxx11-cu126-x86_64-linux/_layers/__init__.py +10 -0
- build/torch214-cxx11-cu126-x86_64-linux/_layers/activation_fn.py +33 -0
- build/torch214-cxx11-cu126-x86_64-linux/_layers/all_to_all.py +54 -0
- build/torch214-cxx11-cu126-x86_64-linux/_layers/arguments.py +101 -0
- build/torch214-cxx11-cu126-x86_64-linux/_layers/common.py +26 -0
build/torch212-cxx11-cu126-x86_64-linux/{_megablocks_cuda_c4d0dc5.abi3.so → _megablocks_cuda_addf474.abi3.so}
RENAMED
|
@@ -1,3 +1,3 @@
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
-
oid sha256:
|
| 3 |
-
size
|
|
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:75eac9fccf5aa5c74e862b0aa580f1c420cfc675df9b8b6048ca2b2c3b71ad66
|
| 3 |
+
size 13818088
|
build/torch212-cxx11-cu126-x86_64-linux/_ops.py
CHANGED
|
@@ -1,9 +1,9 @@
|
|
| 1 |
import torch
|
| 2 |
-
from . import
|
| 3 |
-
ops = torch.ops.
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
-
return f"
|
|
|
|
| 1 |
import torch
|
| 2 |
+
from . import _megablocks_cuda_addf474
|
| 3 |
+
ops = torch.ops._megablocks_cuda_addf474
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
+
return f"_megablocks_cuda_addf474::{op_name}"
|
build/torch212-cxx11-cu126-x86_64-linux/megablocks/__init__.py
DELETED
|
@@ -1,26 +0,0 @@
|
|
| 1 |
-
import ctypes
|
| 2 |
-
import importlib.util
|
| 3 |
-
import sys
|
| 4 |
-
from pathlib import Path
|
| 5 |
-
from types import ModuleType
|
| 6 |
-
|
| 7 |
-
|
| 8 |
-
def _import_from_path(file_path: Path) -> ModuleType:
|
| 9 |
-
# We cannot use the module name as-is, after adding it to `sys.modules`,
|
| 10 |
-
# it would also be used for other imports. So, we make a module name that
|
| 11 |
-
# depends on the path for it to be unique using the hex-encoded hash of
|
| 12 |
-
# the path.
|
| 13 |
-
path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
|
| 14 |
-
module_name = path_hash
|
| 15 |
-
spec = importlib.util.spec_from_file_location(module_name, file_path)
|
| 16 |
-
if spec is None:
|
| 17 |
-
raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
|
| 18 |
-
module = importlib.util.module_from_spec(spec)
|
| 19 |
-
if module is None:
|
| 20 |
-
raise ImportError(f"Cannot load module {module_name} from spec")
|
| 21 |
-
sys.modules[module_name] = module
|
| 22 |
-
spec.loader.exec_module(module) # type: ignore
|
| 23 |
-
return module
|
| 24 |
-
|
| 25 |
-
|
| 26 |
-
globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
build/torch212-cxx11-cu126-x86_64-linux/metadata.json
CHANGED
|
@@ -1,9 +1,10 @@
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
|
|
|
| 7 |
"backend": {
|
| 8 |
"type": "cuda",
|
| 9 |
"archs": [
|
|
@@ -14,7 +15,7 @@
|
|
| 14 |
"8.6",
|
| 15 |
"8.7",
|
| 16 |
"8.9",
|
| 17 |
-
"9.0"
|
| 18 |
]
|
| 19 |
},
|
| 20 |
"digest": {
|
|
@@ -48,8 +49,8 @@
|
|
| 48 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 49 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 50 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 51 |
-
"
|
| 52 |
-
"_ops.py": "
|
| 53 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 54 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 55 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
@@ -61,7 +62,6 @@
|
|
| 61 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 62 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 63 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
| 64 |
-
"megablocks/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY=",
|
| 65 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 66 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 67 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
@@ -106,11 +106,11 @@
|
|
| 106 |
"provenance": {
|
| 107 |
"kernel-builder": {
|
| 108 |
"version": "0.17.0-dev0",
|
| 109 |
-
"
|
| 110 |
"dirty": false
|
| 111 |
},
|
| 112 |
"kernel": {
|
| 113 |
-
"
|
| 114 |
"dirty": false
|
| 115 |
}
|
| 116 |
}
|
|
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
+
"id": "_megablocks_cuda_addf474",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
| 7 |
+
"kernel-depends": [],
|
| 8 |
"backend": {
|
| 9 |
"type": "cuda",
|
| 10 |
"archs": [
|
|
|
|
| 15 |
"8.6",
|
| 16 |
"8.7",
|
| 17 |
"8.9",
|
| 18 |
+
"9.0+PTX"
|
| 19 |
]
|
| 20 |
},
|
| 21 |
"digest": {
|
|
|
|
| 49 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 50 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 51 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 52 |
+
"_megablocks_cuda_addf474.abi3.so": "derJ/M9apcdOhisKpYDxxCDPxnXfm4tgSMorLDtxrWY=",
|
| 53 |
+
"_ops.py": "WXlV4XwtKQjdQ6nap3k9TalaHkK0ypeSTshOrV6Hczo=",
|
| 54 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 55 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 56 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
|
|
| 62 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 63 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 64 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
|
|
|
| 65 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 66 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 67 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
|
|
| 106 |
"provenance": {
|
| 107 |
"kernel-builder": {
|
| 108 |
"version": "0.17.0-dev0",
|
| 109 |
+
"commit": "7d2828e6592de641b44fb8eb896719101a2ab2fb",
|
| 110 |
"dirty": false
|
| 111 |
},
|
| 112 |
"kernel": {
|
| 113 |
+
"commit": "addf4741fafb55b4c101b15985dfaeea7c7c2f3a",
|
| 114 |
"dirty": false
|
| 115 |
}
|
| 116 |
}
|
build/torch212-cxx11-cu126-x86_64-linux/metadata.json.sigstore
CHANGED
|
@@ -1 +1 @@
|
|
| 1 |
-
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSzCCBtGgAwIBAgIUTQP9wTaeeOtp+g3scdQtq9C411owCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI1MTIyMTE3WhcNMjYwODI1MTIzMTE3WjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEPJHV4MhjvJ2wQ8LwG8ho0UIzDoQisDTwmTYSb6o9yaR5vo2Lc5MWf+ps5jmMDZ0koM3SD/zOhuZCPrAI0xoTSaOCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUXmifi2X5OX1prd4IqcLgre4p/6kwHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGM0ZDBkYzU2NjYxYmFkYjljYzQ3YmRlNjQ3ZTA4ZmUzYjRiZjc0NTgwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzI4NDQ2ODk4OTQvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoDjeIycAAAQDAEcwRQIhAPJGRAotH/fo+g0co0NcZOT6HptOzHfwEx9UqJgnqFr/AiB+rqzjXc5Pjiw5Ks+O4AwTKhC2phb1tJOzdWDO3uQfeDAKBggqhkjOPQQDAwNoADBlAjBEB9Up0jyK6OhoRgTTceucPlZmh+Lb6/GzvJRCN/wotx+NzrsIyp0Z9yF28dKwj8kCMQCvpS9tqg9NLWA6Gwuli8N2tq3DmXC8UmXTuX76YBtVXsL/PtEDz8aOOcmxa+wsZhs="}, "tlogEntries":[{"logIndex":"2583302587", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787660477", "inclusionPromise":{"signedEntryTimestamp":"MEQCIFwqhArAhbAHQvG9biEpGXGnX7Kzy/ag5qMjx3PYHsYpAiAkA7lfjxk0WfIAJ95/WZRQlx6LyhnnP1d9EM/6Na9ONg=="}, "inclusionProof":{"logIndex":"2461398325", "rootHash":"ufLjEJeISMOHLeqKARIA/a15/J4rOiSkz3DMv6ZT5jY=", "treeSize":"2461398326", "hashes":["VRANlHr2KmDIZnL7WzDzhmVmb5oPtDMLaG/OimiIQm4=", "HvIDjbtTaqD74X0cy8sx0UPhrGsYXIpOSqe3GREvdUs=", "tWauNUno2kG7EKvel+T+4ks6XYUxIgWBCrom3zwUgLg=", "BQO6Mvbc6dfqTtcGGvXrh9+QB796m4uHnHENTnE52Ao=", "eUK42vZbtTAA+teDm0LcO+qURyI8IPFvT/lz1znHKQ8=", "r8ZmsIkPd8oW6c3MXr06PoZ4nJYw0Vt5ki3ZDt5yOEM=", "M0t6Pyxtbr61GCmNg60qyyc/qJAEEt98xDnkIgI5BIY=", "EY8T0lxQDrmYoWBGpKNGttaHfKmpfqKcWqSNF4XOQKE=", "T4Jz+m9vO48rOpMSsdZSW28/bgHR8prz7JeNX1s7WiQ=", "EXKrSbeG2qm5APD5QSDN66X7UHz9OSofiW9PyNU7fPM=", "cjIeeutzThxIPTlFFHJqa7bAzn3k2FzK8Nrdd56OWvE=", "+pNXWGIXzU1dqSy/aePHKYtSkDnFinI+/bL5HnurJ6w=", "G1R9F5B2KizFA2NyhDznjNKXPOF1SDmbGHz+omLG7TY=", "RZ7XwGuqMr5gDUh2HpUcaXp+AmIiEcPJiMui4GLJE7o=", "mBm+vQtn0C4thMnlTnxfxqowq1dXsPBCKaJ85da3JeU=", "b80/J/68RC8/tx5xRdlKk6MmTpZiVRotquWBrE/z1pI=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2461398326\nufLjEJeISMOHLeqKARIA/a15/J4rOiSkz3DMv6ZT5jY=\n\n— rekor.sigstore.dev wNI9ajBGAiEAy2+bc6au6mGJM2X55oIkc/qkKalnawOHoulZdgZlEL8CIQCqpCQynrlcGYRMNOA5vrrpBR2IYv/hqoPxykwSYH/9uw==\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiIyOGRkODcxNzkyODQwNjVhZDgxNDQzZjZlZTFhM2MzNzkyYjI3NWJmNDljYmRlMTkxNGE4YmVkOWNhMzk1ZDM1In19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJQkV4M3hNWUdYeXB2c0pYZHhYVEoySUFXQnhuUTc2ZUlHbURqcXZ1cUd4TkFpRUFpVXNMZVNEN2JJS2xIZGlWTEhMZE9mUTZ0b2tjam11aUVjTzg3NHdud1hFPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRla05EUW5SSFowRjNTVUpCWjBsVlZGRlFPWGRVWVdWbFQzUndLMmN6YzJOa1VYUnhPVU0wTVRGdmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1RGTlZFbDVUVlJGTTFkb1kwNU5hbGwzVDBSSk1VMVVTWHBOVkVVelYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZRU2toV05FMW9hblpLTW5kUk9FeDNSemhvYnpCVlNYcEViMUZwYzBSVWQyMVVXVk1LWWpadk9YbGhValYyYnpKTVl6Vk5WMllyY0hNMWFtMU5SRm93YTI5Tk0xTkVMM3BQYUhWYVExQnlRVWt3ZUc5VVUyRlBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZZYldsbUNta3lXRFZQV0RGd2NtUTBTWEZqVEdkeVpUUndMelpyZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVplbEpyVFVkU2FrNVVXVEpPYWtacFdWZFNhVTlYVG1wT1JHUnBDbHBIVlRKT1JHUnNUVVJvYlZwVVRtbE9SMHB0VG5wUk1VOUVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2xhUjFVeVRrUmtiRTFFYUcwS1dsUk9hVTVIU20xT2VsRXhUMFJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2tLV2tkVk1rNUVaR3hOUkdodFdsUk9hVTVIU20xT2VsRXhUMFJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZE5NRnBFUW1zS1dYcFZNazVxV1hoWmJVWnJXV3BzYWxsNlVUTlpiVkpzVG1wUk0xcFVRVFJhYlZWNldXcFNhVnBxWXpCT1ZHZDNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWtrMFRrUlJNazlFYXpSUFZGRjJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUVhbVZKZVdOQlFVRlJSQXBCUldOM1VsRkphRUZRU2tkU1FXOTBTQzltYnl0bk1HTnZNRTVqV2s5VU5raHdkRTk2U0daM1JYZzVWWEZLWjI1eFJuSXZRV2xDSzNKeGVtcFlZelZRQ21wcGR6VkxjeXRQTkVGM1ZFdG9Rekp3YUdJeGRFcFBlbVJYUkU4emRWRm1aVVJCUzBKblozRm9hMnBQVUZGUlJFRjNUbTlCUkVKc1FXcENSVUk1VlhBS01HcDVTelpQYUc5U1oxUlVZMlYxWTFCc1dtMW9LMHhpTmk5SGVuWktVa05PTDNkdmRIZ3JUbnB5YzBsNWNEQmFPWGxHTWpoa1MzZHFPR3REVFZGRGRncHdVemwwY1djNVRreFhRVFpIZDNWc2FUaE9NblJ4TTBSdFdFTTRWVzFZVkhWWU56WlpRblJXV0hOTUwxQjBSVVI2T0dGUFQyTnRlR0VyZDNOYWFITTlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyzADAgEAMIICwgYJKoZIhvcNAQcCoIICszCCAq8CAQMxDTALBglghkgBZQMEAgEwgbgGCyqGSIb3DQEJEAEEoIGoBIGlMIGiAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQg3Zr4gvYHXguQ9HpTQa/zCt0frc5KzBdqeyaaP+z4ig4CFQDt4abH42p4ttb62ZXfRiYymbqw6hgPMjAyNjA4MjUxMjIxMTdaMAMCAQGgMqQwMC4xFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEVMBMGA1UEAxMMc2lnc3RvcmUtdHNhoAAxggHcMIIB2AIBATBRMDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAsGCWCGSAFlAwQCAaCB/DAaBgkqhkiG9w0BCQMxDQYLKoZIhvcNAQkQAQQwHAYJKoZIhvcNAQkFMQ8XDTI2MDgyNTEyMjExN1owLwYJKoZIhvcNAQkEMSIEIHn5fQODI/BgWEUrqX3HE5ot0dflrMJBwJfZhlbUV5g1MIGOBgsqhkiG9w0BCRACLzF/MH0wezB5BCCF+Se8B6tiysO0Q1bBDvyBssaIP9p6uebYcNnROs0FtzBVMD2kOzA5MRUwEwYDVQQKEwxzaWdzdG9yZS5kZXYxIDAeBgNVBAMTF3NpZ3N0b3JlLXRzYS1zZWxmc2lnbmVkAhQ6E1QvDJBh7rzBQy/Lio6LKiOLDDAKBggqhkjOPQQDAgRoMGYCMQC9GR6NfulwVg45HhNSg7CnmiCwck7DNLs5vI4vKMtlovk1S+D3qj5FYzWmKh721lECMQD7PEI3N5vc+hrgy86iCZtYEIjZq7k3i0q1/fe3xThkJxFdhFATM9yQJuIdxzsSniA="}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"KN2HF5KEBlrYFEP27ho8N5Kydb9Jy94ZFKi+2co5XTU="}, "signature":"MEUCIBEx3xMYGXypvsJXdxXTJ2IAWBxnQ76eIGmDjqvuqGxNAiEAiUsLeSD7bIKlHdiVLHLdOfQ6tokcjmuiEcO874wnwXE="}}
|
|
|
|
| 1 |
+
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSjCCBtCgAwIBAgIUcnefoa2hroWalI1KhhDd0iQTl64wCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI3MDcxNDEyWhcNMjYwODI3MDcyNDEyWjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEqJgLT/s77Ua9ZAQHF1t9p/h16lG+18NSQCJ8w7QoaONmULYsXQB8H0lNJOSvpLSq4fKJ6RHIMhp91fzAsA+9q6OCBe8wggXrMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUj1e0YQkuCWel+LsKt+YQp0VZbH4wHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGFkZGY0NzQxZmFmYjU1YjRjMTAxYjE1OTg1ZGZhZWVhN2M3YzJmM2EwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzMwNDUzOTU5NTIvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBiQYKKwYBBAHWeQIEAgR7BHkAdwB1AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoEIRtl8AAAQDAEYwRAIgVgqnPHDnwBnUWG9dr1HfzznpQm7T3O4busbrtbCbcYwCIAjg/mbu6vdd48B3A84zr+2LSc5nnWjmmhGpWGN+MkngMAoGCCqGSM49BAMDA2gAMGUCMFuIJE/JYWzlWth5JuVzDrXTzxniaF+wGxShuHzrTnbF9f7n2HwSNfREY3hnL+GMyQIxAPVCL031YIJ7KtOpjTbaGNG2D6yEZwrwao1NVT9xamWbED9p8sQWk9clGPddIdlihQ=="}, "tlogEntries":[{"logIndex":"2614048455", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787814852", "inclusionPromise":{"signedEntryTimestamp":"MEUCIEnatko5JlMhXJ2NzhOlCrNFOrOvUyfGVkJQ25TCeW6pAiEA+3S3PDF7VtCXDElQZmyzLIdnr17zTR/X5vU6Gro6rnQ="}, "inclusionProof":{"logIndex":"2492144193", "rootHash":"OGbwNrJjc+flOi5K+TnVTDy2pX0o9F8o71DTOhk7M/U=", "treeSize":"2492144221", "hashes":["vhQ4CS2xVVo3D8k32qQJfoKnJ18JfxjlbF4ppD7YiSw=", "LEU31UJM50zZ48YlhexjezNCSAxKcShzRxHsW+6AJlc=", "DDFRUYEZguBgANUZaHKRoXIrcKa3VuV+RTZOFvj3bLk=", "z8MO8JCBH3AJGXDd2kcEeFCSNJAvv8n7F4lZ/yt/Dts=", "BoV8hK3bGHabcdDVPwos21LnIWQXUnA17Axo/WSweSA=", "Rtw5GMrUOIGyS5ueZk0ib5Fs8F3TNSxO4icabC7JhpE=", "gxi1eJ8qX1K1iauxvafmRFx1FhOcSvz3IdnMqR1qs9E=", "VFlTxnDU4W+k/WA3ARTvzCiqVt1yxQIMUyUzMCugJ2Q=", "CFNMzmJ8N+uNx37P2asMqvrW9tw4axAVHnXLRO39tFs=", "NypqAbB+mNix4l1ahGeP9fRWiQzZCm4P1664makjjq4=", "8SoaeA7IeJPoIiveQbBfHCWXSDQabIk2PW+FT49NSEQ=", "VGrFt5JvlIUQvoGEdBBxrCe3lYbu581OhgLq/Xdyat4=", "7TfDCaMWgEwSJEZd/e33b24919rF4IDJLmbqNNRNhZs=", "0HfdxDt/zGugZLIuHasrdLEW9s8OaTykSSqBEDFQXdw=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2492144221\nOGbwNrJjc+flOi5K+TnVTDy2pX0o9F8o71DTOhk7M/U=\n\n— rekor.sigstore.dev wNI9ajBFAiEA4pz/k11O+33XODNFNbBNY0NiRCyXiZJeVMPxwMzqlbkCIDX+thsIWjPlhacQBgZRyQGLTk0ILGRkkkvBH1YsJauI\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiI4ZTUxYzRhMjA3NjdkYmIyZmU3OTc1ZDcwNmM1MGM1MjkwM2M0NThkYTBlMzY1ZWE1M2RhYjY3NDlhZjZhNzE0In19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJRlBheUNPdGNXVHIvWkd5Q2lGNXN0SUwwOTB2MnlJRmRjM0FLWk1aWVhxZEFpRUFpZjdZV2hKWTB6a1YyRVdIZ2tkYXp5RS9OK2NWQUo2RFFXVXcwNzlibmRFPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRha05EUW5SRFowRjNTVUpCWjBsVlkyNWxabTloTW1oeWIxZGhiRWt4UzJob1JHUXdhVkZVYkRZMGQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1ROTlJHTjRUa1JGZVZkb1kwNU5hbGwzVDBSSk0wMUVZM2xPUkVWNVYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZ4U21kTVZDOXpOemRWWVRsYVFWRklSakYwT1hBdmFERTJiRWNyTVRoT1UxRkRTamdLZHpkUmIyRlBUbTFWVEZseldGRkNPRWd3YkU1S1QxTjJjRXhUY1RSbVMwbzJVa2hKVFdod09URm1la0Z6UVNzNWNUWlBRMEpsT0hkbloxaHlUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZxTVdVd0NsbFJhM1ZEVjJWc0sweHpTM1FyV1ZGd01GWmFZa2cwZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVpWMUpyV21wUk0wNUVSbTFaVjFwcFRsUldhVTVIVFhoTlJFWnBDazFVVlRWUFJGWnJXbTFHYkZwWFJUTlplbVJxVFcxWmVsbFVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWxOVkZVMVQwUldhMXB0Um13S1dsZEZNMWw2WkdwTmJWbDZXVlJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWtLVFZSVk5VOUVWbXRhYlVac1dsZEZNMWw2WkdwTmJWbDZXVlJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZEdhMXBIV1RBS1RucFJlRnB0Um0xWmFsVXhXV3BTYWsxVVFYaFpha1V4VDFSbk1WcEhXbWhhVjFab1RqSk5NMWw2U20xTk1rVjNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWsxM1RrUlZlazlVVlRWT1ZFbDJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwVVZsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTTjBKSWEwRUtaSGRDTVVGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUZTVkowYkRoQlFVRlJSQXBCUlZsM1VrRkpaMVpuY1c1UVNFUnVkMEp1VlZkSE9XUnlNVWhtZW5wdWNGRnROMVF6VHpSaWRYTmljblJpUTJKaldYZERTVUZxWnk5dFluVTJkbVJrQ2pRNFFqTkJPRFI2Y2lzeVRGTmpOVzV1VjJwdGJXaEhjRmRIVGl0TmEyNW5UVUZ2UjBORGNVZFRUVFE1UWtGTlJFRXlaMEZOUjFWRFRVWjFTVXBGTDBvS1dWZDZiRmQwYURWS2RWWjZSSEpZVkhwNGJtbGhSaXQzUjNoVGFIVkllbkpVYm1KR09XWTNiakpJZDFOT1psSkZXVE5vYmt3clIwMTVVVWw0UVZCV1F3cE1NRE14V1VsS04wdDBUM0JxVkdKaFIwNUhNa1EyZVVWYWQzSjNZVzh4VGxaVU9YaGhiVmRpUlVRNWNEaHpVVmRyT1dOc1IxQmtaRWxrYkdsb1VUMDlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyjADAgEAMIICwQYJKoZIhvcNAQcCoIICsjCCAq4CAQMxDTALBglghkgBZQMEAgEwgbcGCyqGSIb3DQEJEAEEoIGnBIGkMIGhAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgdy1sI3m2IMSHfw1deoZrxwNhv0qVe+YKzpONjj5oBG4CFDBNEb9YDHm1FD33fH6LxCRTM0E6GA8yMDI2MDgyNzA3MTQxMlowAwIBAaAypDAwLjEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MRUwEwYDVQQDEwxzaWdzdG9yZS10c2GgADGCAdwwggHYAgEBMFEwOTEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MSAwHgYDVQQDExdzaWdzdG9yZS10c2Etc2VsZnNpZ25lZAIUOhNULwyQYe68wUMvy4qOiyojiwwwCwYJYIZIAWUDBAIBoIH8MBoGCSqGSIb3DQEJAzENBgsqhkiG9w0BCRABBDAcBgkqhkiG9w0BCQUxDxcNMjYwODI3MDcxNDEyWjAvBgkqhkiG9w0BCQQxIgQgMdZ3f9TM/8ePdAqeloxYU+LMs9Tc4xrwzzu3LvZc1JowgY4GCyqGSIb3DQEJEAIvMX8wfTB7MHkEIIX5J7wHq2LKw7RDVsEO/IGyxog/2nq55thw2dE6zQW3MFUwPaQ7MDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAoGCCqGSM49BAMCBGgwZgIxAIrOlsplb3whBLvuU5JbUVqpvlMe8XeiyxKSmjUgCCWuJk4hDNH6+xi6SMnzMG8l1wIxAOmXiZxrkf8fktq6YAGBuFllN0BuKELXKzFTy6MMgeuJ9tCvOjhttMazMPvyw7HB+g=="}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"jlHEogdn27L+eXXXBsUMUpA8RY2g42XqU9q2dJr2pxQ="}, "signature":"MEUCIFPayCOtcWTr/ZGyCiF5stIL090v2yIFdc3AKZMZYXqdAiEAif7YWhJY0zkV2EWHgkdazyE/N+cVAJ6DQWUw079bndE="}}
|
build/{torch212-cxx11-cu132-x86_64-linux/_megablocks_cuda_c4d0dc5.abi3.so → torch212-cxx11-cu130-x86_64-linux/_megablocks_cuda_addf474.abi3.so}
RENAMED
|
@@ -1,3 +1,3 @@
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
-
oid sha256:
|
| 3 |
-
size
|
|
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:b389c414114473ce7b5e8fe84eb3897115c7e63edb6c2536c212356ee37486bd
|
| 3 |
+
size 14586872
|
build/torch212-cxx11-cu130-x86_64-linux/_ops.py
CHANGED
|
@@ -1,9 +1,9 @@
|
|
| 1 |
import torch
|
| 2 |
-
from . import
|
| 3 |
-
ops = torch.ops.
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
-
return f"
|
|
|
|
| 1 |
import torch
|
| 2 |
+
from . import _megablocks_cuda_addf474
|
| 3 |
+
ops = torch.ops._megablocks_cuda_addf474
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
+
return f"_megablocks_cuda_addf474::{op_name}"
|
build/torch212-cxx11-cu130-x86_64-linux/megablocks/__init__.py
DELETED
|
@@ -1,26 +0,0 @@
|
|
| 1 |
-
import ctypes
|
| 2 |
-
import importlib.util
|
| 3 |
-
import sys
|
| 4 |
-
from pathlib import Path
|
| 5 |
-
from types import ModuleType
|
| 6 |
-
|
| 7 |
-
|
| 8 |
-
def _import_from_path(file_path: Path) -> ModuleType:
|
| 9 |
-
# We cannot use the module name as-is, after adding it to `sys.modules`,
|
| 10 |
-
# it would also be used for other imports. So, we make a module name that
|
| 11 |
-
# depends on the path for it to be unique using the hex-encoded hash of
|
| 12 |
-
# the path.
|
| 13 |
-
path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
|
| 14 |
-
module_name = path_hash
|
| 15 |
-
spec = importlib.util.spec_from_file_location(module_name, file_path)
|
| 16 |
-
if spec is None:
|
| 17 |
-
raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
|
| 18 |
-
module = importlib.util.module_from_spec(spec)
|
| 19 |
-
if module is None:
|
| 20 |
-
raise ImportError(f"Cannot load module {module_name} from spec")
|
| 21 |
-
sys.modules[module_name] = module
|
| 22 |
-
spec.loader.exec_module(module) # type: ignore
|
| 23 |
-
return module
|
| 24 |
-
|
| 25 |
-
|
| 26 |
-
globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
build/torch212-cxx11-cu130-x86_64-linux/metadata.json
CHANGED
|
@@ -1,14 +1,18 @@
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
|
|
|
| 7 |
"backend": {
|
| 8 |
"type": "cuda",
|
| 9 |
"archs": [
|
| 10 |
"10.0",
|
|
|
|
|
|
|
| 11 |
"12.0",
|
|
|
|
| 12 |
"7.5",
|
| 13 |
"8.0",
|
| 14 |
"8.6",
|
|
@@ -48,8 +52,8 @@
|
|
| 48 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 49 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 50 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 51 |
-
"
|
| 52 |
-
"_ops.py": "
|
| 53 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 54 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 55 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
@@ -61,7 +65,6 @@
|
|
| 61 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 62 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 63 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
| 64 |
-
"megablocks/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY=",
|
| 65 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 66 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 67 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
@@ -106,11 +109,11 @@
|
|
| 106 |
"provenance": {
|
| 107 |
"kernel-builder": {
|
| 108 |
"version": "0.17.0-dev0",
|
| 109 |
-
"
|
| 110 |
"dirty": false
|
| 111 |
},
|
| 112 |
"kernel": {
|
| 113 |
-
"
|
| 114 |
"dirty": false
|
| 115 |
}
|
| 116 |
}
|
|
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
+
"id": "_megablocks_cuda_addf474",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
| 7 |
+
"kernel-depends": [],
|
| 8 |
"backend": {
|
| 9 |
"type": "cuda",
|
| 10 |
"archs": [
|
| 11 |
"10.0",
|
| 12 |
+
"10.3",
|
| 13 |
+
"11.0",
|
| 14 |
"12.0",
|
| 15 |
+
"12.1+PTX",
|
| 16 |
"7.5",
|
| 17 |
"8.0",
|
| 18 |
"8.6",
|
|
|
|
| 52 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 53 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 54 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 55 |
+
"_megablocks_cuda_addf474.abi3.so": "s4nEFBFEc857Xo/oTrOJcRXH5j7bbCU2whI1buN0hr0=",
|
| 56 |
+
"_ops.py": "WXlV4XwtKQjdQ6nap3k9TalaHkK0ypeSTshOrV6Hczo=",
|
| 57 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 58 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 59 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
|
|
| 65 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 66 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 67 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
|
|
|
| 68 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 69 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 70 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
|
|
| 109 |
"provenance": {
|
| 110 |
"kernel-builder": {
|
| 111 |
"version": "0.17.0-dev0",
|
| 112 |
+
"commit": "7d2828e6592de641b44fb8eb896719101a2ab2fb",
|
| 113 |
"dirty": false
|
| 114 |
},
|
| 115 |
"kernel": {
|
| 116 |
+
"commit": "addf4741fafb55b4c101b15985dfaeea7c7c2f3a",
|
| 117 |
"dirty": false
|
| 118 |
}
|
| 119 |
}
|
build/torch212-cxx11-cu130-x86_64-linux/metadata.json.sigstore
CHANGED
|
@@ -1 +1 @@
|
|
| 1 |
-
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSzCCBtCgAwIBAgIUJSd8PABgAeO4DR52EF1GLaff97wwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI1MTIyMTExWhcNMjYwODI1MTIzMTExWjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAE2J95rQYUscFEvmng+Txabpy9Lj0wMF/PimL1V6P4s2zZpQ3D7BjImDIfhWkbYvHMrhygi05Wkhim+0xt5oouQ6OCBe8wggXrMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUptKQ2QGeJeH6CPYt8nO+uScnds8wHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGM0ZDBkYzU2NjYxYmFkYjljYzQ3YmRlNjQ3ZTA4ZmUzYjRiZjc0NTgwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzI4NDQ2ODk4OTQvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBiQYKKwYBBAHWeQIEAgR7BHkAdwB1AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoDjeDocAAAQDAEYwRAIgWfVz6msxL2/OLXywayXxyN3KQNSvgR9bgBEc7ExLXdkCIGNdn4O880wqZp54Q9VpzrJiR/dTGq2y7FR9N/ibNWvdMAoGCCqGSM49BAMDA2kAMGYCMQCoo+jPr8Se4pyJ+0XE95hXL+oyXF0XYu/97P7CoQkVnNz2gEya98dJ/nJF/8uoRzgCMQDd7PPDdEeXiupUhD+BMPTgDzrPFBlVaErHjqXPcKKDanuZwYgfa/NnwLddlkPwOM8="}, "tlogEntries":[{"logIndex":"2583302524", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787660472", "inclusionPromise":{"signedEntryTimestamp":"MEUCID1/KL7+8Sn5TlVyvNAwioEK69P6m5YjZakMUiv5nLZvAiEAlW8nULrUO9h2vPsZwA6T1LKgVXaD2EA3gYREf4Ywp+s="}, "inclusionProof":{"logIndex":"2461398262", "rootHash":"2pXLX5oZTZ/BWtgov4MXYlbWCXrXtoYb3I8T6XBsuig=", "treeSize":"2461398263", "hashes":["mnJetiTJfE1DvJ7lhBRN8OaQX3jLQeNGwI+QLXp0AAM=", "Lh/MtYSxJ8LFybPOQU/g+Frcanyv8Lw29x4fY6wojJc=", "HSg8Qe+EVcgHW9HM8VhTc/rLdK4ZgsiJwWaUtkCDg40=", "dOVbn4NZnSCPsjEt7IgIPnRcYFBIESrnojSV4l++OhI=", "9i5IUFCaI3RtgG9ZfyQpY4loPB/G4SrqefjRx2QlesY=", "s/NH5Bz76m3CS1MP63DS1Iy77fx+9fsdHWF98zIEtG4=", "r8ZmsIkPd8oW6c3MXr06PoZ4nJYw0Vt5ki3ZDt5yOEM=", "M0t6Pyxtbr61GCmNg60qyyc/qJAEEt98xDnkIgI5BIY=", "EY8T0lxQDrmYoWBGpKNGttaHfKmpfqKcWqSNF4XOQKE=", "T4Jz+m9vO48rOpMSsdZSW28/bgHR8prz7JeNX1s7WiQ=", "EXKrSbeG2qm5APD5QSDN66X7UHz9OSofiW9PyNU7fPM=", "cjIeeutzThxIPTlFFHJqa7bAzn3k2FzK8Nrdd56OWvE=", "+pNXWGIXzU1dqSy/aePHKYtSkDnFinI+/bL5HnurJ6w=", "G1R9F5B2KizFA2NyhDznjNKXPOF1SDmbGHz+omLG7TY=", "RZ7XwGuqMr5gDUh2HpUcaXp+AmIiEcPJiMui4GLJE7o=", "mBm+vQtn0C4thMnlTnxfxqowq1dXsPBCKaJ85da3JeU=", "b80/J/68RC8/tx5xRdlKk6MmTpZiVRotquWBrE/z1pI=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2461398263\n2pXLX5oZTZ/BWtgov4MXYlbWCXrXtoYb3I8T6XBsuig=\n\n— rekor.sigstore.dev wNI9ajBEAiAq0xstSQYZOi7qccIojmwFVRzfYnlEejjc7NuwX0OxdwIgAQ4/9zdGDLj4el6CbJOci3IC2YMVwKWxfTHGeZ/05kQ=\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiJlN2ZlMzg0OTI4ODY1MTI4ZjNiMTc2NTkwMjI1YzcxMGY3YjVjZDhkNmRhOWFlMWFhODlmN2Y0Y2I1ODUyOGJmIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJQ3RzNHhPcVUvYVdwU2Q5K3lvRXBnamk4ZTY0S0tza3pid3I5ZHBqUkFqa0FpRUE4YmxoUTgwZ0FXMWNubUZKSGJFdS9tTjNOVlh0amdZcTMyUTFuRjNGVjlFPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRla05EUW5SRFowRjNTVUpCWjBsVlNsTmtPRkJCUW1kQlpVODBSRkkxTWtWR01VZE1ZV1ptT1RkM2QwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1RGTlZFbDVUVlJGZUZkb1kwNU5hbGwzVDBSSk1VMVVTWHBOVkVWNFYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVV5U2prMWNsRlpWWE5qUmtWMmJXNW5LMVI0WVdKd2VUbE1hakIzVFVZdlVHbHRUREVLVmpaUU5ITXllbHB3VVRORU4wSnFTVzFFU1dab1YydGlXWFpJVFhKb2VXZHBNRFZYYTJocGJTc3dlSFExYjI5MVVUWlBRMEpsT0hkbloxaHlUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZ3ZEV0UkNqSlJSMlZLWlVnMlExQlpkRGh1VHl0MVUyTnVaSE00ZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVplbEpyVFVkU2FrNVVXVEpPYWtacFdWZFNhVTlYVG1wT1JHUnBDbHBIVlRKT1JHUnNUVVJvYlZwVVRtbE9SMHB0VG5wUk1VOUVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2xhUjFVeVRrUmtiRTFFYUcwS1dsUk9hVTVIU20xT2VsRXhUMFJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2tLV2tkVk1rNUVaR3hOUkdodFdsUk9hVTVIU20xT2VsRXhUMFJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZE5NRnBFUW1zS1dYcFZNazVxV1hoWmJVWnJXV3BzYWxsNlVUTlpiVkpzVG1wUk0xcFVRVFJhYlZWNldXcFNhVnBxWXpCT1ZHZDNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWtrMFRrUlJNazlFYXpSUFZGRjJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwVVZsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTTjBKSWEwRUtaSGRDTVVGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUVhbVZFYjJOQlFVRlJSQXBCUlZsM1VrRkpaMWRtVm5vMmJYTjRUREl2VDB4WWVYZGhlVmg0ZVU0elMxRk9VM1puVWpsaVowSkZZemRGZUV4WVpHdERTVWRPWkc0MFR6ZzRNSGR4Q2xwd05UUlJPVlp3ZW5KS2FWSXZaRlJIY1RKNU4wWlNPVTR2YVdKT1YzWmtUVUZ2UjBORGNVZFRUVFE1UWtGTlJFRXlhMEZOUjFsRFRWRkRiMjhyYWxBS2NqaFRaVFJ3ZVVvck1GaEZPVFZvV0V3cmIzbFlSakJZV1hVdk9UZFFOME52VVd0V2JrNTZNbWRGZVdFNU9HUktMMjVLUmk4NGRXOVNlbWREVFZGRVpBbzNVRkJFWkVWbFdHbDFjRlZvUkN0Q1RWQlVaMFI2Y2xCR1FteFdZVVZ5U0dweFdGQmpTMHRFWVc1MVduZFpaMlpoTDA1dWQweGtaR3hyVUhkUFRUZzlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyjADAgEAMIICwQYJKoZIhvcNAQcCoIICsjCCAq4CAQMxDTALBglghkgBZQMEAgEwgbgGCyqGSIb3DQEJEAEEoIGoBIGlMIGiAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgbqaIhtv8tLzbp0MAtPYDCkxi33gaCUYaaabJdr6QlKgCFQDt+ch1DN+Vj7pBE+G0P7QGTeilCRgPMjAyNjA4MjUxMjIxMTJaMAMCAQGgMqQwMC4xFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEVMBMGA1UEAxMMc2lnc3RvcmUtdHNhoAAxggHbMIIB1wIBATBRMDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAsGCWCGSAFlAwQCAaCB/DAaBgkqhkiG9w0BCQMxDQYLKoZIhvcNAQkQAQQwHAYJKoZIhvcNAQkFMQ8XDTI2MDgyNTEyMjExMlowLwYJKoZIhvcNAQkEMSIEIPV/JPEvcqAyNDTd/+BlkFTb7Y0K9+yO04ncRdnvH9omMIGOBgsqhkiG9w0BCRACLzF/MH0wezB5BCCF+Se8B6tiysO0Q1bBDvyBssaIP9p6uebYcNnROs0FtzBVMD2kOzA5MRUwEwYDVQQKEwxzaWdzdG9yZS5kZXYxIDAeBgNVBAMTF3NpZ3N0b3JlLXRzYS1zZWxmc2lnbmVkAhQ6E1QvDJBh7rzBQy/Lio6LKiOLDDAKBggqhkjOPQQDAgRnMGUCMQC+7AsRpZOTkeH34vW6OtA0R/RonZqK8Xjw1a+dGVAo66xxxZZQEuoiwaP3UGXNb/cCMEQZ5CLn5qpxfBtepi1chVLXH8KtmODxOeH7q9tnli4NFxz36Nqm1seveqvAIlo1cA=="}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"5/44SSiGUSjzsXZZAiXHEPe1zY1tqa4aqJ9/TLWFKL8="}, "signature":"MEUCICts4xOqU/aWpSd9+yoEpgji8e64KKskzbwr9dpjRAjkAiEA8blhQ80gAW1cnmFJHbEu/mN3NVXtjgYq32Q1nF3FV9E="}}
|
|
|
|
| 1 |
+
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSzCCBtGgAwIBAgIUA5vul1TIVVVD/55ZQsCMvclTsPUwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI3MDcxNDEzWhcNMjYwODI3MDcyNDEzWjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEVsmxRDo5AYOQSQTAsN2aOorB4owKn6xe+XSzgWPglS4aVOohNjIFEYxRMkqzqbLIUJofcHvWu8shwCHi4JmWTqOCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUsDNkj11Wgu9wXqXUpF8F/w4f0wcwHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGFkZGY0NzQxZmFmYjU1YjRjMTAxYjE1OTg1ZGZhZWVhN2M3YzJmM2EwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzMwNDUzOTU5NTIvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoEIRvWoAAAQDAEcwRQIgcR/4wLDiD9qUy4y2vs5kFgcpq3kesrIJjw/f6qC8SuICIQCQNMl1mFo+raHGdMB/QQOkCMRB/d5/fRndBzqLmt7acDAKBggqhkjOPQQDAwNoADBlAjEA/dfyWGtIrwMFzbux13DZsA2hLHUD7kRNRACuN25oBY/2BGqzCB1292Z2oSMLGJa7AjBOiQuqr2faYk9oCfREmjKm8SrbxHWb8WuHAIOgFxGPwADaAHQY/zpoV1v6l5u2blY="}, "tlogEntries":[{"logIndex":"2614048847", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787814854", "inclusionPromise":{"signedEntryTimestamp":"MEQCIHe4p5RPWriaW3uM0phVLA+qZ4hBe+Ywgg82OvTwnOGxAiBuzfj9q2eRheokCj2W33ae37fNGXztKcRriN1wvcjUfg=="}, "inclusionProof":{"logIndex":"2492144585", "rootHash":"GpTg7G2pAAx6ltSe5p5w+mLW5RAZyDa/fxDDYSwnGn0=", "treeSize":"2492144599", "hashes":["Del7EEbYya1uFE6Q2kAS3OZq3gwg3dljdLkc01DaSZY=", "5sL2HZeS6NipYIJ4ciW3HjdMkIB30MyPQQRBPD4VaKI=", "8W861Mu2l0lr8MJykHoaF+w8FUpt2oO2MdJmZMKDGSA=", "bKnOUb0K3WsY+a8UEqANcRu2zO9a82a7lX7Ddi4P21c=", "t+Lp+CPy++8B/TPyzeRJUGruZvhM02wHEpZ6NSeU5FY=", "grNFx/bqJovsvgsk3uMs3EWSG8h43kjyg6PE/8HaWgM=", "izo8eXXW3X4Ups0JGdW/otWtJOUWzCS879a5SRzjnkg=", "SWVtAPePcNx5Tmz8BPDzSaRncMca+mcKIfZKQxBl1/s=", "gxi1eJ8qX1K1iauxvafmRFx1FhOcSvz3IdnMqR1qs9E=", "VFlTxnDU4W+k/WA3ARTvzCiqVt1yxQIMUyUzMCugJ2Q=", "CFNMzmJ8N+uNx37P2asMqvrW9tw4axAVHnXLRO39tFs=", "NypqAbB+mNix4l1ahGeP9fRWiQzZCm4P1664makjjq4=", "8SoaeA7IeJPoIiveQbBfHCWXSDQabIk2PW+FT49NSEQ=", "VGrFt5JvlIUQvoGEdBBxrCe3lYbu581OhgLq/Xdyat4=", "7TfDCaMWgEwSJEZd/e33b24919rF4IDJLmbqNNRNhZs=", "0HfdxDt/zGugZLIuHasrdLEW9s8OaTykSSqBEDFQXdw=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2492144599\nGpTg7G2pAAx6ltSe5p5w+mLW5RAZyDa/fxDDYSwnGn0=\n\n— rekor.sigstore.dev wNI9ajBFAiBxt2lmDypDniaus9sV/HRm7QaAQVOoQtAxiks24mYxvAIhAJBcxCC2pqSDxOdGVGar+cUrIT3Ta9O19cw+35PZR7hQ\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiJjOWJlNTllYjMyMWQ2Y2E5ZGE4NzgxMTIxY2RmMjQwMjMzNDU1NTE2NTc3ZDUzMzY0ZmNjODUxNWJkZTVhNDUxIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJUUM1bS9tZ1FQODk5WjBNMWFPWGQ1NDBwdmY1WmhjZk1zZjVWaGticEhMZkNBSWdXcCtabUdnVEFocG1lQkNaZ0tISWkxUlIwTUVkQmFDOEVWRVpQYUo4S2NJPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRla05EUW5SSFowRjNTVUpCWjBsVlFUVjJkV3d4VkVsV1ZsWkVMelUxV2xGelEwMTJZMnhVYzFCVmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1ROTlJHTjRUa1JGZWxkb1kwNU5hbGwzVDBSSk0wMUVZM2xPUkVWNlYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZXYzIxNFVrUnZOVUZaVDFGVFVWUkJjMDR5WVU5dmNrSTBiM2RMYmpaNFpTdFlVM29LWjFkUVoyeFROR0ZXVDI5b1RtcEpSa1ZaZUZKTmEzRjZjV0pNU1ZWS2IyWmpTSFpYZFRoemFIZERTR2swU20xWFZIRlBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZ6UkU1ckNtb3hNVmRuZFRsM1dIRllWWEJHT0VZdmR6Um1NSGRqZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVpWMUpyV21wUk0wNUVSbTFaVjFwcFRsUldhVTVIVFhoTlJFWnBDazFVVlRWUFJGWnJXbTFHYkZwWFJUTlplbVJxVFcxWmVsbFVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWxOVkZVMVQwUldhMXB0Um13S1dsZEZNMWw2WkdwTmJWbDZXVlJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWtLVFZSVk5VOUVWbXRhYlVac1dsZEZNMWw2WkdwTmJWbDZXVlJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZEdhMXBIV1RBS1RucFJlRnB0Um0xWmFsVXhXV3BTYWsxVVFYaFpha1V4VDFSbk1WcEhXbWhhVjFab1RqSk5NMWw2U20xTk1rVjNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWsxM1RrUlZlazlVVlRWT1ZFbDJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUZTVkoyVjI5QlFVRlJSQXBCUldOM1VsRkpaMk5TTHpSM1RFUnBSRGx4VlhrMGVUSjJjelZyUm1kamNIRXphMlZ6Y2tsS2FuY3Zaalp4UXpoVGRVbERTVkZEVVU1TmJERnRSbThyQ25KaFNFZGtUVUl2VVZGUGEwTk5Va0l2WkRVdlpsSnVaRUo2Y1V4dGREZGhZMFJCUzBKblozRm9hMnBQVUZGUlJFRjNUbTlCUkVKc1FXcEZRUzlrWm5rS1YwZDBTWEozVFVaNlluVjRNVE5FV25OQk1taE1TRlZFTjJ0U1RsSkJRM1ZPTWpWdlFsa3ZNa0pIY1hwRFFqRXlPVEphTW05VFRVeEhTbUUzUVdwQ1R3cHBVWFZ4Y2pKbVlWbHJPVzlEWmxKRmJXcExiVGhUY21KNFNGZGlPRmQxU0VGSlQyZEdlRWRRZDBGRVlVRklVVmt2ZW5CdlZqRjJObXcxZFRKaWJGazlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyTADAgEAMIICwAYJKoZIhvcNAQcCoIICsTCCAq0CAQMxDTALBglghkgBZQMEAgEwgbcGCyqGSIb3DQEJEAEEoIGnBIGkMIGhAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgivTPiz46ImUVRh/dqA/JKYgdKFOkokSgBmdjgH2WLIECFFTGQ0mtdpm+MoVyxznY6glXNZ/uGA8yMDI2MDgyNzA3MTQxNFowAwIBAaAypDAwLjEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MRUwEwYDVQQDEwxzaWdzdG9yZS10c2GgADGCAdswggHXAgEBMFEwOTEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MSAwHgYDVQQDExdzaWdzdG9yZS10c2Etc2VsZnNpZ25lZAIUOhNULwyQYe68wUMvy4qOiyojiwwwCwYJYIZIAWUDBAIBoIH8MBoGCSqGSIb3DQEJAzENBgsqhkiG9w0BCRABBDAcBgkqhkiG9w0BCQUxDxcNMjYwODI3MDcxNDE0WjAvBgkqhkiG9w0BCQQxIgQgLKg7YOS45PuElnTBE7vy6grALFrzeI6kw9/l+MXo6P8wgY4GCyqGSIb3DQEJEAIvMX8wfTB7MHkEIIX5J7wHq2LKw7RDVsEO/IGyxog/2nq55thw2dE6zQW3MFUwPaQ7MDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAoGCCqGSM49BAMCBGcwZQIxAN4sMIDl23RF+zZffpzr7zVVJvpcQsFr4/Xr1pG1KcudoghTp531eDpJEshCwwDihAIwXCxgv9HRlvtVb/dQP9pSvOJZ83gnQeIm87q5WPCh9N+BVpOyYuX/rsqINPRHe/2g"}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"yb5Z6zIdbKnah4ESHN8kAjNFVRZXfVM2T8yFFb3lpFE="}, "signature":"MEUCIQC5m/mgQP899Z0M1aOXd540pvf5ZhcfMsf5VhkbpHLfCAIgWp+ZmGgTAhpmeBCZgKHIi1RR0MEdBaC8EVEZPaJ8KcI="}}
|
build/{torch212-cxx11-cu130-x86_64-linux/_megablocks_cuda_c4d0dc5.abi3.so → torch212-cxx11-cu132-x86_64-linux/_megablocks_cuda_addf474.abi3.so}
RENAMED
|
@@ -1,3 +1,3 @@
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
-
oid sha256:
|
| 3 |
-
size
|
|
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:d6f0251bf07e0c3d494c14363a5b387c25fc59b308c80752270a76ee177a72f5
|
| 3 |
+
size 14986840
|
build/torch212-cxx11-cu132-x86_64-linux/_ops.py
CHANGED
|
@@ -1,9 +1,9 @@
|
|
| 1 |
import torch
|
| 2 |
-
from . import
|
| 3 |
-
ops = torch.ops.
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
-
return f"
|
|
|
|
| 1 |
import torch
|
| 2 |
+
from . import _megablocks_cuda_addf474
|
| 3 |
+
ops = torch.ops._megablocks_cuda_addf474
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
+
return f"_megablocks_cuda_addf474::{op_name}"
|
build/torch212-cxx11-cu132-x86_64-linux/megablocks/__init__.py
DELETED
|
@@ -1,26 +0,0 @@
|
|
| 1 |
-
import ctypes
|
| 2 |
-
import importlib.util
|
| 3 |
-
import sys
|
| 4 |
-
from pathlib import Path
|
| 5 |
-
from types import ModuleType
|
| 6 |
-
|
| 7 |
-
|
| 8 |
-
def _import_from_path(file_path: Path) -> ModuleType:
|
| 9 |
-
# We cannot use the module name as-is, after adding it to `sys.modules`,
|
| 10 |
-
# it would also be used for other imports. So, we make a module name that
|
| 11 |
-
# depends on the path for it to be unique using the hex-encoded hash of
|
| 12 |
-
# the path.
|
| 13 |
-
path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
|
| 14 |
-
module_name = path_hash
|
| 15 |
-
spec = importlib.util.spec_from_file_location(module_name, file_path)
|
| 16 |
-
if spec is None:
|
| 17 |
-
raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
|
| 18 |
-
module = importlib.util.module_from_spec(spec)
|
| 19 |
-
if module is None:
|
| 20 |
-
raise ImportError(f"Cannot load module {module_name} from spec")
|
| 21 |
-
sys.modules[module_name] = module
|
| 22 |
-
spec.loader.exec_module(module) # type: ignore
|
| 23 |
-
return module
|
| 24 |
-
|
| 25 |
-
|
| 26 |
-
globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
build/torch212-cxx11-cu132-x86_64-linux/metadata.json
CHANGED
|
@@ -1,14 +1,18 @@
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
|
|
|
| 7 |
"backend": {
|
| 8 |
"type": "cuda",
|
| 9 |
"archs": [
|
| 10 |
"10.0",
|
|
|
|
|
|
|
| 11 |
"12.0",
|
|
|
|
| 12 |
"7.5",
|
| 13 |
"8.0",
|
| 14 |
"8.6",
|
|
@@ -48,8 +52,8 @@
|
|
| 48 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 49 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 50 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 51 |
-
"
|
| 52 |
-
"_ops.py": "
|
| 53 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 54 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 55 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
@@ -61,7 +65,6 @@
|
|
| 61 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 62 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 63 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
| 64 |
-
"megablocks/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY=",
|
| 65 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 66 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 67 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
@@ -106,11 +109,11 @@
|
|
| 106 |
"provenance": {
|
| 107 |
"kernel-builder": {
|
| 108 |
"version": "0.17.0-dev0",
|
| 109 |
-
"
|
| 110 |
"dirty": false
|
| 111 |
},
|
| 112 |
"kernel": {
|
| 113 |
-
"
|
| 114 |
"dirty": false
|
| 115 |
}
|
| 116 |
}
|
|
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
+
"id": "_megablocks_cuda_addf474",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
| 7 |
+
"kernel-depends": [],
|
| 8 |
"backend": {
|
| 9 |
"type": "cuda",
|
| 10 |
"archs": [
|
| 11 |
"10.0",
|
| 12 |
+
"10.3",
|
| 13 |
+
"11.0",
|
| 14 |
"12.0",
|
| 15 |
+
"12.1+PTX",
|
| 16 |
"7.5",
|
| 17 |
"8.0",
|
| 18 |
"8.6",
|
|
|
|
| 52 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 53 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 54 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 55 |
+
"_megablocks_cuda_addf474.abi3.so": "1vAlG/B+DD1JTBQ2Ols4fCX8WbMIyAdSJwp27hd6cvU=",
|
| 56 |
+
"_ops.py": "WXlV4XwtKQjdQ6nap3k9TalaHkK0ypeSTshOrV6Hczo=",
|
| 57 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 58 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 59 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
|
|
| 65 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 66 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 67 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
|
|
|
| 68 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 69 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 70 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
|
|
| 109 |
"provenance": {
|
| 110 |
"kernel-builder": {
|
| 111 |
"version": "0.17.0-dev0",
|
| 112 |
+
"commit": "7d2828e6592de641b44fb8eb896719101a2ab2fb",
|
| 113 |
"dirty": false
|
| 114 |
},
|
| 115 |
"kernel": {
|
| 116 |
+
"commit": "addf4741fafb55b4c101b15985dfaeea7c7c2f3a",
|
| 117 |
"dirty": false
|
| 118 |
}
|
| 119 |
}
|
build/torch212-cxx11-cu132-x86_64-linux/metadata.json.sigstore
CHANGED
|
@@ -1 +1 @@
|
|
| 1 |
-
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSTCCBtCgAwIBAgIUOHsK45h2ipRMvVeaYZD9Guc/o+YwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI1MTIyMTEzWhcNMjYwODI1MTIzMTEzWjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEwhAxJFNYMtNxoPBYzmXY1tZH/5P3ixKgQy1erUTFhE+XzHlDIfF66u3qpKGGN3b3zQTI+seDAopqe4L2F+W6SqOCBe8wggXrMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUYsVomHv+UP1q+6iomqYhMO+nRLIwHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGM0ZDBkYzU2NjYxYmFkYjljYzQ3YmRlNjQ3ZTA4ZmUzYjRiZjc0NTgwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzI4NDQ2ODk4OTQvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBiQYKKwYBBAHWeQIEAgR7BHkAdwB1AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoDjeFTkAAAQDAEYwRAIgHpfpogDCG6+DnDRUVo3Cvfswjjg85Fn0xaVr1DlVHSECIE5+H/mpf0H51UrS7Gep82qlBLN4t7bY3M5qRv1hvfEbMAoGCCqGSM49BAMDA2cAMGQCMCH9dPE9TAtitu85ZilE+xvqEqQtKVnS1ecYBQrLUvXLmFN6Qsjkx0A7CDjPMsF3WgIwbpeg5+7b7RzyLu8V2hP8VJTynFsiI8lXfY6CJov8ZPdWXk7+BYHshpkwXzIvZxQX"}, "tlogEntries":[{"logIndex":"2583302548", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787660473", "inclusionPromise":{"signedEntryTimestamp":"MEUCIQC81EpDHYqwD0Ft814fa9bgywabWZa1rmGEEbdi6ypdSwIgZS33SeJwbkil3sLdO+EK9OcVYynzONQJxEDc++4blyI="}, "inclusionProof":{"logIndex":"2461398286", "rootHash":"Nd9xPnfnKztd4KuBSM1XAUEh09WalV2yViS27323Qto=", "treeSize":"2461398287", "hashes":["f03VBwyMMrtujA9uh+znFCJ+YJnwPonyHC+iDq1K4h4=", "MVjt7QvYVWoDUKMZYe5kWdnj5kTP6cmEbStyAuhnZ+E=", "HXuvjlv+weMklWPM3GC0X93XS19fBuxDsVxUTRlaKFU=", "eUK42vZbtTAA+teDm0LcO+qURyI8IPFvT/lz1znHKQ8=", "r8ZmsIkPd8oW6c3MXr06PoZ4nJYw0Vt5ki3ZDt5yOEM=", "M0t6Pyxtbr61GCmNg60qyyc/qJAEEt98xDnkIgI5BIY=", "EY8T0lxQDrmYoWBGpKNGttaHfKmpfqKcWqSNF4XOQKE=", "T4Jz+m9vO48rOpMSsdZSW28/bgHR8prz7JeNX1s7WiQ=", "EXKrSbeG2qm5APD5QSDN66X7UHz9OSofiW9PyNU7fPM=", "cjIeeutzThxIPTlFFHJqa7bAzn3k2FzK8Nrdd56OWvE=", "+pNXWGIXzU1dqSy/aePHKYtSkDnFinI+/bL5HnurJ6w=", "G1R9F5B2KizFA2NyhDznjNKXPOF1SDmbGHz+omLG7TY=", "RZ7XwGuqMr5gDUh2HpUcaXp+AmIiEcPJiMui4GLJE7o=", "mBm+vQtn0C4thMnlTnxfxqowq1dXsPBCKaJ85da3JeU=", "b80/J/68RC8/tx5xRdlKk6MmTpZiVRotquWBrE/z1pI=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2461398287\nNd9xPnfnKztd4KuBSM1XAUEh09WalV2yViS27323Qto=\n\n— rekor.sigstore.dev wNI9ajBGAiEAowOyAt2Zb5k4VPUV70DaxxZgHbHEk5IiK8BEzCFFs9oCIQDQbBghR365IJIaOMaHhO1aNufrIrpfrvKbxER0I+6/1g==\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiJjZGEyYWUxOTE4N2M5MTdhNjczMzE3MjEzMTNiMThmMzVkMGM5ZTZmZmJjNGM4MTI5Zjk5MTczYTA1YWVlMTY1In19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJQWtnRStqUXEyL09qU21CZTl6c3Rydm94SEVLSDIzREVkcjg3M0pTU0NuUkFpRUFqV25GTlBacGY0WlBNSE1KVFQ3RHRXTjZjdE1GemRMQTY1ZW8vQVdrMTFvPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRWRU5EUW5SRFowRjNTVUpCWjBsVlQwaHpTelExYURKcGNGSk5kbFpsWVZsYVJEbEhkV012Ynl0WmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1RGTlZFbDVUVlJGZWxkb1kwNU5hbGwzVDBSSk1VMVVTWHBOVkVWNlYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVYzYUVGNFNrWk9XVTEwVG5odlVFSlplbTFZV1RGMFdrZ3ZOVkF6YVhoTFoxRjVNV1VLY2xWVVJtaEZLMWg2U0d4RVNXWkdOaloxTTNGd1MwZEhUak5pTTNwUlZFa3JjMlZFUVc5d2NXVTBUREpHSzFjMlUzRlBRMEpsT0hkbloxaHlUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZaYzFadkNtMUlkaXRWVURGeEt6WnBiMjF4V1doTlR5dHVVa3hKZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVplbEpyVFVkU2FrNVVXVEpPYWtacFdWZFNhVTlYVG1wT1JHUnBDbHBIVlRKT1JHUnNUVVJvYlZwVVRtbE9SMHB0VG5wUk1VOUVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2xhUjFVeVRrUmtiRTFFYUcwS1dsUk9hVTVIU20xT2VsRXhUMFJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2tLV2tkVk1rNUVaR3hOUkdodFdsUk9hVTVIU20xT2VsRXhUMFJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZE5NRnBFUW1zS1dYcFZNazVxV1hoWmJVWnJXV3BzYWxsNlVUTlpiVkpzVG1wUk0xcFVRVFJhYlZWNldXcFNhVnBxWXpCT1ZHZDNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWtrMFRrUlJNazlFYXpSUFZGRjJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwVVZsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTTjBKSWEwRUtaSGRDTVVGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUVhbVZHVkd0QlFVRlJSQXBCUlZsM1VrRkpaMGh3Wm5CdlowUkRSellyUkc1RVVsVldiek5EZG1aemQycHFaemcxUm00d2VHRldjakZFYkZaSVUwVkRTVVUxSzBndmJYQm1NRWcxQ2pGVmNsTTNSMlZ3T0RKeGJFSk1UalIwTjJKWk0wMDFjVkoyTVdoMlprVmlUVUZ2UjBORGNVZFRUVFE1UWtGTlJFRXlZMEZOUjFGRFRVTklPV1JRUlRrS1ZFRjBhWFIxT0RWYWFXeEZLM2gyY1VWeFVYUkxWbTVUTVdWaldVSlJja3hWZGxoTWJVWk9ObEZ6YW10NE1FRTNRMFJxVUUxelJqTlhaMGwzWW5CbFp3bzFLemRpTjFKNmVVeDFPRll5YUZBNFZrcFVlVzVHYzJsSk9HeFlabGsyUTBwdmRqaGFVR1JYV0dzM0swSlpTSE5vY0d0M1dIcEpkbHA0VVZnS0xTMHRMUzFGVGtRZ1EwVlNWRWxHU1VOQlZFVXRMUzB0TFFvPSJ9fX19"}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyTADAgEAMIICwAYJKoZIhvcNAQcCoIICsTCCAq0CAQMxDTALBglghkgBZQMEAgEwgbcGCyqGSIb3DQEJEAEEoIGnBIGkMIGhAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgIVfsLoyRiTisbRWNoXaZi4aOrfyvcYDrzbHIRZwfkx8CFFkI2OD0ue+GJevJ9BqXxzyQc4/HGA8yMDI2MDgyNTEyMjExM1owAwIBAaAypDAwLjEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MRUwEwYDVQQDEwxzaWdzdG9yZS10c2GgADGCAdswggHXAgEBMFEwOTEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MSAwHgYDVQQDExdzaWdzdG9yZS10c2Etc2VsZnNpZ25lZAIUOhNULwyQYe68wUMvy4qOiyojiwwwCwYJYIZIAWUDBAIBoIH8MBoGCSqGSIb3DQEJAzENBgsqhkiG9w0BCRABBDAcBgkqhkiG9w0BCQUxDxcNMjYwODI1MTIyMTEzWjAvBgkqhkiG9w0BCQQxIgQg73jlbewycgDn6ZZtcuFL6G47SQNo1LEi6cvo2v8RA6UwgY4GCyqGSIb3DQEJEAIvMX8wfTB7MHkEIIX5J7wHq2LKw7RDVsEO/IGyxog/2nq55thw2dE6zQW3MFUwPaQ7MDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAoGCCqGSM49BAMCBGcwZQIwHJK4VKQvu4D8/jY5NxiK+tLnv7SjCF3UvSugLWpdplcPmNXLTawo8u67zCb0r0ofAjEA/HNmVSHIv25cz/JiecPsdW5IvaxbEO9sYtL1+g5CMbKuZSTYb7GKm+YrS7GoFQRg"}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"zaKuGRh8kXpnMxchMTsY810Mnm/7xMgSn5kXOgWu4WU="}, "signature":"MEUCIAkgE+jQq2/OjSmBe9zstrvoxHEKH23DEdr873JSSCnRAiEAjWnFNPZpf4ZPMHMJTT7DtWN6ctMFzdLA65eo/AWk11o="}}
|
|
|
|
| 1 |
+
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSzCCBtGgAwIBAgIUTZAc00cHnQe/D7dPNb1Q61JdLRcwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI3MDcxNDA3WhcNMjYwODI3MDcyNDA3WjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAER6YYvzSORzAOvMUg91NjQN/meDWhZESTyVOZvpYArdAJb6XpRxyVKCP2NTiCQb4S67w+SY9aCYbIvU57VJKICKOCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUP2uJ7tbXEHAHclfmgpCp2iOIXMgwHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGFkZGY0NzQxZmFmYjU1YjRjMTAxYjE1OTg1ZGZhZWVhN2M3YzJmM2EwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzMwNDUzOTU5NTIvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoEIRoyUAAAQDAEcwRQIhANfV92DP21CqpyxEhaIYTKT80B+7DpLmkYIVJ/jqf3tLAiA+sKwiSftQXcOtUSFXjbpIVG/buDUxcIygJrUH9SMGqjAKBggqhkjOPQQDAwNoADBlAjB4+0MVVKnyueeobqse6wpT5yPDhp3HKMbrXdP6ARODY8CURXf/BmrpkLQkMdILBrYCMQDqp0crYZVYeF7P7H/ZjCEWt/laFosGerr4MQafJsxUxnQOvslUyqjcBO9PywyCmhk="}, "tlogEntries":[{"logIndex":"2614047458", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787814847", "inclusionPromise":{"signedEntryTimestamp":"MEQCIF8DaUEUunXnJ+c+NNaVDJYJDPNlO+5ARYoXx/JQxom2AiA1/fS0ErekdpIbpEs+yb8MhRsggClHw/WAGXRFtBoYWg=="}, "inclusionProof":{"logIndex":"2492143196", "rootHash":"elS9QbwVGQqQJHInKI82veZ0COxI5mIm59byPy5PjBw=", "treeSize":"2492143250", "hashes":["3lEfJ3x+2iIiESY9vPI0YvNsiXwoBsw3autJ/n3SGsY=", "WCBemMt4uq8mCZe4bZaHUw1LM+iEc3lo8mAc28L890o=", "ssYcXlLFYUdrfTI+7BhrXFd5V+J7xUklg5BJOdqO4lk=", "UFJAxFXqb/5K0mZZYDspMHmhFDT9m8wcuAOHu0aRqGI=", "R2v+lhN5WjL5xRYlra5RCdfYhV7+9NcKyleW6RMdRTA=", "qETILEsd+PvuF2Ml2mgjSHHcC/s5RjoNDn1C4V/ade8=", "hoqaFa+3T1Dpc6l587gp4kC8JxrVg1UxF1bRVVCC4Dc=", "nFudgbhky7nmBjpXigFw5/wyhomku3S8kQnvmOoP0us=", "74YOSoxcaARV3QmRlp2HcIEFoCumNC8/rDR31M3b3Zk=", "KIag3bHIhHqGqEJl+No8RWATx3rcwFMD/MdDkBYHPmQ=", "CFNMzmJ8N+uNx37P2asMqvrW9tw4axAVHnXLRO39tFs=", "NypqAbB+mNix4l1ahGeP9fRWiQzZCm4P1664makjjq4=", "8SoaeA7IeJPoIiveQbBfHCWXSDQabIk2PW+FT49NSEQ=", "VGrFt5JvlIUQvoGEdBBxrCe3lYbu581OhgLq/Xdyat4=", "7TfDCaMWgEwSJEZd/e33b24919rF4IDJLmbqNNRNhZs=", "0HfdxDt/zGugZLIuHasrdLEW9s8OaTykSSqBEDFQXdw=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2492143250\nelS9QbwVGQqQJHInKI82veZ0COxI5mIm59byPy5PjBw=\n\n— rekor.sigstore.dev wNI9ajBEAiBuapr9w5qEQSdRQyW3vAbWlWgOlHUq6egp0rb3DBhMsAIgdzVnoj2ySLgjrBC5UJzbcoYmQHpRVWG/kWiEliX2rl8=\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiJkMDUxZTAwMzQxYzQ5YTY5Mzc5ZTVlNDg4MGE1ODU4Y2ViYmIyZDE4ZmRlYjIxNzdmOWQyZjBhNzMyYzI4M2Y1In19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJRkN1bmhhZXMwMkxnU3JNREJYNytBb09yYmFFd29IaDYwZTNEa2U0QzMvcUFpRUEvWldiTW82RTZMQTVWbU1MbzYrL3U5cXFsZmlRdXgyVDh2TjJiWVFKSW1VPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRla05EUW5SSFowRjNTVUpCWjBsVlZGcEJZekF3WTBodVVXVXZSRGRrVUU1aU1WRTJNVXBrVEZKamQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1ROTlJHTjRUa1JCTTFkb1kwNU5hbGwzVDBSSk0wMUVZM2xPUkVFelYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZTTmxsWmRucFRUMUo2UVU5MlRWVm5PVEZPYWxGT0wyMWxSRmRvV2tWVFZIbFdUMW9LZG5CWlFYSmtRVXBpTmxod1VuaDVWa3REVURKT1ZHbERVV0kwVXpZM2R5dFRXVGxoUTFsaVNYWlZOVGRXU2t0SlEwdFBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZRTW5WS0NqZDBZbGhGU0VGSVkyeG1iV2R3UTNBeWFVOUpXRTFuZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVpWMUpyV21wUk0wNUVSbTFaVjFwcFRsUldhVTVIVFhoTlJFWnBDazFVVlRWUFJGWnJXbTFHYkZwWFJUTlplbVJxVFcxWmVsbFVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWxOVkZVMVQwUldhMXB0Um13S1dsZEZNMWw2WkdwTmJWbDZXVlJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWtLVFZSVk5VOUVWbXRhYlVac1dsZEZNMWw2WkdwTmJWbDZXVlJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZEdhMXBIV1RBS1RucFJlRnB0Um0xWmFsVXhXV3BTYWsxVVFYaFpha1V4VDFSbk1WcEhXbWhhVjFab1RqSk5NMWw2U20xTk1rVjNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWsxM1RrUlZlazlVVlRWT1ZFbDJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUZTVkp2ZVZWQlFVRlJSQXBCUldOM1VsRkphRUZPWmxZNU1rUlFNakZEY1hCNWVFVm9ZVWxaVkV0VU9EQkNLemRFY0V4dGExbEpWa292YW5GbU0zUk1RV2xCSzNOTGQybFRablJSQ2xoalQzUlZVMFpZYW1Kd1NWWkhMMkoxUkZWNFkwbDVaMHB5VlVnNVUwMUhjV3BCUzBKblozRm9hMnBQVUZGUlJFRjNUbTlCUkVKc1FXcENOQ3N3VFZZS1ZrdHVlWFZsWlc5aWNYTmxObmR3VkRWNVVFUm9jRE5JUzAxaWNsaGtVRFpCVWs5RVdUaERWVkpZWmk5Q2JYSndhMHhSYTAxa1NVeENjbGxEVFZGRWNRcHdNR055V1ZwV1dXVkdOMUEzU0M5YWFrTkZWM1F2YkdGR2IzTkhaWEp5TkUxUllXWktjM2hWZUc1UlQzWnpiRlY1Y1dwalFrODVVSGwzZVVOdGFHczlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyjADAgEAMIICwQYJKoZIhvcNAQcCoIICsjCCAq4CAQMxDTALBglghkgBZQMEAgEwgbgGCyqGSIb3DQEJEAEEoIGoBIGlMIGiAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgYrdu60lF/YEA9+445V4IwNdNGco3ByUoydVZAZpdHvYCFQCSs1zQcj+u6HVZoauDaeq5o3JgTBgPMjAyNjA4MjcwNzE0MDdaMAMCAQGgMqQwMC4xFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEVMBMGA1UEAxMMc2lnc3RvcmUtdHNhoAAxggHbMIIB1wIBATBRMDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAsGCWCGSAFlAwQCAaCB/DAaBgkqhkiG9w0BCQMxDQYLKoZIhvcNAQkQAQQwHAYJKoZIhvcNAQkFMQ8XDTI2MDgyNzA3MTQwN1owLwYJKoZIhvcNAQkEMSIEILeVojiyiC30P6ymBGoC1BDrJXn76/DsnZrjdFCgp6NtMIGOBgsqhkiG9w0BCRACLzF/MH0wezB5BCCF+Se8B6tiysO0Q1bBDvyBssaIP9p6uebYcNnROs0FtzBVMD2kOzA5MRUwEwYDVQQKEwxzaWdzdG9yZS5kZXYxIDAeBgNVBAMTF3NpZ3N0b3JlLXRzYS1zZWxmc2lnbmVkAhQ6E1QvDJBh7rzBQy/Lio6LKiOLDDAKBggqhkjOPQQDAgRnMGUCMQCArlEMkaTZxDanNxKF8yn42j/WsWxiPTTPcq3UxP9oVd2s9AJ/LSexVl45gmAYDZ4CMBw0Ld8ZV6RIWwy5lP5g0QaBucSmfUK6JEO2hVy2OM9IHFRqMaBodls1sDJFOMKoQg=="}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"0FHgA0HEmmk3nl5IgKWFjOu7LRj96yF3+dLwpzLCg/U="}, "signature":"MEUCIFCunhaes02LgSrMDBX7+AoOrbaEwoHh60e3Dke4C3/qAiEA/ZWbMo6E6LA5VmMLo6+/u9qqlfiQux2T8vN2bYQJImU="}}
|
build/torch213-cxx11-cu126-x86_64-linux/{_megablocks_cuda_c4d0dc5.abi3.so → _megablocks_cuda_addf474.abi3.so}
RENAMED
|
@@ -1,3 +1,3 @@
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
-
oid sha256:
|
| 3 |
-
size
|
|
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:6069e18a24d5fb1be504d649ce432a8052e1a2bae83896b9723ad581137559bd
|
| 3 |
+
size 13822376
|
build/torch213-cxx11-cu126-x86_64-linux/_ops.py
CHANGED
|
@@ -1,9 +1,9 @@
|
|
| 1 |
import torch
|
| 2 |
-
from . import
|
| 3 |
-
ops = torch.ops.
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
-
return f"
|
|
|
|
| 1 |
import torch
|
| 2 |
+
from . import _megablocks_cuda_addf474
|
| 3 |
+
ops = torch.ops._megablocks_cuda_addf474
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
+
return f"_megablocks_cuda_addf474::{op_name}"
|
build/torch213-cxx11-cu126-x86_64-linux/megablocks/__init__.py
DELETED
|
@@ -1,26 +0,0 @@
|
|
| 1 |
-
import ctypes
|
| 2 |
-
import importlib.util
|
| 3 |
-
import sys
|
| 4 |
-
from pathlib import Path
|
| 5 |
-
from types import ModuleType
|
| 6 |
-
|
| 7 |
-
|
| 8 |
-
def _import_from_path(file_path: Path) -> ModuleType:
|
| 9 |
-
# We cannot use the module name as-is, after adding it to `sys.modules`,
|
| 10 |
-
# it would also be used for other imports. So, we make a module name that
|
| 11 |
-
# depends on the path for it to be unique using the hex-encoded hash of
|
| 12 |
-
# the path.
|
| 13 |
-
path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
|
| 14 |
-
module_name = path_hash
|
| 15 |
-
spec = importlib.util.spec_from_file_location(module_name, file_path)
|
| 16 |
-
if spec is None:
|
| 17 |
-
raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
|
| 18 |
-
module = importlib.util.module_from_spec(spec)
|
| 19 |
-
if module is None:
|
| 20 |
-
raise ImportError(f"Cannot load module {module_name} from spec")
|
| 21 |
-
sys.modules[module_name] = module
|
| 22 |
-
spec.loader.exec_module(module) # type: ignore
|
| 23 |
-
return module
|
| 24 |
-
|
| 25 |
-
|
| 26 |
-
globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
build/torch213-cxx11-cu126-x86_64-linux/metadata.json
CHANGED
|
@@ -1,9 +1,10 @@
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
|
|
|
| 7 |
"backend": {
|
| 8 |
"type": "cuda",
|
| 9 |
"archs": [
|
|
@@ -14,7 +15,7 @@
|
|
| 14 |
"8.6",
|
| 15 |
"8.7",
|
| 16 |
"8.9",
|
| 17 |
-
"9.0"
|
| 18 |
]
|
| 19 |
},
|
| 20 |
"digest": {
|
|
@@ -48,8 +49,8 @@
|
|
| 48 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 49 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 50 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 51 |
-
"
|
| 52 |
-
"_ops.py": "
|
| 53 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 54 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 55 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
@@ -61,7 +62,6 @@
|
|
| 61 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 62 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 63 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
| 64 |
-
"megablocks/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY=",
|
| 65 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 66 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 67 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
@@ -106,11 +106,11 @@
|
|
| 106 |
"provenance": {
|
| 107 |
"kernel-builder": {
|
| 108 |
"version": "0.17.0-dev0",
|
| 109 |
-
"
|
| 110 |
"dirty": false
|
| 111 |
},
|
| 112 |
"kernel": {
|
| 113 |
-
"
|
| 114 |
"dirty": false
|
| 115 |
}
|
| 116 |
}
|
|
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
+
"id": "_megablocks_cuda_addf474",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
| 7 |
+
"kernel-depends": [],
|
| 8 |
"backend": {
|
| 9 |
"type": "cuda",
|
| 10 |
"archs": [
|
|
|
|
| 15 |
"8.6",
|
| 16 |
"8.7",
|
| 17 |
"8.9",
|
| 18 |
+
"9.0+PTX"
|
| 19 |
]
|
| 20 |
},
|
| 21 |
"digest": {
|
|
|
|
| 49 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 50 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 51 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 52 |
+
"_megablocks_cuda_addf474.abi3.so": "YGnhiiTV+xvlBNZJzkMqgFLhorroOJa5cjrVgRN1Wb0=",
|
| 53 |
+
"_ops.py": "WXlV4XwtKQjdQ6nap3k9TalaHkK0ypeSTshOrV6Hczo=",
|
| 54 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 55 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 56 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
|
|
| 62 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 63 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 64 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
|
|
|
| 65 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 66 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 67 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
|
|
| 106 |
"provenance": {
|
| 107 |
"kernel-builder": {
|
| 108 |
"version": "0.17.0-dev0",
|
| 109 |
+
"commit": "7d2828e6592de641b44fb8eb896719101a2ab2fb",
|
| 110 |
"dirty": false
|
| 111 |
},
|
| 112 |
"kernel": {
|
| 113 |
+
"commit": "addf4741fafb55b4c101b15985dfaeea7c7c2f3a",
|
| 114 |
"dirty": false
|
| 115 |
}
|
| 116 |
}
|
build/torch213-cxx11-cu126-x86_64-linux/metadata.json.sigstore
CHANGED
|
@@ -1 +1 @@
|
|
| 1 |
-
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSjCCBtCgAwIBAgIUPKexlqSPtfwNIuqQcSRd2Uk9IoMwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI1MTIyMTE1WhcNMjYwODI1MTIzMTE1WjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAE4o/8NCXKwvP+gL8EUDN9sggFCqFLCGbzcW4Bo/vCJRQGqWaIklQxd0rjc0yLT3bKeurramzfYL72OH+HMMltMaOCBe8wggXrMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQU4bfvGCLTJT31yr+EjuftwCM98xIwHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGM0ZDBkYzU2NjYxYmFkYjljYzQ3YmRlNjQ3ZTA4ZmUzYjRiZjc0NTgwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzI4NDQ2ODk4OTQvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBiQYKKwYBBAHWeQIEAgR7BHkAdwB1AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoDjeGpgAAAQDAEYwRAIgDQa0Ph1G8Ofv2HreWSKuOSVhSqrSPt6LjRa3HSbEDpcCIFudiw3vzHAun7tCNZg4ceUqfMfrp5tHQqJlEIF6WwWoMAoGCCqGSM49BAMDA2gAMGUCME0BcZCgWkgG0sxuuFOH0t3sFUvYw2GFT/pTBz/Npoh/C0V3UlIi1dvsT2FJwmJrWgIxAO5oAdpf7FkE4Fnf1OuJmrn14JJGEXAp6VprfG1ryrk32S4cTqFm8QzD2iv7EwRJvQ=="}, "tlogEntries":[{"logIndex":"2583302562", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787660475", "inclusionPromise":{"signedEntryTimestamp":"MEYCIQCMWp0psl6m849JTCxD8eYL5YE2PyAUhGC+093vNkJYhgIhANb3ZQBeXt9g2SdEY5e9At+Dm1YeiiwiBJPNwvKatXiI"}, "inclusionProof":{"logIndex":"2461398300", "rootHash":"MBJWZZdAgl5oAXtH7j6jZjGCmgQgxMYKm0nN+UNCmWc=", "treeSize":"2461398304", "hashes":["jlhtoFrVQec0Vb8sv+SpX9+8lOxKnkm6h0z72BSV20U=", "D3IhsQOOBII4ZIpbUfPcqF16+2UwsBKnfyMMsxfelVg=", "L+rX4UOu5EUK9rONsM+nzp0zMzsTFe7vryA0IWQQwqk=", "4QEWTGXLb0zUJxNlZdxkMX7QcytuFcYSeJ1ls224brY=", "dJdYtMdqZPZWIUgGrN0c7iVv8Th0woL+vXNyI9SrD3o=", "eUK42vZbtTAA+teDm0LcO+qURyI8IPFvT/lz1znHKQ8=", "r8ZmsIkPd8oW6c3MXr06PoZ4nJYw0Vt5ki3ZDt5yOEM=", "M0t6Pyxtbr61GCmNg60qyyc/qJAEEt98xDnkIgI5BIY=", "EY8T0lxQDrmYoWBGpKNGttaHfKmpfqKcWqSNF4XOQKE=", "T4Jz+m9vO48rOpMSsdZSW28/bgHR8prz7JeNX1s7WiQ=", "EXKrSbeG2qm5APD5QSDN66X7UHz9OSofiW9PyNU7fPM=", "cjIeeutzThxIPTlFFHJqa7bAzn3k2FzK8Nrdd56OWvE=", "+pNXWGIXzU1dqSy/aePHKYtSkDnFinI+/bL5HnurJ6w=", "G1R9F5B2KizFA2NyhDznjNKXPOF1SDmbGHz+omLG7TY=", "RZ7XwGuqMr5gDUh2HpUcaXp+AmIiEcPJiMui4GLJE7o=", "mBm+vQtn0C4thMnlTnxfxqowq1dXsPBCKaJ85da3JeU=", "b80/J/68RC8/tx5xRdlKk6MmTpZiVRotquWBrE/z1pI=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2461398304\nMBJWZZdAgl5oAXtH7j6jZjGCmgQgxMYKm0nN+UNCmWc=\n\n— rekor.sigstore.dev wNI9ajBFAiEAnvWMyDABFH9IZD0DKx3tw2snxsw1IhgzTWRLVNQNZg4CIDDGFSGfhl369hxfe51XGP4qD6c4Ib7X0S3cUTqhg+Hz\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiI0M2Q0NGZiZTg3YzJhOWQwZTJiYjIwZmNlYTI4NzAyNTlkZDE2YzczMTY2NmVhYWQ0NjAzZTBjMjQ2NzQ5NDVkIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJUUQzSTBPVzBPQ2hyRkVuR0RtMnA5SGhrN25mM1Y1RVJBaXBpbE9pTWtFaEhnSWdBbFhZV0NvdEVpcFFkd2dZVFBIMUJKcDR6VXpWNG9uQzFzN3lhRXRrM3lZPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRha05EUW5SRFowRjNTVUpCWjBsVlVFdGxlR3h4VTFCMFpuZE9TWFZ4VVdOVFVtUXlWV3M1U1c5TmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1RGTlZFbDVUVlJGTVZkb1kwNU5hbGwzVDBSSk1VMVVTWHBOVkVVeFYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVUwYnk4NFRrTllTM2QyVUN0blREaEZWVVJPT1hOblowWkRjVVpNUTBkaWVtTlhORUlLYnk5MlEwcFNVVWR4VjJGSmEyeFJlR1F3Y21wak1IbE1WRE5pUzJWMWNuSmhiWHBtV1V3M01rOUlLMGhOVFd4MFRXRlBRMEpsT0hkbloxaHlUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlUwWW1aMkNrZERURlJLVkRNeGVYSXJSV3AxWm5SM1EwMDVPSGhKZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVplbEpyVFVkU2FrNVVXVEpPYWtacFdWZFNhVTlYVG1wT1JHUnBDbHBIVlRKT1JHUnNUVVJvYlZwVVRtbE9SMHB0VG5wUk1VOUVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2xhUjFVeVRrUmtiRTFFYUcwS1dsUk9hVTVIU20xT2VsRXhUMFJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2tLV2tkVk1rNUVaR3hOUkdodFdsUk9hVTVIU20xT2VsRXhUMFJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZE5NRnBFUW1zS1dYcFZNazVxV1hoWmJVWnJXV3BzYWxsNlVUTlpiVkpzVG1wUk0xcFVRVFJhYlZWNldXcFNhVnBxWXpCT1ZHZDNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWtrMFRrUlJNazlFYXpSUFZGRjJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwVVZsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTTjBKSWEwRUtaSGRDTVVGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUVhbVZIY0dkQlFVRlJSQXBCUlZsM1VrRkpaMFJSWVRCUWFERkhPRTltZGpKSWNtVlhVMHQxVDFOV2FGTnhjbE5RZERaTWFsSmhNMGhUWWtWRWNHTkRTVVoxWkdsM00zWjZTRUYxQ200M2RFTk9XbWMwWTJWVmNXWk5abkp3TlhSSVVYRktiRVZKUmpaWGQxZHZUVUZ2UjBORGNVZFRUVFE1UWtGTlJFRXlaMEZOUjFWRFRVVXdRbU5hUTJjS1YydG5SekJ6ZUhWMVJrOUlNSFF6YzBaVmRsbDNNa2RHVkM5d1ZFSjZMMDV3YjJndlF6QldNMVZzU1dreFpIWnpWREpHU25kdFNuSlhaMGw0UVU4MWJ3cEJaSEJtTjBaclJUUkdibVl4VDNWS2JYSnVNVFJLU2tkRldFRndObFp3Y21aSE1YSjVjbXN6TWxNMFkxUnhSbTA0VVhwRU1tbDJOMFYzVWtwMlVUMDlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyjADAgEAMIICwQYJKoZIhvcNAQcCoIICsjCCAq4CAQMxDTALBglghkgBZQMEAgEwgbcGCyqGSIb3DQEJEAEEoIGnBIGkMIGhAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgbmwUpq5yS+9FyOoq5FWeEdmwXYnYqxjOZLCfymM3KP0CFBQEnRhqvI60FFKgzIl+uz/txpqPGA8yMDI2MDgyNTEyMjExNVowAwIBAaAypDAwLjEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MRUwEwYDVQQDEwxzaWdzdG9yZS10c2GgADGCAdwwggHYAgEBMFEwOTEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MSAwHgYDVQQDExdzaWdzdG9yZS10c2Etc2VsZnNpZ25lZAIUOhNULwyQYe68wUMvy4qOiyojiwwwCwYJYIZIAWUDBAIBoIH8MBoGCSqGSIb3DQEJAzENBgsqhkiG9w0BCRABBDAcBgkqhkiG9w0BCQUxDxcNMjYwODI1MTIyMTE1WjAvBgkqhkiG9w0BCQQxIgQgVskgydghYfTyg9W7Vmh3NCFVm/2CYRk064SEPT4vuuQwgY4GCyqGSIb3DQEJEAIvMX8wfTB7MHkEIIX5J7wHq2LKw7RDVsEO/IGyxog/2nq55thw2dE6zQW3MFUwPaQ7MDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAoGCCqGSM49BAMCBGgwZgIxAKCzAw/Bg+pX2+4Y85xqfcsyHsTgfVPwvne0VtMMXZb0Y8Fv+glXD2MNnH7MPrR8oAIxAOHVQHO/zJaR8mDMi7qSAZH3CwTv4+UEyDdiIv/Gk/AmGzd364ndeEbVJQ9MIcHFfw=="}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"Q9RPvofCqdDiuyD86ihwJZ3RbHMWZuqtRgPgwkZ0lF0="}, "signature":"MEUCIQD3I0OW0OChrFEnGDm2p9Hhk7nf3V5ERAipilOiMkEhHgIgAlXYWCotEipQdwgYTPH1BJp4zUzV4onC1s7yaEtk3yY="}}
|
|
|
|
| 1 |
+
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSjCCBtGgAwIBAgIUHsSigXA8V+9sG5OzX86VzEeerF0wCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI3MDcxNDA5WhcNMjYwODI3MDcyNDA5WjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEMhLZCSS9724Q/wP1hrYddGnFsHiNkgWOHI+fgFnFwTeCnuaVIR/vSJaC4+9M7udTneTewkUfnfMDFbkPLMC4Z6OCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUQs0U90i88+tsgNSfVZ9RebGnhJ8wHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGFkZGY0NzQxZmFmYjU1YjRjMTAxYjE1OTg1ZGZhZWVhN2M3YzJmM2EwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzMwNDUzOTU5NTIvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoEIRq+YAAAQDAEcwRQIhAOBCYzk9TKa16WftgdDTo5TIJO38ezm5XssKT02txOtBAiBA1EzzIp7Wuq4xPciL1wOwocge+Kyvkx389dXJ0X6RszAKBggqhkjOPQQDAwNnADBkAjAz49tG/pLiwDCm21g9OsHK36n74gBWNgrouMrt4K9HkNKbNWbKbRueP0UM2KZ81OICMBTfSq82r8VzSR69S2bC68Jpvs+D4nTKruV2r+mJ8oM0h3kaJvNB0QPqK/N0GKtm9g=="}, "tlogEntries":[{"logIndex":"2614047887", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787814849", "inclusionPromise":{"signedEntryTimestamp":"MEUCIQC8BmGgNUkm6T/yk1EbJ1NNQ3zrbkJFri7dw3BnlR0e4AIgW85WxEA8YWinXDzfHW25WI/bckkoLtebiAI5L5ivkFI="}, "inclusionProof":{"logIndex":"2492143625", "rootHash":"Z2V6ugp4s7v/3fyWoaVbvtWRKINfIgX1d4DwdnD88YE=", "treeSize":"2492143669", "hashes":["PsU39qYVVAlIp8wI9cxVkqzhdSwCLR+IMX1yb9ZcpJM=", "ePBoPtZHQHI6mGYjaI+Wnq06ShZj6wmXAHA9GiCrr6k=", "4yD1bGiJH93uIYVeh36jqXbecUqJvdMXrsNhhhQq4p4=", "prsuuxeDzdF0aRkSpCTbvJSJ7M1oATAN3hDlk6Okuas=", "KU/ZzA5Z52HbJmBQqm5JPA5I4EERvsuiceOXfei63RU=", "rLjjfxTiWdnJsLWof2wJoZ7bWiVDuGBe8vdD8VYk6jE=", "VFlTxnDU4W+k/WA3ARTvzCiqVt1yxQIMUyUzMCugJ2Q=", "CFNMzmJ8N+uNx37P2asMqvrW9tw4axAVHnXLRO39tFs=", "NypqAbB+mNix4l1ahGeP9fRWiQzZCm4P1664makjjq4=", "8SoaeA7IeJPoIiveQbBfHCWXSDQabIk2PW+FT49NSEQ=", "VGrFt5JvlIUQvoGEdBBxrCe3lYbu581OhgLq/Xdyat4=", "7TfDCaMWgEwSJEZd/e33b24919rF4IDJLmbqNNRNhZs=", "0HfdxDt/zGugZLIuHasrdLEW9s8OaTykSSqBEDFQXdw=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2492143669\nZ2V6ugp4s7v/3fyWoaVbvtWRKINfIgX1d4DwdnD88YE=\n\n— rekor.sigstore.dev wNI9ajBEAiAeemS0veSYb/ZhP0KXnRIewuVB0BCQwhxGI0Be5VIOLQIgWyRQoHecQBAppQqcv3qGcrdUekeQSEzRoH6kGxhTr9A=\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiJlNGIyMzhhNTAyNzg5MWU1YWM3YjRkZDRhMTQ2N2Y3MDIwMzY5Y2M1OThmNGJlZDhiOGY5YThhMzc4ZGM4YTdkIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJREd0UEw2cGttUVc2bDlYUVdTbXRmamRZTDA1RitRdW84dWk4OHQvUGNwVkFpRUFvMm1wZUJwOUhYQkdGSEtLRFRFTTBtTXJ3NG5HeDVFeEdmcFhyZnNMcWkwPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRha05EUW5SSFowRjNTVUpCWjBsVlNITlRhV2RZUVRoV0t6bHpSelZQZWxnNE5sWjZSV1ZsY2tZd2QwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1ROTlJHTjRUa1JCTlZkb1kwNU5hbGwzVDBSSk0wMUVZM2xPUkVFMVYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZOYUV4YVExTlRPVGN5TkZFdmQxQXhhSEpaWkdSSGJrWnpTR2xPYTJkWFQwaEpLMllLWjBadVJuZFVaVU51ZFdGV1NWSXZkbE5LWVVNMEt6bE5OM1ZrVkc1bFZHVjNhMVZtYm1aTlJFWmlhMUJNVFVNMFdqWlBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZSY3pCVkNqa3dhVGc0SzNSelowNVRabFphT1ZKbFlrZHVhRW80ZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVpWMUpyV21wUk0wNUVSbTFaVjFwcFRsUldhVTVIVFhoTlJFWnBDazFVVlRWUFJGWnJXbTFHYkZwWFJUTlplbVJxVFcxWmVsbFVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWxOVkZVMVQwUldhMXB0Um13S1dsZEZNMWw2WkdwTmJWbDZXVlJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWtLVFZSVk5VOUVWbXRhYlVac1dsZEZNMWw2WkdwTmJWbDZXVlJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZEdhMXBIV1RBS1RucFJlRnB0Um0xWmFsVXhXV3BTYWsxVVFYaFpha1V4VDFSbk1WcEhXbWhhVjFab1RqSk5NMWw2U20xTk1rVjNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWsxM1RrUlZlazlVVlRWT1ZFbDJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUZTVkp4SzFsQlFVRlJSQXBCUldOM1VsRkphRUZQUWtOWmVtczVWRXRoTVRaWFpuUm5aRVJVYnpWVVNVcFBNemhsZW0wMVdITnpTMVF3TW5SNFQzUkNRV2xDUVRGRmVucEpjRGRYQ25WeE5IaFFZMmxNTVhkUGQyOWpaMlVyUzNsMmEzZ3pPRGxrV0Vvd1dEWlNjM3BCUzBKblozRm9hMnBQVUZGUlJFRjNUbTVCUkVKclFXcEJlalE1ZEVjS0wzQk1hWGRFUTIweU1XYzVUM05JU3pNMmJqYzBaMEpYVG1keWIzVk5jblEwU3psSWEwNUxZazVYWWt0aVVuVmxVREJWVFRKTFdqZ3hUMGxEVFVKVVpncFRjVGd5Y2poV2VsTlNOamxUTW1KRE5qaEtjSFp6SzBRMGJsUkxjblZXTW5JcmJVbzRiMDB3YUROcllVcDJUa0l3VVZCeFN5OU9NRWRMZEcwNVp6MDlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyTADAgEAMIICwAYJKoZIhvcNAQcCoIICsTCCAq0CAQMxDTALBglghkgBZQMEAgEwgbgGCyqGSIb3DQEJEAEEoIGoBIGlMIGiAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgolMZ3liJt017qRUx/yicdjS/xSe/M030X9M90yGLfbYCFQDz6/FLZLAyE6g0wOg+ZueTePcdixgPMjAyNjA4MjcwNzE0MDlaMAMCAQGgMqQwMC4xFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEVMBMGA1UEAxMMc2lnc3RvcmUtdHNhoAAxggHaMIIB1gIBATBRMDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAsGCWCGSAFlAwQCAaCB/DAaBgkqhkiG9w0BCQMxDQYLKoZIhvcNAQkQAQQwHAYJKoZIhvcNAQkFMQ8XDTI2MDgyNzA3MTQwOVowLwYJKoZIhvcNAQkEMSIEINS5YE9VEE5i99un4RvM8i9NvLFVkkAY+OP5b0KJHScVMIGOBgsqhkiG9w0BCRACLzF/MH0wezB5BCCF+Se8B6tiysO0Q1bBDvyBssaIP9p6uebYcNnROs0FtzBVMD2kOzA5MRUwEwYDVQQKEwxzaWdzdG9yZS5kZXYxIDAeBgNVBAMTF3NpZ3N0b3JlLXRzYS1zZWxmc2lnbmVkAhQ6E1QvDJBh7rzBQy/Lio6LKiOLDDAKBggqhkjOPQQDAgRmMGQCMF4AAJ+VQLVtO4m/N+3el75uOvQH/PDfAq82TsURJkMxE4zm2zK+VXXp+jsgcJCTVQIwf1NmqhkhKJx1dwlmF8M5q8d6FwUelOy9j+0af9mFexoxdw2AJ/SmiKCi4UfhLzDB"}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"5LI4pQJ4keWse03UoUZ/cCA2nMWY9L7YuPmoo3jcin0="}, "signature":"MEUCIDGtPL6pkmQW6l9XQWSmtfjdYL05F+Quo8ui88t/PcpVAiEAo2mpeBp9HXBGFHKKDTEM0mMrw4nGx5ExGfpXrfsLqi0="}}
|
build/torch213-cxx11-cu130-x86_64-linux/_megablocks_cuda_addf474.abi3.so
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:e6cd703144e65044ed60688b698eb888911ff348ec9cc2d205565210431ef06d
|
| 3 |
+
size 14587064
|
build/torch213-cxx11-cu130-x86_64-linux/_megablocks_cuda_c4d0dc5.abi3.so
DELETED
|
@@ -1,3 +0,0 @@
|
|
| 1 |
-
version https://git-lfs.github.com/spec/v1
|
| 2 |
-
oid sha256:da57ed156333deb2511dea109ad49184f980f95a5cd13eff0e0a178a817f27ab
|
| 3 |
-
size 10113368
|
|
|
|
|
|
|
|
|
|
|
|
build/torch213-cxx11-cu130-x86_64-linux/_ops.py
CHANGED
|
@@ -1,9 +1,9 @@
|
|
| 1 |
import torch
|
| 2 |
-
from . import
|
| 3 |
-
ops = torch.ops.
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
-
return f"
|
|
|
|
| 1 |
import torch
|
| 2 |
+
from . import _megablocks_cuda_addf474
|
| 3 |
+
ops = torch.ops._megablocks_cuda_addf474
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
+
return f"_megablocks_cuda_addf474::{op_name}"
|
build/torch213-cxx11-cu130-x86_64-linux/megablocks/__init__.py
DELETED
|
@@ -1,26 +0,0 @@
|
|
| 1 |
-
import ctypes
|
| 2 |
-
import importlib.util
|
| 3 |
-
import sys
|
| 4 |
-
from pathlib import Path
|
| 5 |
-
from types import ModuleType
|
| 6 |
-
|
| 7 |
-
|
| 8 |
-
def _import_from_path(file_path: Path) -> ModuleType:
|
| 9 |
-
# We cannot use the module name as-is, after adding it to `sys.modules`,
|
| 10 |
-
# it would also be used for other imports. So, we make a module name that
|
| 11 |
-
# depends on the path for it to be unique using the hex-encoded hash of
|
| 12 |
-
# the path.
|
| 13 |
-
path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
|
| 14 |
-
module_name = path_hash
|
| 15 |
-
spec = importlib.util.spec_from_file_location(module_name, file_path)
|
| 16 |
-
if spec is None:
|
| 17 |
-
raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
|
| 18 |
-
module = importlib.util.module_from_spec(spec)
|
| 19 |
-
if module is None:
|
| 20 |
-
raise ImportError(f"Cannot load module {module_name} from spec")
|
| 21 |
-
sys.modules[module_name] = module
|
| 22 |
-
spec.loader.exec_module(module) # type: ignore
|
| 23 |
-
return module
|
| 24 |
-
|
| 25 |
-
|
| 26 |
-
globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
build/torch213-cxx11-cu130-x86_64-linux/metadata.json
CHANGED
|
@@ -1,14 +1,18 @@
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
|
|
|
| 7 |
"backend": {
|
| 8 |
"type": "cuda",
|
| 9 |
"archs": [
|
| 10 |
"10.0",
|
|
|
|
|
|
|
| 11 |
"12.0",
|
|
|
|
| 12 |
"7.5",
|
| 13 |
"8.0",
|
| 14 |
"8.6",
|
|
@@ -48,8 +52,8 @@
|
|
| 48 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 49 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 50 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 51 |
-
"
|
| 52 |
-
"_ops.py": "
|
| 53 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 54 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 55 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
@@ -61,7 +65,6 @@
|
|
| 61 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 62 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 63 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
| 64 |
-
"megablocks/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY=",
|
| 65 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 66 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 67 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
@@ -106,11 +109,11 @@
|
|
| 106 |
"provenance": {
|
| 107 |
"kernel-builder": {
|
| 108 |
"version": "0.17.0-dev0",
|
| 109 |
-
"
|
| 110 |
"dirty": false
|
| 111 |
},
|
| 112 |
"kernel": {
|
| 113 |
-
"
|
| 114 |
"dirty": false
|
| 115 |
}
|
| 116 |
}
|
|
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
+
"id": "_megablocks_cuda_addf474",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
| 7 |
+
"kernel-depends": [],
|
| 8 |
"backend": {
|
| 9 |
"type": "cuda",
|
| 10 |
"archs": [
|
| 11 |
"10.0",
|
| 12 |
+
"10.3",
|
| 13 |
+
"11.0",
|
| 14 |
"12.0",
|
| 15 |
+
"12.1+PTX",
|
| 16 |
"7.5",
|
| 17 |
"8.0",
|
| 18 |
"8.6",
|
|
|
|
| 52 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 53 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 54 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 55 |
+
"_megablocks_cuda_addf474.abi3.so": "5s1wMUTmUETtYGiLaY64iJEf80jsnMLSBVZSEEMe8G0=",
|
| 56 |
+
"_ops.py": "WXlV4XwtKQjdQ6nap3k9TalaHkK0ypeSTshOrV6Hczo=",
|
| 57 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 58 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 59 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
|
|
| 65 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 66 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 67 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
|
|
|
| 68 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 69 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 70 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
|
|
| 109 |
"provenance": {
|
| 110 |
"kernel-builder": {
|
| 111 |
"version": "0.17.0-dev0",
|
| 112 |
+
"commit": "7d2828e6592de641b44fb8eb896719101a2ab2fb",
|
| 113 |
"dirty": false
|
| 114 |
},
|
| 115 |
"kernel": {
|
| 116 |
+
"commit": "addf4741fafb55b4c101b15985dfaeea7c7c2f3a",
|
| 117 |
"dirty": false
|
| 118 |
}
|
| 119 |
}
|
build/torch213-cxx11-cu130-x86_64-linux/metadata.json.sigstore
CHANGED
|
@@ -1 +1 @@
|
|
| 1 |
-
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHTDCCBtGgAwIBAgIUQ2eE3qLcPa9m8ynfT7fYq7P3T2EwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI1MTIyMTE2WhcNMjYwODI1MTIzMTE2WjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEzgKIXs9vLnMRHRNUmzJvBwN30AXyzzEGJdjmRfvgWZkeriOjjDWIWE8uKi2fz54uwjhcGUiefG55pZYCHRBdfaOCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUrufvz8pLglIx0xd0HauzenxtyWYwHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGM0ZDBkYzU2NjYxYmFkYjljYzQ3YmRlNjQ3ZTA4ZmUzYjRiZjc0NTgwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzI4NDQ2ODk4OTQvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoDjeINwAAAQDAEcwRQIgFM+inYsHBq6IHyTDVTLkzJyzC4mWb1w/X6ZtQWl33HsCIQDe9htRenp1LMRVh9YcYOlRTJKyUbH9A3YnTWfF+h6KsDAKBggqhkjOPQQDAwNpADBmAjEAhRLMp9RgMBeoBjZufmOYUPMktpnPWSMb0VwmEBQRyR1qmqzTO/FSmIpDufLKXtnwAjEA3c59oijmhBLIApFdGgjfhwHvOCxskuKk92bCgwglCvOMZ5FZZtAltWfwGgCGKRKt"}, "tlogEntries":[{"logIndex":"2583302581", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787660476", "inclusionPromise":{"signedEntryTimestamp":"MEUCIQCu5Bw3jdm5PVhJfG1OK/dtP6rFN6XikGK929u60UAJ9QIgdxqksyWuLTskuB5Anbxzj2vv2gKM0twxNewmsHVd1As="}, "inclusionProof":{"logIndex":"2461398319", "rootHash":"tr2Lavb3x9CxvbKytRnxcBRgWo5+GXc540zdYrOiGYU=", "treeSize":"2461398320", "hashes":["5eEexfxfd2NSQvGhcfUn+u2yM9QNOQG1BpmPcpmDWcM=", "uk6tvzglAkgiSfxRJTSKG1TUXds2l1v1uMya5p14L1w=", "AWtcY2WymsIk4LueyDgSQ3fersFRSvrng4fW2anym5c=", "cAJp2HKQxitB324dS3hL9um825VCbKzKrjNKb5AEbZg=", "BQO6Mvbc6dfqTtcGGvXrh9+QB796m4uHnHENTnE52Ao=", "eUK42vZbtTAA+teDm0LcO+qURyI8IPFvT/lz1znHKQ8=", "r8ZmsIkPd8oW6c3MXr06PoZ4nJYw0Vt5ki3ZDt5yOEM=", "M0t6Pyxtbr61GCmNg60qyyc/qJAEEt98xDnkIgI5BIY=", "EY8T0lxQDrmYoWBGpKNGttaHfKmpfqKcWqSNF4XOQKE=", "T4Jz+m9vO48rOpMSsdZSW28/bgHR8prz7JeNX1s7WiQ=", "EXKrSbeG2qm5APD5QSDN66X7UHz9OSofiW9PyNU7fPM=", "cjIeeutzThxIPTlFFHJqa7bAzn3k2FzK8Nrdd56OWvE=", "+pNXWGIXzU1dqSy/aePHKYtSkDnFinI+/bL5HnurJ6w=", "G1R9F5B2KizFA2NyhDznjNKXPOF1SDmbGHz+omLG7TY=", "RZ7XwGuqMr5gDUh2HpUcaXp+AmIiEcPJiMui4GLJE7o=", "mBm+vQtn0C4thMnlTnxfxqowq1dXsPBCKaJ85da3JeU=", "b80/J/68RC8/tx5xRdlKk6MmTpZiVRotquWBrE/z1pI=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2461398320\ntr2Lavb3x9CxvbKytRnxcBRgWo5+GXc540zdYrOiGYU=\n\n— rekor.sigstore.dev wNI9ajBFAiEA4QzUPjnDD+99iJdJiXx8XFQfUXnk4l4I2AhaoqPiA/cCIHzgjuUe64ZA/r5AFnUT7BJabeD0sTDdzQRPYMPYfKEf\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiJiMDg2ZTliMTlkNjRhZjIzYzVlYWE3ODk2NjUyNzNhNGI3MTBkMzk5NzljYzNiZmVmNTRmYWI2NzQ2NGZmMDdmIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FUUNJRWt5UlA5cUVDc3dNaC9XVjN5NVVjeUJ3RzRuaU5EUVU1dHh6TGJUbjRzS0FpQnRJcWJYbG9KQzVLempOeGU5VTd0bWFpaG1lRjVjVHVNVjhYRGVRZ0pGYnc9PSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFVSRU5EUW5SSFowRjNTVUpCWjBsVlVUSmxSVE54VEdOUVlUbHRPSGx1WmxRM1psbHhOMUF6VkRKRmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1RGTlZFbDVUVlJGTWxkb1kwNU5hbGwzVDBSSk1VMVVTWHBOVkVVeVYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVY2WjB0SldITTVka3h1VFZKSVVrNVZiWHBLZGtKM1RqTXdRVmg1ZW5wRlIwcGthbTBLVW1aMloxZGFhMlZ5YVU5cWFrUlhTVmRGT0hWTGFUSm1lalUwZFhkcWFHTkhWV2xsWmtjMU5YQmFXVU5JVWtKa1ptRlBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZ5ZFdaMkNubzRjRXhuYkVsNE1IaGtNRWhoZFhwbGJuaDBlVmRaZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVplbEpyVFVkU2FrNVVXVEpPYWtacFdWZFNhVTlYVG1wT1JHUnBDbHBIVlRKT1JHUnNUVVJvYlZwVVRtbE9SMHB0VG5wUk1VOUVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2xhUjFVeVRrUmtiRTFFYUcwS1dsUk9hVTVIU20xT2VsRXhUMFJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2tLV2tkVk1rNUVaR3hOUkdodFdsUk9hVTVIU20xT2VsRXhUMFJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZE5NRnBFUW1zS1dYcFZNazVxV1hoWmJVWnJXV3BzYWxsNlVUTlpiVkpzVG1wUk0xcFVRVFJhYlZWNldXcFNhVnBxWXpCT1ZHZDNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWtrMFRrUlJNazlFYXpSUFZGRjJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUVhbVZKVG5kQlFVRlJSQXBCUldOM1VsRkpaMFpOSzJsdVdYTklRbkUyU1VoNVZFUldWRXhyZWtwNWVrTTBiVmRpTVhjdldEWmFkRkZYYkRNelNITkRTVkZFWlRsb2RGSmxibkF4Q2t4TlVsWm9PVmxqV1U5c1VsUktTM2xWWWtnNVFUTlpibFJYWmtZcmFEWkxjMFJCUzBKblozRm9hMnBQVUZGUlJFRjNUbkJCUkVKdFFXcEZRV2hTVEUwS2NEbFNaMDFDWlc5Q2FscDFabTFQV1ZWUVRXdDBjRzVRVjFOTllqQldkMjFGUWxGU2VWSXhjVzF4ZWxSUEwwWlRiVWx3UkhWbVRFdFlkRzUzUVdwRlFRb3pZelU1YjJscWJXaENURWxCY0Vaa1IyZHFabWgzU0haUFEzaHphM1ZMYXpreVlrTm5kMmRzUTNaUFRWbzFSbHBhZEVGc2RGZG1kMGRuUTBkTFVrdDBDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyTADAgEAMIICwAYJKoZIhvcNAQcCoIICsTCCAq0CAQMxDTALBglghkgBZQMEAgEwgbcGCyqGSIb3DQEJEAEEoIGnBIGkMIGhAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgoEKiCwnbszWkpyH7aZhp2zp98sIl6+SfdYuchsJCQKMCFCqAIYa0upn0e8tNjfzKIZUnsWheGA8yMDI2MDgyNTEyMjExNlowAwIBAaAypDAwLjEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MRUwEwYDVQQDEwxzaWdzdG9yZS10c2GgADGCAdswggHXAgEBMFEwOTEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MSAwHgYDVQQDExdzaWdzdG9yZS10c2Etc2VsZnNpZ25lZAIUOhNULwyQYe68wUMvy4qOiyojiwwwCwYJYIZIAWUDBAIBoIH8MBoGCSqGSIb3DQEJAzENBgsqhkiG9w0BCRABBDAcBgkqhkiG9w0BCQUxDxcNMjYwODI1MTIyMTE2WjAvBgkqhkiG9w0BCQQxIgQg+cVzcDs30QEUy6HEa5YVdCqI4/EuLnAKSu72arVC5IowgY4GCyqGSIb3DQEJEAIvMX8wfTB7MHkEIIX5J7wHq2LKw7RDVsEO/IGyxog/2nq55thw2dE6zQW3MFUwPaQ7MDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAoGCCqGSM49BAMCBGcwZQIxAOkiCa9zmmYJqPbwZuCcmeY/rDBjY9Oa+N+l89R9ieJlW7nxNr4XZBejfshbl+xPwwIweWwRFMFw1x9xL2WCxXhppsJDFlytRM6r31+UxHsRxQ4ZKsNUmlj4DAQy3/p0wJSM"}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"sIbpsZ1kryPF6qeJZlJzpLcQ05l5zDv+9U+rZ0ZP8H8="}, "signature":"MEQCIEkyRP9qECswMh/WV3y5UcyBwG4niNDQU5txzLbTn4sKAiBtIqbXloJC5KzjNxe9U7tmaihmeF5cTuMV8XDeQgJFbw=="}}
|
|
|
|
| 1 |
+
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSjCCBtGgAwIBAgIUWnYS1zF2bbuoocIpT7LpeSwcVAwwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI3MDcxNDExWhcNMjYwODI3MDcyNDExWjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEm56kStkBeKppo36Zy0/O9+ud2mot2lYBInPSztG8Swme3t7G6gR5iZFewIZUMqFpcLFEdz/a7FgHQa8BPkBFHqOCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUt64F/s2UHJxlGH/TF4auwMvkOb4wHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGFkZGY0NzQxZmFmYjU1YjRjMTAxYjE1OTg1ZGZhZWVhN2M3YzJmM2EwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzMwNDUzOTU5NTIvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoEIRsr4AAAQDAEcwRQIgBd2usNrf/RatExbhKbwl4cPi1tkR3u6HuHBy9dZk6bICIQC7GnHowpAicnoI10ohWCVODenPMLHYwdMV3BghpxfueDAKBggqhkjOPQQDAwNnADBkAjAdeVmivmvK5SCERzLssJT2uc6jn3RMYd5igYOR+G7endAevyFfek2WwwQD1pzJ6+ICMF8fpahkZBn/+gFkiC7qTuSLnF0gGuhDAV6oLkUBqxQNAKbaphG8VZr2mLyxkbkSCA=="}, "tlogEntries":[{"logIndex":"2614048283", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787814851", "inclusionPromise":{"signedEntryTimestamp":"MEUCIEn0dbAX2fh5XVkw9oAnpFNO3GnfeaICyj8JsmYQEdmEAiEAxTfJk5S5kdZXYm09zgjGvKPo6FZQOFyZL2lfd3LndLk="}, "inclusionProof":{"logIndex":"2492144021", "rootHash":"4HEIrictycn6kLT2dBzTnuExmbY1AvWX7mu2EWSzBPY=", "treeSize":"2492144050", "hashes":["03bt7RcBIn3/6pISiUTvoNE6lbQ2UfKmbt2VDhsz1YY=", "YRkfubZKGlRqrQ82R/CF6pMCvRRXVHo/0zYh1OSyTSA=", "cnAywHbi1+mwmV5RxPg+Pfv3PX1BvJyZY0Hei/vdWJw=", "UdCfK0MFItMl/zmZGwFNehWO8i44SumIVuCDPjXT6f8=", "R0VuQNDYNOJty64v4++YKvtG8or5QMzi+v0iWc4y0j4=", "BTWs7zmg6x5lZcg8LJf5XGk6absoEVvidn27LXPsi7c=", "vw4cIKgV5P3+Cz+P5vqtQZc0JkF6TtbPAAowXhOW8KE=", "IlLZ05IZbs5lgJDWCYLdKfpUKQLmCeTy+T6G3jHt+Yc=", "VFlTxnDU4W+k/WA3ARTvzCiqVt1yxQIMUyUzMCugJ2Q=", "CFNMzmJ8N+uNx37P2asMqvrW9tw4axAVHnXLRO39tFs=", "NypqAbB+mNix4l1ahGeP9fRWiQzZCm4P1664makjjq4=", "8SoaeA7IeJPoIiveQbBfHCWXSDQabIk2PW+FT49NSEQ=", "VGrFt5JvlIUQvoGEdBBxrCe3lYbu581OhgLq/Xdyat4=", "7TfDCaMWgEwSJEZd/e33b24919rF4IDJLmbqNNRNhZs=", "0HfdxDt/zGugZLIuHasrdLEW9s8OaTykSSqBEDFQXdw=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2492144050\n4HEIrictycn6kLT2dBzTnuExmbY1AvWX7mu2EWSzBPY=\n\n— rekor.sigstore.dev wNI9ajBFAiEA7Xwme3VhZr+aF8t+9Ei+tWLArAAtQnkh0/s7kf+/T8kCIE8UbaWMTzAEqUbU5E5IYN4jGPBCFNQ9nRpO2p3ZV0Nl\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiJmNjEwZGM5YWU2YThmZGMzODRlZTRjZDcwYTlhOThmYjcyYjFkYThjMmJiNjI4Yjc0NGJiYzAzZjg3MmY1OGNlIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJR1BxT1BCeGV6M1pteU5YWkRqN0FSMklPcVBsbHJsVDRndGJaNnpSUDFtL0FpRUE3bXc4WlRTd2tEYUNNRkJML3NXM2lMZ3NNMzF6Q2U2SVdKOVlkRmduZHBBPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRha05EUW5SSFowRjNTVUpCWjBsVlYyNVpVekY2UmpKaVluVnZiMk5KY0ZRM1RIQmxVM2RqVmtGM2QwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1ROTlJHTjRUa1JGZUZkb1kwNU5hbGwzVDBSSk0wMUVZM2xPUkVWNFYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZ0TlRaclUzUnJRbVZMY0hCdk16WmFlVEF2VHprcmRXUXliVzkwTW14WlFrbHVVRk1LZW5SSE9GTjNiV1V6ZERkSE5tZFNOV2xhUm1WM1NWcFZUWEZHY0dOTVJrVmtlaTloTjBablNGRmhPRUpRYTBKR1NIRlBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlYwTmpSR0NpOXpNbFZJU25oc1IwZ3ZWRVkwWVhWM1RYWnJUMkkwZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVpWMUpyV21wUk0wNUVSbTFaVjFwcFRsUldhVTVIVFhoTlJFWnBDazFVVlRWUFJGWnJXbTFHYkZwWFJUTlplbVJxVFcxWmVsbFVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWxOVkZVMVQwUldhMXB0Um13S1dsZEZNMWw2WkdwTmJWbDZXVlJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWtLVFZSVk5VOUVWbXRhYlVac1dsZEZNMWw2WkdwTmJWbDZXVlJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZEdhMXBIV1RBS1RucFJlRnB0Um0xWmFsVXhXV3BTYWsxVVFYaFpha1V4VDFSbk1WcEhXbWhhVjFab1RqSk5NMWw2U20xTk1rVjNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWsxM1RrUlZlazlVVlRWT1ZFbDJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUZTVkp6Y2pSQlFVRlJSQXBCUldOM1VsRkpaMEprTW5WelRuSm1MMUpoZEVWNFltaExZbmRzTkdOUWFURjBhMUl6ZFRaSWRVaENlVGxrV21zMllrbERTVkZETjBkdVNHOTNjRUZwQ21OdWIwa3hNRzlvVjBOV1QwUmxibEJOVEVoWmQyUk5Wak5DWjJod2VHWjFaVVJCUzBKblozRm9hMnBQVUZGUlJFRjNUbTVCUkVKclFXcEJaR1ZXYldrS2RtMTJTelZUUTBWU2VreHpjMHBVTW5Wak5tcHVNMUpOV1dRMWFXZFpUMUlyUnpkbGJtUkJaWFo1Um1abGF6SlhkM2RSUkRGd2VrbzJLMGxEVFVZNFpncHdZV2hyV2tKdUx5dG5SbXRwUXpkeFZIVlRURzVHTUdkSGRXaEVRVlkyYjB4clZVSnhlRkZPUVV0aVlYQm9SemhXV25JeWJVeDVlR3RpYTFORFFUMDlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyzADAgEAMIICwgYJKoZIhvcNAQcCoIICszCCAq8CAQMxDTALBglghkgBZQMEAgEwgbgGCyqGSIb3DQEJEAEEoIGoBIGlMIGiAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgdIBY9FmMYs0v5LBhLCqPnft0U+RHZP0V8fhzOV1lPBgCFQCpRsmhsENgA9qMO6Zh2nQidUaKlhgPMjAyNjA4MjcwNzE0MTFaMAMCAQGgMqQwMC4xFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEVMBMGA1UEAxMMc2lnc3RvcmUtdHNhoAAxggHcMIIB2AIBATBRMDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAsGCWCGSAFlAwQCAaCB/DAaBgkqhkiG9w0BCQMxDQYLKoZIhvcNAQkQAQQwHAYJKoZIhvcNAQkFMQ8XDTI2MDgyNzA3MTQxMVowLwYJKoZIhvcNAQkEMSIEIMlODRHd427Crp8kSiHfaVC73mmpMRNr82e0Qj+M5l5TMIGOBgsqhkiG9w0BCRACLzF/MH0wezB5BCCF+Se8B6tiysO0Q1bBDvyBssaIP9p6uebYcNnROs0FtzBVMD2kOzA5MRUwEwYDVQQKEwxzaWdzdG9yZS5kZXYxIDAeBgNVBAMTF3NpZ3N0b3JlLXRzYS1zZWxmc2lnbmVkAhQ6E1QvDJBh7rzBQy/Lio6LKiOLDDAKBggqhkjOPQQDAgRoMGYCMQDKrcjprxpxpA6mZBufEqfOX+aROCy11HnmvAxiTTWsRa1zXfUB8BgdwGlSzoQiQKsCMQDJ8/8g9terKbehxFPdwH9YgJNyUFiTXAjC89qXN3sESjC7WKLnMocGPbtmTrxV4TA="}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"9hDcmuao/cOE7kzXCpqY+3Kx2owrtii3RLvAP4cvWM4="}, "signature":"MEUCIGPqOPBxez3ZmyNXZDj7AR2IOqPllrlT4gtbZ6zRP1m/AiEA7mw8ZTSwkDaCMFBL/sW3iLgsM31zCe6IWJ9YdFgndpA="}}
|
build/torch213-cxx11-cu132-x86_64-linux/_megablocks_cuda_addf474.abi3.so
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:6c590af63c62c28e8ae434d47b9ab628a6e96deeb87b3c1280753fab22101ef4
|
| 3 |
+
size 14987032
|
build/torch213-cxx11-cu132-x86_64-linux/_megablocks_cuda_c4d0dc5.abi3.so
DELETED
|
@@ -1,3 +0,0 @@
|
|
| 1 |
-
version https://git-lfs.github.com/spec/v1
|
| 2 |
-
oid sha256:884567461acf708fef2b49b8ab9efcf462d0cb8b0751e808386d69b925779146
|
| 3 |
-
size 10402704
|
|
|
|
|
|
|
|
|
|
|
|
build/torch213-cxx11-cu132-x86_64-linux/_ops.py
CHANGED
|
@@ -1,9 +1,9 @@
|
|
| 1 |
import torch
|
| 2 |
-
from . import
|
| 3 |
-
ops = torch.ops.
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
-
return f"
|
|
|
|
| 1 |
import torch
|
| 2 |
+
from . import _megablocks_cuda_addf474
|
| 3 |
+
ops = torch.ops._megablocks_cuda_addf474
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
+
return f"_megablocks_cuda_addf474::{op_name}"
|
build/torch213-cxx11-cu132-x86_64-linux/megablocks/__init__.py
DELETED
|
@@ -1,26 +0,0 @@
|
|
| 1 |
-
import ctypes
|
| 2 |
-
import importlib.util
|
| 3 |
-
import sys
|
| 4 |
-
from pathlib import Path
|
| 5 |
-
from types import ModuleType
|
| 6 |
-
|
| 7 |
-
|
| 8 |
-
def _import_from_path(file_path: Path) -> ModuleType:
|
| 9 |
-
# We cannot use the module name as-is, after adding it to `sys.modules`,
|
| 10 |
-
# it would also be used for other imports. So, we make a module name that
|
| 11 |
-
# depends on the path for it to be unique using the hex-encoded hash of
|
| 12 |
-
# the path.
|
| 13 |
-
path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
|
| 14 |
-
module_name = path_hash
|
| 15 |
-
spec = importlib.util.spec_from_file_location(module_name, file_path)
|
| 16 |
-
if spec is None:
|
| 17 |
-
raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
|
| 18 |
-
module = importlib.util.module_from_spec(spec)
|
| 19 |
-
if module is None:
|
| 20 |
-
raise ImportError(f"Cannot load module {module_name} from spec")
|
| 21 |
-
sys.modules[module_name] = module
|
| 22 |
-
spec.loader.exec_module(module) # type: ignore
|
| 23 |
-
return module
|
| 24 |
-
|
| 25 |
-
|
| 26 |
-
globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
build/torch213-cxx11-cu132-x86_64-linux/metadata.json
CHANGED
|
@@ -1,14 +1,18 @@
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
|
|
|
| 7 |
"backend": {
|
| 8 |
"type": "cuda",
|
| 9 |
"archs": [
|
| 10 |
"10.0",
|
|
|
|
|
|
|
| 11 |
"12.0",
|
|
|
|
| 12 |
"7.5",
|
| 13 |
"8.0",
|
| 14 |
"8.6",
|
|
@@ -48,8 +52,8 @@
|
|
| 48 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 49 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 50 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 51 |
-
"
|
| 52 |
-
"_ops.py": "
|
| 53 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 54 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 55 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
@@ -61,7 +65,6 @@
|
|
| 61 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 62 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 63 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
| 64 |
-
"megablocks/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY=",
|
| 65 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 66 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 67 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
@@ -106,11 +109,11 @@
|
|
| 106 |
"provenance": {
|
| 107 |
"kernel-builder": {
|
| 108 |
"version": "0.17.0-dev0",
|
| 109 |
-
"
|
| 110 |
"dirty": false
|
| 111 |
},
|
| 112 |
"kernel": {
|
| 113 |
-
"
|
| 114 |
"dirty": false
|
| 115 |
}
|
| 116 |
}
|
|
|
|
| 1 |
{
|
| 2 |
"name": "megablocks",
|
| 3 |
+
"id": "_megablocks_cuda_addf474",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
| 7 |
+
"kernel-depends": [],
|
| 8 |
"backend": {
|
| 9 |
"type": "cuda",
|
| 10 |
"archs": [
|
| 11 |
"10.0",
|
| 12 |
+
"10.3",
|
| 13 |
+
"11.0",
|
| 14 |
"12.0",
|
| 15 |
+
"12.1+PTX",
|
| 16 |
"7.5",
|
| 17 |
"8.0",
|
| 18 |
"8.6",
|
|
|
|
| 52 |
"_layers/mpu.py": "BX61/pmqRCVeEMnzpuK8x5NpITBJyZfOawGnDegW1O8=",
|
| 53 |
"_layers/router.py": "GKC/5//Uf43HT6zXlBTStCMTml2YXn7Jbgq9EgTgDC0=",
|
| 54 |
"_layers/sharedexpert_registry.py": "GfwHKjbC7lBXYq9kNifp3HudTdZKSqwsFxGxFLc2krw=",
|
| 55 |
+
"_megablocks_cuda_addf474.abi3.so": "bFkK9jxiwo6K5DTUe5q2KKbpbe64ezwSgHU/qyIQHvQ=",
|
| 56 |
+
"_ops.py": "WXlV4XwtKQjdQ6nap3k9TalaHkK0ypeSTshOrV6Hczo=",
|
| 57 |
"_version.py": "W3l9oUrnmvfBoIy06Lv0y2ezTQhNU04ug/txHcdWJD0=",
|
| 58 |
"backend/__init__.py": "4bunAqqjwE93nym8CigHbqVAal7UaYGdONb8lGRfrek=",
|
| 59 |
"backend/kernels.py": "2UFNTlmSDk3S5y6UW14MJMPVVsVhmozxQ5wFDxtZHhQ=",
|
|
|
|
| 65 |
"grouped_gemm/ops.py": "Kj7yRgv7afB5SVwXIkN9RNIF2kxD2qrZ0dVsKX5ph2s=",
|
| 66 |
"grouped_gemm_util.py": "bVPrdsEtBVJJ+v/bNrGNf/QHNlVL4VxHg92HvaL3oq4=",
|
| 67 |
"layers.py": "mNTlX2oKJyHBJllRd24lPNdGLa2H9cEUgVLvkZJpoNo=",
|
|
|
|
| 68 |
"ops/__init__.py": "MHQsiMEgkKeZCa8TAYFMGDFnwFIgijfAN7DhMUt/48U=",
|
| 69 |
"ops/all_to_all_benchmark.py": "sxXSCSDdsh+62IWpKNTyBphLTngY1jlgJ1WgNooSEO8=",
|
| 70 |
"ops/binned_gather.py": "pkWne+ThfiTR/bMbDs2QDimKxAI4klqD1FdvhytE7lU=",
|
|
|
|
| 109 |
"provenance": {
|
| 110 |
"kernel-builder": {
|
| 111 |
"version": "0.17.0-dev0",
|
| 112 |
+
"commit": "7d2828e6592de641b44fb8eb896719101a2ab2fb",
|
| 113 |
"dirty": false
|
| 114 |
},
|
| 115 |
"kernel": {
|
| 116 |
+
"commit": "addf4741fafb55b4c101b15985dfaeea7c7c2f3a",
|
| 117 |
"dirty": false
|
| 118 |
}
|
| 119 |
}
|
build/torch213-cxx11-cu132-x86_64-linux/metadata.json.sigstore
CHANGED
|
@@ -1 +1 @@
|
|
| 1 |
-
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSjCCBtGgAwIBAgIUF5DyOMVZVai+Z+yjWIintpmoRvkwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI1MTIyMTE3WhcNMjYwODI1MTIzMTE3WjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEkHVEgyv/5Z2hPIosGrIjxsDY2SgPvIR3I7ryRAoqPOHxeVnwdMijkh4IOHrrMUTIey5ZtlypNeaGwP9Lop9ds6OCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUc/f8L/9W3JqtmFkrNeFy34wjvBMwHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYzRkMGRjNTY2NjFiYWRiOWNjNDdiZGU2NDdlMDhmZTNiNGJmNzQ1ODAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGM0ZDBkYzU2NjYxYmFkYjljYzQ3YmRlNjQ3ZTA4ZmUzYjRiZjc0NTgwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzI4NDQ2ODk4OTQvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoDjeJfYAAAQDAEcwRQIhANvjiRk8H2pUQ1yRN0WIED12UWYHAXymSWQdCA3y+ws/AiAX8lgfooBaUxCltPodrVvBtRYSjEMKn8BdnyXCq4HZWDAKBggqhkjOPQQDAwNnADBkAjA1Tsb8pF1dsG2SidoVomoiipbkqA2Df7LK2NSud1N1i5LvM9Hc+mKpXZfBrTka8+ICMCLzQER1kuNIVc4SiD3L2y1+zW7b2e7oyjbtyZd5eEMjAp8g6HN7H+5QtNu2LiigpQ=="}, "tlogEntries":[{"logIndex":"2583302595", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787660478", "inclusionPromise":{"signedEntryTimestamp":"MEQCID+pplq9/9Kem/px3K+yzWfxjNldVc1fXu6ibh6ufELtAiB7aSn0+iqpy8/wqDa8z+6Ke93A5O9moJSzqkNlZmffLg=="}, "inclusionProof":{"logIndex":"2461398333", "rootHash":"EY2FE8ANXSU2qIj9h4gwrGVFAFad8BQ7eG4+YYD8o5g=", "treeSize":"2461398335", "hashes":["NvMTyutpVc0DFj+m3vEFgR3arquLsDAvYiBtzLizk3w=", "9Rrkg3PlpCmq8cQdSinyYoeLRN0QHQUiXh2sYc8v58k=", "sQkWK7cKmGuj3OiwNkTN/wftR/IJgvbixs2HoJRzOSE=", "gvW/0ym4UwAFu3jSc7Cs5+ZNTAVvHLe1MzxNHz+npCQ=", "tWauNUno2kG7EKvel+T+4ks6XYUxIgWBCrom3zwUgLg=", "BQO6Mvbc6dfqTtcGGvXrh9+QB796m4uHnHENTnE52Ao=", "eUK42vZbtTAA+teDm0LcO+qURyI8IPFvT/lz1znHKQ8=", "r8ZmsIkPd8oW6c3MXr06PoZ4nJYw0Vt5ki3ZDt5yOEM=", "M0t6Pyxtbr61GCmNg60qyyc/qJAEEt98xDnkIgI5BIY=", "EY8T0lxQDrmYoWBGpKNGttaHfKmpfqKcWqSNF4XOQKE=", "T4Jz+m9vO48rOpMSsdZSW28/bgHR8prz7JeNX1s7WiQ=", "EXKrSbeG2qm5APD5QSDN66X7UHz9OSofiW9PyNU7fPM=", "cjIeeutzThxIPTlFFHJqa7bAzn3k2FzK8Nrdd56OWvE=", "+pNXWGIXzU1dqSy/aePHKYtSkDnFinI+/bL5HnurJ6w=", "G1R9F5B2KizFA2NyhDznjNKXPOF1SDmbGHz+omLG7TY=", "RZ7XwGuqMr5gDUh2HpUcaXp+AmIiEcPJiMui4GLJE7o=", "mBm+vQtn0C4thMnlTnxfxqowq1dXsPBCKaJ85da3JeU=", "b80/J/68RC8/tx5xRdlKk6MmTpZiVRotquWBrE/z1pI=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2461398335\nEY2FE8ANXSU2qIj9h4gwrGVFAFad8BQ7eG4+YYD8o5g=\n\n— rekor.sigstore.dev wNI9ajBFAiAa42WA5Z85UvzhojuwnU3XSYXiO3MWSjARjVSwxiyodgIhAK6zi0m6uDXzhG+B0OnBNXYG0V85ll9nsO9fpzHotkvL\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiI1MTEzYWE1M2Y5ZmZjZmNlOTBiMjQxNzc1YmMyZTk0NzhlNjgzOTNkZTI1MWM0M2E5YjBkZjYzZmY4NDY4ZTBiIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJUUNFeGZJOE1oVVpVUVAyNm1SUFBHck1jdktTYWJkei9OOTY2UzJHMTB3MHJ3SWdNWnhyRkRDV3JIK3dYQzFSWE9xRmJ4NHlMVXRPNTlmOU1oelkrVXBBck1FPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRha05EUW5SSFowRjNTVUpCWjBsVlJqVkVlVTlOVmxwV1lXa3JXaXQ1YWxkSmFXNTBjRzF2VW5acmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1RGTlZFbDVUVlJGTTFkb1kwNU5hbGwzVDBSSk1VMVVTWHBOVkVVelYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZyU0ZaRlozbDJMelZhTW1oUVNXOXpSM0pKYW5oelJGa3lVMmRRZGtsU00wazNjbmtLVWtGdmNWQlBTSGhsVm01M1pFMXBhbXRvTkVsUFNISnlUVlZVU1dWNU5WcDBiSGx3VG1WaFIzZFFPVXh2Y0Rsa2N6WlBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZqTDJZNENrd3ZPVmN6U25GMGJVWnJjazVsUm5rek5IZHFka0pOZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVplbEpyVFVkU2FrNVVXVEpPYWtacFdWZFNhVTlYVG1wT1JHUnBDbHBIVlRKT1JHUnNUVVJvYlZwVVRtbE9SMHB0VG5wUk1VOUVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2xhUjFVeVRrUmtiRTFFYUcwS1dsUk9hVTVIU20xT2VsRXhUMFJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWw2VW10TlIxSnFUbFJaTWs1cVJtbFpWMUpwVDFkT2FrNUVaR2tLV2tkVk1rNUVaR3hOUkdodFdsUk9hVTVIU20xT2VsRXhUMFJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZE5NRnBFUW1zS1dYcFZNazVxV1hoWmJVWnJXV3BzYWxsNlVUTlpiVkpzVG1wUk0xcFVRVFJhYlZWNldXcFNhVnBxWXpCT1ZHZDNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWtrMFRrUlJNazlFYXpSUFZGRjJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUVhbVZLWmxsQlFVRlJSQXBCUldOM1VsRkphRUZPZG1wcFVtczRTREp3VlZFeGVWSk9NRmRKUlVReE1sVlhXVWhCV0hsdFUxZFJaRU5CTTNrcmQzTXZRV2xCV0Roc1oyWnZiMEpoQ2xWNFEyeDBVRzlrY2xaMlFuUlNXVk5xUlUxTGJqaENaRzU1V0VOeE5FaGFWMFJCUzBKblozRm9hMnBQVUZGUlJFRjNUbTVCUkVKclFXcEJNVlJ6WWpnS2NFWXhaSE5ITWxOcFpHOVdiMjF2YVdsd1ltdHhRVEpFWmpkTVN6Sk9VM1ZrTVU0eGFUVk1kazA1U0dNcmJVdHdXRnBtUW5KVWEyRTRLMGxEVFVOTWVncFJSVkl4YTNWT1NWWmpORk5wUkROTU1ua3hLM3BYTjJJeVpUZHZlV3BpZEhsYVpEVmxSVTFxUVhBNFp6WklUamRJS3pWUmRFNTFNa3hwYVdkd1VUMDlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyzADAgEAMIICwgYJKoZIhvcNAQcCoIICszCCAq8CAQMxDTALBglghkgBZQMEAgEwgbgGCyqGSIb3DQEJEAEEoIGoBIGlMIGiAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgxeIVPKNrrYtXMZsmfNtCwnPOGdfR9NvP7Njnw/SqP6MCFQDNfmKF4pJ1EFJSHeo/A1ts0Ts+lhgPMjAyNjA4MjUxMjIxMThaMAMCAQGgMqQwMC4xFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEVMBMGA1UEAxMMc2lnc3RvcmUtdHNhoAAxggHcMIIB2AIBATBRMDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAsGCWCGSAFlAwQCAaCB/DAaBgkqhkiG9w0BCQMxDQYLKoZIhvcNAQkQAQQwHAYJKoZIhvcNAQkFMQ8XDTI2MDgyNTEyMjExOFowLwYJKoZIhvcNAQkEMSIEIGzt5sr2xY/srtfvIQZ81xvXl5YTCSlFvVU+o9lpfx7gMIGOBgsqhkiG9w0BCRACLzF/MH0wezB5BCCF+Se8B6tiysO0Q1bBDvyBssaIP9p6uebYcNnROs0FtzBVMD2kOzA5MRUwEwYDVQQKEwxzaWdzdG9yZS5kZXYxIDAeBgNVBAMTF3NpZ3N0b3JlLXRzYS1zZWxmc2lnbmVkAhQ6E1QvDJBh7rzBQy/Lio6LKiOLDDAKBggqhkjOPQQDAgRoMGYCMQCevyIlSlcjbg7xk5rgexkfxPSSyTmJ3wrlT+JZNwtNaFBSBXkykScNlp6s5eUM3O4CMQDORxw7I/QNIH+FgoN30YR6Hyj4k7S5eGdsAcYa5aO3fPRRgfVoDMIPWbAjr4GWaI4="}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"UROqU/n/z86QskF3W8LpR45oOT3iUcQ6mw32P/hGjgs="}, "signature":"MEUCIQCExfI8MhUZUQP26mRPPGrMcvKSabdz/N966S2G10w0rwIgMZxrFDCWrH+wXC1RXOqFbx4yLUtO59f9MhzY+UpArME="}}
|
|
|
|
| 1 |
+
{"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSjCCBtGgAwIBAgIUa1rY5fShGcQfedCk3dbQrDcT6powCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwODI3MDcxNDEzWhcNMjYwODI3MDcyNDEzWjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEng6lK/BeN0Awm0cVKGhiN34YyvwCdQlrFCMUiSm4XwCyjcntR4ZG/r36WpT+YgIt32nKbmIfahanBg7uJrM15KOCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUdekMREaT+z0MtF46KfKv7+4mviswHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoYWRkZjQ3NDFmYWZiNTViNGMxMDFiMTU5ODVkZmFlZWE3YzdjMmYzYTAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKGFkZGY0NzQxZmFmYjU1YjRjMTAxYjE1OTg1ZGZhZWVhN2M3YzJmM2EwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMzMwNDUzOTU5NTIvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABoEIRue0AAAQDAEcwRQIhAIqqj4AYFmsf/8LDKF4+6xLDIuF68PEQtTHFDvkGK7i+AiBmobPTu8utL4OC2WTCHCTvZikUlROtWFhagvnFeMrWuzAKBggqhkjOPQQDAwNnADBkAjAmLpAKlJyTYX6DA8NaipyZo8i8e1WN151ZU9n7bga9NWLTTF5W3fQMq75QU4aVXnsCMAWo1SiNV4VwLzu0OxxrRE7pZhXOYRupc7d9ahrZhStaU6ooj5ljFtBVOQ3Dcc/XoA=="}, "tlogEntries":[{"logIndex":"2614048644", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1787814853", "inclusionPromise":{"signedEntryTimestamp":"MEUCIHeyRrGFYMLsZ90B4SZ1SP2fZIk5Cri+iPTdCK8wu3syAiEAzGhDlpxrQGZH/4nO/q3/GqXKsAa51gs3/0ahYNPQYRY="}, "inclusionProof":{"logIndex":"2492144382", "rootHash":"515/sFUGPvAeLi6ZM+9gYRAEMYvbTqAOGMePSp4q638=", "treeSize":"2492144419", "hashes":["SZ1QcuHxZpQv12e6F9U5CdtheLQTUnQIZyZ20RW65Ts=", "ClF8WxFR/BX5FNWojg2xA+RSMLCsfFQQ9l0Bz0+UgBI=", "vKY14mV1p23Vw0FaAXydfZlP6nyOMx4zyIi92OOUJAY=", "RzUpHZ++ehuR0OeOKXhJd+5ReeSc7pEJenUy8m/gNQI=", "E0xAnhM1lhenY84RQEwRv6usiyPDP07f2cag248BYSs=", "1mpQ4BfD6S1Tnix5+RB2ABkEzwcXNQmWDrc00zgqZ7g=", "ZAEPc4SjH4mAouX2IuTCEkLhyPlBL0ZdQEHa5DWwDk8=", "/86f/+8yplBih+rHWRuNE/xNx7zHslKJzBzCQZwKfno=", "q0gGRQ2vDBFvDwUIaO6HzxFrMWyjCLgxprzhA3BxcNU=", "gxi1eJ8qX1K1iauxvafmRFx1FhOcSvz3IdnMqR1qs9E=", "VFlTxnDU4W+k/WA3ARTvzCiqVt1yxQIMUyUzMCugJ2Q=", "CFNMzmJ8N+uNx37P2asMqvrW9tw4axAVHnXLRO39tFs=", "NypqAbB+mNix4l1ahGeP9fRWiQzZCm4P1664makjjq4=", "8SoaeA7IeJPoIiveQbBfHCWXSDQabIk2PW+FT49NSEQ=", "VGrFt5JvlIUQvoGEdBBxrCe3lYbu581OhgLq/Xdyat4=", "7TfDCaMWgEwSJEZd/e33b24919rF4IDJLmbqNNRNhZs=", "0HfdxDt/zGugZLIuHasrdLEW9s8OaTykSSqBEDFQXdw=", "SndbMKVtcTenAkwi2JBfGzD+mhexp1qJbRIY+A1JRIU=", "xH/DCseLHr9eKoYT8qsORZK7zVdEGYWHuVtsVrD95wY="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2492144419\n515/sFUGPvAeLi6ZM+9gYRAEMYvbTqAOGMePSp4q638=\n\n— rekor.sigstore.dev wNI9ajBGAiEA5nM5F5tXv13DDJ6sw6nLDAGUPHNRQ+782WwtX4qJD8gCIQDv3EYlHEWkD+77fTXLyPUEPWkBsgD1PSmFIZT61MqmTA==\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiI0YzJlNzJlOTU3NjRjNjMxYjk3YThhYTRjN2VlNThlNTU3NWE1NDJkZDhhMjY5MDUwMzE4YmY3ZTgyMTFlNzEzIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FUUNJRHpEbW5aN1dYZ3E5bmhOTFBYeGV0NGptZW0vWHhLOVhrRFRoMFhheEJJaEFpQTdxU2t3ZFhSUUROV09IU1BiN1RCaTRLVExCZXNuakVzU09XQ1BRTnkraWc9PSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRha05EUW5SSFowRjNTVUpCWjBsVllURnlXVFZtVTJoSFkxRm1aV1JEYXpOa1lsRnlSR05VTm5CdmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDlFU1ROTlJHTjRUa1JGZWxkb1kwNU5hbGwzVDBSSk0wMUVZM2xPUkVWNlYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZ1Wnpac1N5OUNaVTR3UVhkdE1HTldTMGRvYVU0ek5GbDVkbmREWkZGc2NrWkRUVlVLYVZOdE5GaDNRM2xxWTI1MFVqUmFSeTl5TXpaWGNGUXJXV2RKZERNeWJrdGliVWxtWVdoaGJrSm5OM1ZLY2sweE5VdFBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZrWld0TkNsSkZZVlFyZWpCTmRFWTBOa3RtUzNZM0t6UnRkbWx6ZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVpWMUpyV21wUk0wNUVSbTFaVjFwcFRsUldhVTVIVFhoTlJFWnBDazFVVlRWUFJGWnJXbTFHYkZwWFJUTlplbVJxVFcxWmVsbFVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWxOVkZVMVQwUldhMXB0Um13S1dsZEZNMWw2WkdwTmJWbDZXVlJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMWxYVW10YWFsRXpUa1JHYlZsWFdtbE9WRlpwVGtkTmVFMUVSbWtLVFZSVk5VOUVWbXRhYlVac1dsZEZNMWw2WkdwTmJWbDZXVlJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwZEdhMXBIV1RBS1RucFJlRnB0Um0xWmFsVXhXV3BTYWsxVVFYaFpha1V4VDFSbk1WcEhXbWhhVjFab1RqSk5NMWw2U20xTk1rVjNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOZWsxM1RrUlZlazlVVlRWT1ZFbDJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtOUZTVkoxWlRCQlFVRlJSQXBCUldOM1VsRkphRUZKY1hGcU5FRlpSbTF6Wmk4NFRFUkxSalFyTm5oTVJFbDFSalk0VUVWUmRGUklSa1IyYTBkTE4ya3JRV2xDYlc5aVVGUjFPSFYwQ2t3MFQwTXlWMVJEU0VOVWRscHBhMVZzVWs5MFYwWm9ZV2QyYmtabFRYSlhkWHBCUzBKblozRm9hMnBQVUZGUlJFRjNUbTVCUkVKclFXcEJiVXh3UVVzS2JFcDVWRmxZTmtSQk9FNWhhWEI1V204NGFUaGxNVmRPTVRVeFdsVTViamRpWjJFNVRsZE1WRlJHTlZjelpsRk5jVGMxVVZVMFlWWllibk5EVFVGWGJ3b3hVMmxPVmpSV2QweDZkVEJQZUhoeVVrVTNjRnBvV0U5WlVuVndZemRrT1dGb2NscG9VM1JoVlRadmIybzFiR3BHZEVKV1QxRXpSR05qTDFodlFUMDlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyTADAgEAMIICwAYJKoZIhvcNAQcCoIICsTCCAq0CAQMxDTALBglghkgBZQMEAgEwgbgGCyqGSIb3DQEJEAEEoIGoBIGlMIGiAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgvmVnRfBJfjpCzbvskgeP5s87wALCyJooza5/zgrFVjECFQCEXItenUdwVo50Wjw+xiMatUl/ExgPMjAyNjA4MjcwNzE0MTNaMAMCAQGgMqQwMC4xFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEVMBMGA1UEAxMMc2lnc3RvcmUtdHNhoAAxggHaMIIB1gIBATBRMDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAsGCWCGSAFlAwQCAaCB/DAaBgkqhkiG9w0BCQMxDQYLKoZIhvcNAQkQAQQwHAYJKoZIhvcNAQkFMQ8XDTI2MDgyNzA3MTQxM1owLwYJKoZIhvcNAQkEMSIEICx6Sjp1NQ6cpS0ZI0oxnWT7hPyku6xwYf86o+4xGnWmMIGOBgsqhkiG9w0BCRACLzF/MH0wezB5BCCF+Se8B6tiysO0Q1bBDvyBssaIP9p6uebYcNnROs0FtzBVMD2kOzA5MRUwEwYDVQQKEwxzaWdzdG9yZS5kZXYxIDAeBgNVBAMTF3NpZ3N0b3JlLXRzYS1zZWxmc2lnbmVkAhQ6E1QvDJBh7rzBQy/Lio6LKiOLDDAKBggqhkjOPQQDAgRmMGQCMETS2x6QVopD2v0Vi6wlWDvPWquJjEzn/khJSRm20btt/mRkhksKm0jiMtqG+Vf/2wIwPq8H1AGa3vIXqYnaP2YWfnNkaU8/y9jT9OVQLP/nNq3Xd9X0HC/Oh/vcLX13YjrB"}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"TC5y6VdkxjG5eoqkx+5Y5VdaVC3YomkFAxi/foIR5xM="}, "signature":"MEQCIDzDmnZ7WXgq9nhNLPXxet4jmem/XxK9XkDTh0XaxBIhAiA7qSkwdXRQDNWOHSPb7TBi4KTLBesnjEsSOWCPQNy+ig=="}}
|
build/torch214-cxx11-cu126-x86_64-linux/__init__.py
ADDED
|
@@ -0,0 +1,205 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2024 Databricks
|
| 2 |
+
# SPDX-License-Identifier: Apache-2.0
|
| 3 |
+
|
| 4 |
+
import torch
|
| 5 |
+
|
| 6 |
+
# Stable alias: bare `ops` is shadowed by `from . import layers` below.
|
| 7 |
+
from ._ops import ops as _compiled_ops
|
| 8 |
+
from . import ops
|
| 9 |
+
|
| 10 |
+
from .grouped_gemm import backend as gg_backend
|
| 11 |
+
from .grouped_gemm import ops as gg_ops
|
| 12 |
+
|
| 13 |
+
|
| 14 |
+
from ._layers.arguments import Arguments
|
| 15 |
+
from ._layers.dmoe import ParallelDroplessMLP, dMoE
|
| 16 |
+
from ._layers.glu import SparseGLU
|
| 17 |
+
from ._layers.mlp import MLP, SparseMLP
|
| 18 |
+
from ._layers.moe import MoE, ParallelMLP, get_load_balancing_loss
|
| 19 |
+
|
| 20 |
+
from . import layers
|
| 21 |
+
|
| 22 |
+
# This section contains the direct kernel exports (not inlcuded in the original code)
|
| 23 |
+
def exclusive_cumsum(x: torch.Tensor, dim: int, out: torch.Tensor) -> torch.Tensor:
|
| 24 |
+
"""
|
| 25 |
+
Compute exclusive cumulative sum along the specified dimension.
|
| 26 |
+
|
| 27 |
+
Args:
|
| 28 |
+
x: Input tensor
|
| 29 |
+
dim: Dimension along which to compute cumsum
|
| 30 |
+
out: Output tensor (modified in-place)
|
| 31 |
+
|
| 32 |
+
Returns:
|
| 33 |
+
The output tensor
|
| 34 |
+
"""
|
| 35 |
+
result = ops.exclusive_cumsum(x, dim)
|
| 36 |
+
out.copy_(result)
|
| 37 |
+
return out
|
| 38 |
+
|
| 39 |
+
|
| 40 |
+
def inclusive_cumsum(x: torch.Tensor, dim: int, out: torch.Tensor) -> torch.Tensor:
|
| 41 |
+
"""
|
| 42 |
+
Compute inclusive cumulative sum along the specified dimension.
|
| 43 |
+
|
| 44 |
+
Args:
|
| 45 |
+
x: Input tensor
|
| 46 |
+
dim: Dimension along which to compute cumsum
|
| 47 |
+
out: Output tensor (modified in-place)
|
| 48 |
+
|
| 49 |
+
Returns:
|
| 50 |
+
The output tensor
|
| 51 |
+
"""
|
| 52 |
+
result = ops.inclusive_cumsum(x, dim)
|
| 53 |
+
out.copy_(result)
|
| 54 |
+
return out
|
| 55 |
+
|
| 56 |
+
|
| 57 |
+
def histogram(x: torch.Tensor, num_bins: int) -> torch.Tensor:
|
| 58 |
+
"""
|
| 59 |
+
Compute histogram of input tensor values.
|
| 60 |
+
|
| 61 |
+
Args:
|
| 62 |
+
x: Input tensor
|
| 63 |
+
num_bins: Number of histogram bins
|
| 64 |
+
|
| 65 |
+
Returns:
|
| 66 |
+
Histogram tensor with counts for each bin
|
| 67 |
+
"""
|
| 68 |
+
return ops.histogram(x, num_bins)
|
| 69 |
+
|
| 70 |
+
|
| 71 |
+
def indices(
|
| 72 |
+
padded_bins: torch.Tensor,
|
| 73 |
+
block_size: int,
|
| 74 |
+
output_block_rows: int,
|
| 75 |
+
output_block_columns: int,
|
| 76 |
+
) -> torch.Tensor:
|
| 77 |
+
"""
|
| 78 |
+
Construct indices from padded bins for sparse operations.
|
| 79 |
+
|
| 80 |
+
Args:
|
| 81 |
+
padded_bins: Tensor containing bin boundaries
|
| 82 |
+
block_size: Size of each block
|
| 83 |
+
output_block_rows: Number of rows in output blocks
|
| 84 |
+
output_block_columns: Number of columns in output blocks
|
| 85 |
+
|
| 86 |
+
Returns:
|
| 87 |
+
Tensor containing constructed indices
|
| 88 |
+
"""
|
| 89 |
+
return ops.indices(padded_bins, block_size, output_block_rows, output_block_columns)
|
| 90 |
+
|
| 91 |
+
|
| 92 |
+
def replicate_forward(
|
| 93 |
+
x: torch.Tensor, bins: torch.Tensor, out: torch.Tensor
|
| 94 |
+
) -> torch.Tensor:
|
| 95 |
+
"""
|
| 96 |
+
Forward pass of replicate operation - replicate values according to bin sizes.
|
| 97 |
+
|
| 98 |
+
Args:
|
| 99 |
+
x: Input tensor with values to replicate
|
| 100 |
+
bins: Tensor containing bin sizes
|
| 101 |
+
out: Output tensor (modified in-place)
|
| 102 |
+
|
| 103 |
+
Returns:
|
| 104 |
+
The output tensor
|
| 105 |
+
"""
|
| 106 |
+
return ops.replicate_forward(x, bins, out)
|
| 107 |
+
|
| 108 |
+
|
| 109 |
+
def replicate_backward(
|
| 110 |
+
grad: torch.Tensor, bins: torch.Tensor, out: torch.Tensor
|
| 111 |
+
) -> torch.Tensor:
|
| 112 |
+
"""
|
| 113 |
+
Backward pass of replicate operation - reduce gradients back to bins.
|
| 114 |
+
|
| 115 |
+
Args:
|
| 116 |
+
grad: Gradient tensor to reduce
|
| 117 |
+
bins: Tensor containing bin sizes
|
| 118 |
+
out: Output tensor (modified in-place)
|
| 119 |
+
|
| 120 |
+
Returns:
|
| 121 |
+
The output tensor
|
| 122 |
+
"""
|
| 123 |
+
return ops.replicate_backward(grad, bins, out)
|
| 124 |
+
|
| 125 |
+
|
| 126 |
+
def sort(
|
| 127 |
+
x: torch.Tensor, end_bit: int, x_out: torch.Tensor, iota_out: torch.Tensor
|
| 128 |
+
) -> torch.Tensor:
|
| 129 |
+
"""
|
| 130 |
+
Radix sort with index tracking.
|
| 131 |
+
|
| 132 |
+
Args:
|
| 133 |
+
x: Input tensor to sort
|
| 134 |
+
end_bit: Number of bits to consider in sorting
|
| 135 |
+
x_out: Output tensor for sorted values
|
| 136 |
+
iota_out: Output tensor for sorted indices
|
| 137 |
+
|
| 138 |
+
Returns:
|
| 139 |
+
The sorted values tensor
|
| 140 |
+
"""
|
| 141 |
+
_compiled_ops.sort(x, end_bit, x_out, iota_out)
|
| 142 |
+
return x_out
|
| 143 |
+
|
| 144 |
+
|
| 145 |
+
# Convenience functions for common use cases
|
| 146 |
+
def cumsum(x: torch.Tensor, dim: int = -1, exclusive: bool = False) -> torch.Tensor:
|
| 147 |
+
"""
|
| 148 |
+
Compute cumulative sum with automatic output allocation.
|
| 149 |
+
|
| 150 |
+
Args:
|
| 151 |
+
x: Input tensor
|
| 152 |
+
dim: Dimension along which to compute cumsum (default: last dimension)
|
| 153 |
+
exclusive: Whether to compute exclusive (True) or inclusive (False) cumsum
|
| 154 |
+
|
| 155 |
+
Returns:
|
| 156 |
+
New tensor containing the cumulative sum
|
| 157 |
+
"""
|
| 158 |
+
out = torch.empty_like(x)
|
| 159 |
+
if exclusive:
|
| 160 |
+
return exclusive_cumsum(x, dim, out)
|
| 161 |
+
else:
|
| 162 |
+
return inclusive_cumsum(x, dim, out)
|
| 163 |
+
|
| 164 |
+
|
| 165 |
+
def argsort(x: torch.Tensor, end_bit: int = 32) -> tuple[torch.Tensor, torch.Tensor]:
|
| 166 |
+
"""
|
| 167 |
+
Sort tensor and return both sorted values and indices.
|
| 168 |
+
|
| 169 |
+
Args:
|
| 170 |
+
x: Input tensor to sort
|
| 171 |
+
end_bit: Number of bits to consider in sorting
|
| 172 |
+
|
| 173 |
+
Returns:
|
| 174 |
+
Tuple of (sorted_values, sorted_indices)
|
| 175 |
+
"""
|
| 176 |
+
x_out = torch.empty_like(x)
|
| 177 |
+
iota_out = torch.empty_like(x)
|
| 178 |
+
sort(x, end_bit, x_out, iota_out)
|
| 179 |
+
return x_out, iota_out
|
| 180 |
+
|
| 181 |
+
|
| 182 |
+
# Export public API
|
| 183 |
+
__all__ = [
|
| 184 |
+
"MyReplacementLayer",
|
| 185 |
+
# Direct kernel exports
|
| 186 |
+
"exclusive_cumsum",
|
| 187 |
+
"inclusive_cumsum",
|
| 188 |
+
"histogram",
|
| 189 |
+
"indices",
|
| 190 |
+
"replicate_forward",
|
| 191 |
+
"replicate_backward",
|
| 192 |
+
"sort",
|
| 193 |
+
"cumsum",
|
| 194 |
+
"argsort",
|
| 195 |
+
# Original exports
|
| 196 |
+
"Arguments",
|
| 197 |
+
"ParallelDroplessMLP",
|
| 198 |
+
"dMoE",
|
| 199 |
+
"SparseGLU",
|
| 200 |
+
"MLP",
|
| 201 |
+
"SparseMLP",
|
| 202 |
+
"MoE",
|
| 203 |
+
"ParallelMLP",
|
| 204 |
+
"get_load_balancing_loss",
|
| 205 |
+
]
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/__init__.py
ADDED
|
File without changes
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/_triton_kernels/__init__.py
ADDED
|
File without changes
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/_triton_kernels/gmm.py
ADDED
|
@@ -0,0 +1,574 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# SPDX-License-Identifier: MIT
|
| 2 |
+
# Copyright (C) 2025-2026, Advanced Micro Devices, Inc. All rights reserved.
|
| 3 |
+
|
| 4 |
+
|
| 5 |
+
# Imports.
|
| 6 |
+
# ------------------------------------------------------------------------------
|
| 7 |
+
|
| 8 |
+
# Python standard library
|
| 9 |
+
import functools
|
| 10 |
+
|
| 11 |
+
# Triton
|
| 12 |
+
import triton
|
| 13 |
+
import triton.language as tl
|
| 14 |
+
|
| 15 |
+
# AITER
|
| 16 |
+
from ..configs import CONFIGS as _CONFIGS
|
| 17 |
+
from ..utils._triton import arch_info
|
| 18 |
+
from ..utils._triton.pid_preprocessing import pid_grid, remap_xcd
|
| 19 |
+
|
| 20 |
+
# Kernel config.
|
| 21 |
+
# ------------------------------------------------------------------------------
|
| 22 |
+
|
| 23 |
+
|
| 24 |
+
@functools.lru_cache()
|
| 25 |
+
def get_config(
|
| 26 |
+
gmm_type: str, M: int, K: int, N: int, G: int, accumulate: bool = False
|
| 27 |
+
) -> dict[str, int]:
|
| 28 |
+
assert gmm_type in {
|
| 29 |
+
"gmm",
|
| 30 |
+
"ptgmm",
|
| 31 |
+
"nptgmm",
|
| 32 |
+
}, f"'{gmm_type}' is an invalid GMM variant."
|
| 33 |
+
dev = arch_info.get_arch()
|
| 34 |
+
assert (
|
| 35 |
+
dev in _CONFIGS
|
| 36 |
+
), f"No GMM configuration tuned for arch '{dev}'. Supported: {sorted(_CONFIGS)}."
|
| 37 |
+
arch_configs = _CONFIGS[dev]
|
| 38 |
+
assert (
|
| 39 |
+
"default" in arch_configs[gmm_type]
|
| 40 |
+
), "Default configuration is absent."
|
| 41 |
+
key = "accumulate" if accumulate else "default"
|
| 42 |
+
return arch_configs[gmm_type][key]
|
| 43 |
+
|
| 44 |
+
|
| 45 |
+
# Common code shared by GMM and TGMM kernels.
|
| 46 |
+
# ------------------------------------------------------------------------------
|
| 47 |
+
|
| 48 |
+
|
| 49 |
+
# XCD remapping followed by 1D PID to 2D grid mapping.
|
| 50 |
+
@triton.jit
|
| 51 |
+
def _remap_xcd_tile_grid(
|
| 52 |
+
tile_in_mm,
|
| 53 |
+
num_row_tiles,
|
| 54 |
+
num_col_tiles,
|
| 55 |
+
GROUP_SIZE: tl.constexpr = 1,
|
| 56 |
+
NUM_XCDS: tl.constexpr = 8,
|
| 57 |
+
):
|
| 58 |
+
return pid_grid(
|
| 59 |
+
remap_xcd(tile_in_mm, num_row_tiles * num_col_tiles, NUM_XCDS=NUM_XCDS),
|
| 60 |
+
num_row_tiles,
|
| 61 |
+
num_col_tiles,
|
| 62 |
+
GROUP_SIZE_M=GROUP_SIZE,
|
| 63 |
+
)
|
| 64 |
+
|
| 65 |
+
|
| 66 |
+
# GMM kernel.
|
| 67 |
+
# ------------------------------------------------------------------------------
|
| 68 |
+
|
| 69 |
+
|
| 70 |
+
@triton.heuristics(
|
| 71 |
+
{
|
| 72 |
+
"K_DIVISIBLE_BY_BLOCK_SIZE_K": lambda META: META["K"] % META["BLOCK_SIZE_K"]
|
| 73 |
+
== 0,
|
| 74 |
+
}
|
| 75 |
+
)
|
| 76 |
+
@triton.jit
|
| 77 |
+
def gmm_kernel(
|
| 78 |
+
# Tensor pointers:
|
| 79 |
+
lhs_ptr,
|
| 80 |
+
rhs_ptr,
|
| 81 |
+
group_sizes_ptr,
|
| 82 |
+
out_ptr,
|
| 83 |
+
bias_ptr,
|
| 84 |
+
# Tensor shapes:
|
| 85 |
+
M: int,
|
| 86 |
+
K: int,
|
| 87 |
+
N: int,
|
| 88 |
+
G: int,
|
| 89 |
+
# Meta-parameters:
|
| 90 |
+
TRANS_RHS: tl.constexpr,
|
| 91 |
+
BLOCK_SIZE_M: tl.constexpr,
|
| 92 |
+
BLOCK_SIZE_K: tl.constexpr,
|
| 93 |
+
BLOCK_SIZE_N: tl.constexpr,
|
| 94 |
+
K_DIVISIBLE_BY_BLOCK_SIZE_K: tl.constexpr,
|
| 95 |
+
GROUP_SIZE: tl.constexpr,
|
| 96 |
+
GRID_DIM: tl.constexpr,
|
| 97 |
+
USE_BIAS: tl.constexpr,
|
| 98 |
+
):
|
| 99 |
+
tl.assume(M > 0)
|
| 100 |
+
tl.assume(K > 0)
|
| 101 |
+
tl.assume(N > 0)
|
| 102 |
+
tl.assume(G > 0)
|
| 103 |
+
|
| 104 |
+
num_n_tiles = tl.cdiv(N, BLOCK_SIZE_N)
|
| 105 |
+
tl.device_assert(num_n_tiles > 0, "num_n_tiles <= 0")
|
| 106 |
+
|
| 107 |
+
# Current tile. Each program computes multiple tiles of each group.
|
| 108 |
+
tile = tl.program_id(0)
|
| 109 |
+
tl.device_assert(tile >= 0, "tile < 0 (at initialization)")
|
| 110 |
+
|
| 111 |
+
# Tile limit of last MM problem (inclusive).
|
| 112 |
+
last_mm_tile = 0
|
| 113 |
+
|
| 114 |
+
# Last input row of lhs and output row of out. Each group reads some rows of
|
| 115 |
+
# lhs and writes some rows to out.
|
| 116 |
+
last_m = 0
|
| 117 |
+
|
| 118 |
+
# Loop through all (m, K, N) MM problems:
|
| 119 |
+
# (m, K) x (K, N) = (m, N)
|
| 120 |
+
# sum(m) = M
|
| 121 |
+
for g in range(G):
|
| 122 |
+
# Get m dimension of current MM problem.
|
| 123 |
+
m = tl.load(group_sizes_ptr + g)
|
| 124 |
+
# m can be zero if group is empty
|
| 125 |
+
tl.device_assert(m >= 0, "m < 0")
|
| 126 |
+
|
| 127 |
+
num_m_tiles = tl.cdiv(m, BLOCK_SIZE_M)
|
| 128 |
+
# num_m_tiles can be zero if group is empty
|
| 129 |
+
tl.device_assert(num_m_tiles >= 0, "num_m_tiles < 0")
|
| 130 |
+
|
| 131 |
+
num_tiles = num_m_tiles * num_n_tiles
|
| 132 |
+
# num_tiles can be zero if group is empty
|
| 133 |
+
tl.device_assert(num_tiles >= 0, "num_tiles < 0")
|
| 134 |
+
|
| 135 |
+
# Loop through tiles of current MM problem.
|
| 136 |
+
while tile >= last_mm_tile and tile < last_mm_tile + num_tiles:
|
| 137 |
+
# Figure out tile coordinates in current MM problem.
|
| 138 |
+
tile_in_mm = tile - last_mm_tile
|
| 139 |
+
tl.device_assert(tile_in_mm >= 0, "tile_in_mm < 0")
|
| 140 |
+
|
| 141 |
+
tile_m, tile_n = _remap_xcd_tile_grid(
|
| 142 |
+
tile_in_mm, num_m_tiles, num_n_tiles, GROUP_SIZE=GROUP_SIZE
|
| 143 |
+
)
|
| 144 |
+
|
| 145 |
+
# Do regular MM:
|
| 146 |
+
|
| 147 |
+
tl.device_assert(tile_m * BLOCK_SIZE_M >= 0, "tile_m * BLOCK_SIZE_M < 0")
|
| 148 |
+
tl.device_assert(tile_n * BLOCK_SIZE_N >= 0, "tile_n * BLOCK_SIZE_N < 0")
|
| 149 |
+
|
| 150 |
+
offs_lhs_m = (
|
| 151 |
+
tile_m.to(tl.int64) * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
|
| 152 |
+
) % m
|
| 153 |
+
offs_rhs_n = (
|
| 154 |
+
tile_n.to(tl.int64) * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
|
| 155 |
+
) % N
|
| 156 |
+
offs_k = tl.arange(0, BLOCK_SIZE_K).to(tl.int64)
|
| 157 |
+
|
| 158 |
+
lhs_ptrs = lhs_ptr + (last_m + offs_lhs_m[:, None]) * K + offs_k[None, :]
|
| 159 |
+
|
| 160 |
+
if TRANS_RHS:
|
| 161 |
+
rhs_ptrs = (
|
| 162 |
+
rhs_ptr
|
| 163 |
+
+ g.to(tl.int64) * K * N
|
| 164 |
+
+ offs_k[:, None]
|
| 165 |
+
+ offs_rhs_n[None, :] * K
|
| 166 |
+
)
|
| 167 |
+
else:
|
| 168 |
+
rhs_ptrs = (
|
| 169 |
+
rhs_ptr
|
| 170 |
+
+ g.to(tl.int64) * K * N
|
| 171 |
+
+ offs_k[:, None] * N
|
| 172 |
+
+ offs_rhs_n[None, :]
|
| 173 |
+
)
|
| 174 |
+
|
| 175 |
+
acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
|
| 176 |
+
|
| 177 |
+
for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
|
| 178 |
+
if K_DIVISIBLE_BY_BLOCK_SIZE_K:
|
| 179 |
+
lhs = tl.load(lhs_ptrs)
|
| 180 |
+
rhs = tl.load(rhs_ptrs)
|
| 181 |
+
else:
|
| 182 |
+
k_mask_limit = K - k * BLOCK_SIZE_K
|
| 183 |
+
lhs = tl.load(
|
| 184 |
+
lhs_ptrs, mask=offs_k[None, :] < k_mask_limit, other=0
|
| 185 |
+
)
|
| 186 |
+
rhs = tl.load(
|
| 187 |
+
rhs_ptrs, mask=offs_k[:, None] < k_mask_limit, other=0
|
| 188 |
+
)
|
| 189 |
+
|
| 190 |
+
acc = tl.dot(lhs, rhs, acc=acc)
|
| 191 |
+
|
| 192 |
+
lhs_ptrs += BLOCK_SIZE_K
|
| 193 |
+
|
| 194 |
+
if TRANS_RHS:
|
| 195 |
+
rhs_ptrs += BLOCK_SIZE_K
|
| 196 |
+
else:
|
| 197 |
+
rhs_ptrs += BLOCK_SIZE_K * N
|
| 198 |
+
|
| 199 |
+
# Add bias if enabled
|
| 200 |
+
if USE_BIAS:
|
| 201 |
+
offs_bias_n = tile_n.to(tl.int64) * BLOCK_SIZE_N + tl.arange(
|
| 202 |
+
0, BLOCK_SIZE_N
|
| 203 |
+
)
|
| 204 |
+
bias_ptrs = bias_ptr + g.to(tl.int64) * N + offs_bias_n
|
| 205 |
+
bias = tl.load(bias_ptrs, mask=offs_bias_n < N, other=0.0)
|
| 206 |
+
# Convert bias to float32 to match accumulator precision
|
| 207 |
+
bias = bias.to(tl.float32)
|
| 208 |
+
# Broadcast bias across M dimension and add in float32
|
| 209 |
+
acc += bias[None, :]
|
| 210 |
+
|
| 211 |
+
# Convert to output dtype after all computations
|
| 212 |
+
acc = acc.to(out_ptr.type.element_ty)
|
| 213 |
+
|
| 214 |
+
offs_out_m = tile_m.to(tl.int64) * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
|
| 215 |
+
offs_out_n = tile_n.to(tl.int64) * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
|
| 216 |
+
|
| 217 |
+
out_ptrs = (
|
| 218 |
+
out_ptr + (last_m + offs_out_m[:, None]) * N + offs_out_n[None, :]
|
| 219 |
+
)
|
| 220 |
+
|
| 221 |
+
tl.store(
|
| 222 |
+
out_ptrs,
|
| 223 |
+
acc,
|
| 224 |
+
mask=(offs_out_m[:, None] < m) & (offs_out_n[None, :] < N),
|
| 225 |
+
)
|
| 226 |
+
|
| 227 |
+
# Go to the next tile by advancing number of programs.
|
| 228 |
+
tile += GRID_DIM
|
| 229 |
+
tl.device_assert(tile > 0, "tile <= 0 (at update)")
|
| 230 |
+
|
| 231 |
+
# Get ready to go to the next MM problem.
|
| 232 |
+
|
| 233 |
+
last_mm_tile += num_tiles
|
| 234 |
+
# last_mm_tile can be zero if group 0 is skipped
|
| 235 |
+
tl.device_assert(last_mm_tile >= 0, "last_mm_tile < 0 (at update)")
|
| 236 |
+
|
| 237 |
+
last_m += m
|
| 238 |
+
# last_m can be zero if group 0 is skipped
|
| 239 |
+
tl.device_assert(last_m >= 0, "last_m < 0 (at update)")
|
| 240 |
+
tl.device_assert(last_m <= M, "last_m > M (at update)")
|
| 241 |
+
|
| 242 |
+
|
| 243 |
+
# Persistent TGMM kernel.
|
| 244 |
+
# ------------------------------------------------------------------------------
|
| 245 |
+
|
| 246 |
+
|
| 247 |
+
@triton.jit
|
| 248 |
+
def tgmm_persistent_kernel(
|
| 249 |
+
# Tensor pointers:
|
| 250 |
+
lhs_ptr,
|
| 251 |
+
rhs_ptr,
|
| 252 |
+
group_sizes_ptr,
|
| 253 |
+
out_ptr,
|
| 254 |
+
bias_grad_ptr,
|
| 255 |
+
# Tensor shapes:
|
| 256 |
+
M: int,
|
| 257 |
+
K: int,
|
| 258 |
+
N: int,
|
| 259 |
+
G: int,
|
| 260 |
+
# Meta-parameters:
|
| 261 |
+
TRANS_LHS: tl.constexpr,
|
| 262 |
+
BLOCK_SIZE_M: tl.constexpr,
|
| 263 |
+
BLOCK_SIZE_K: tl.constexpr,
|
| 264 |
+
BLOCK_SIZE_N: tl.constexpr,
|
| 265 |
+
GROUP_SIZE: tl.constexpr,
|
| 266 |
+
GRID_DIM: tl.constexpr,
|
| 267 |
+
COMPUTE_BIAS_GRAD: tl.constexpr,
|
| 268 |
+
ACCUMULATE: tl.constexpr,
|
| 269 |
+
):
|
| 270 |
+
tl.assume(M > 0)
|
| 271 |
+
tl.assume(K > 0)
|
| 272 |
+
tl.assume(N > 0)
|
| 273 |
+
tl.assume(G > 0)
|
| 274 |
+
|
| 275 |
+
num_k_tiles = tl.cdiv(K, BLOCK_SIZE_K)
|
| 276 |
+
tl.device_assert(num_k_tiles > 0, "num_k_tiles <= 0")
|
| 277 |
+
|
| 278 |
+
num_n_tiles = tl.cdiv(N, BLOCK_SIZE_N)
|
| 279 |
+
tl.device_assert(num_n_tiles > 0, "num_n_tiles <= 0")
|
| 280 |
+
|
| 281 |
+
num_tiles = num_k_tiles * num_n_tiles
|
| 282 |
+
tl.device_assert(num_tiles > 0, "num_tiles <= 0")
|
| 283 |
+
|
| 284 |
+
# Current tile. Each program computes multiple tiles of each group.
|
| 285 |
+
tile = tl.program_id(0)
|
| 286 |
+
tl.device_assert(tile >= 0, "tile < 0 (at initialization)")
|
| 287 |
+
|
| 288 |
+
# Tile limit of last MM problem (inclusive).
|
| 289 |
+
last_mm_tile = 0
|
| 290 |
+
|
| 291 |
+
# Last input column of lhs and input row of rhs. Each group reads some
|
| 292 |
+
# columns of lhs and some rows of rhs.
|
| 293 |
+
last_m = 0
|
| 294 |
+
|
| 295 |
+
# Loop through all (K, m, N) MM problems:
|
| 296 |
+
# (K, m) x (m, N) = (K, N)
|
| 297 |
+
# sum(m) = M
|
| 298 |
+
for g in range(G):
|
| 299 |
+
# Get m dimension of current MM problem.
|
| 300 |
+
m = tl.load(group_sizes_ptr + g)
|
| 301 |
+
# m can be zero if group is empty
|
| 302 |
+
tl.device_assert(m >= 0, "m < 0")
|
| 303 |
+
|
| 304 |
+
# Loop through tiles of current MM problem.
|
| 305 |
+
while tile >= last_mm_tile and tile < last_mm_tile + num_tiles:
|
| 306 |
+
# Figure out tile coordinates in current MM problem.
|
| 307 |
+
tile_in_mm = tile - last_mm_tile
|
| 308 |
+
tl.device_assert(tile_in_mm >= 0, "tile_in_mm < 0")
|
| 309 |
+
|
| 310 |
+
tile_k, tile_n = _remap_xcd_tile_grid(
|
| 311 |
+
tile_in_mm, num_k_tiles, num_n_tiles, GROUP_SIZE=GROUP_SIZE
|
| 312 |
+
)
|
| 313 |
+
|
| 314 |
+
# Do regular MM:
|
| 315 |
+
|
| 316 |
+
tl.device_assert(tile_k * BLOCK_SIZE_K >= 0, "tile_k * BLOCK_SIZE_K < 0")
|
| 317 |
+
tl.device_assert(tile_n * BLOCK_SIZE_N >= 0, "tile_n * BLOCK_SIZE_N < 0")
|
| 318 |
+
|
| 319 |
+
offs_lhs_k = (
|
| 320 |
+
tile_k.to(tl.int64) * BLOCK_SIZE_K + tl.arange(0, BLOCK_SIZE_K)
|
| 321 |
+
) % K
|
| 322 |
+
offs_rhs_n = (
|
| 323 |
+
tile_n.to(tl.int64) * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
|
| 324 |
+
) % N
|
| 325 |
+
offs_m = tl.arange(0, BLOCK_SIZE_M).to(tl.int64)
|
| 326 |
+
|
| 327 |
+
if TRANS_LHS:
|
| 328 |
+
lhs_ptrs = (
|
| 329 |
+
lhs_ptr + offs_lhs_k[:, None] + (last_m + offs_m[None, :]) * K
|
| 330 |
+
)
|
| 331 |
+
else:
|
| 332 |
+
lhs_ptrs = (
|
| 333 |
+
lhs_ptr + offs_lhs_k[:, None] * M + (last_m + offs_m[None, :])
|
| 334 |
+
)
|
| 335 |
+
|
| 336 |
+
rhs_ptrs = rhs_ptr + (last_m + offs_m[:, None]) * N + offs_rhs_n[None, :]
|
| 337 |
+
|
| 338 |
+
loop_m = tl.cdiv(m, BLOCK_SIZE_M)
|
| 339 |
+
m_divisible_by_block_m = m % BLOCK_SIZE_M == 0
|
| 340 |
+
if not m_divisible_by_block_m:
|
| 341 |
+
loop_m -= 1
|
| 342 |
+
|
| 343 |
+
acc = tl.zeros((BLOCK_SIZE_K, BLOCK_SIZE_N), dtype=tl.float32)
|
| 344 |
+
|
| 345 |
+
# Initialize bias accumulator
|
| 346 |
+
bias_acc = tl.zeros((BLOCK_SIZE_K,), dtype=tl.float32)
|
| 347 |
+
|
| 348 |
+
for _ in range(0, loop_m):
|
| 349 |
+
lhs = tl.load(lhs_ptrs)
|
| 350 |
+
rhs = tl.load(rhs_ptrs)
|
| 351 |
+
|
| 352 |
+
acc = tl.dot(lhs, rhs, acc=acc)
|
| 353 |
+
|
| 354 |
+
# Accumulate for bias gradient: sum lhs across M dimension
|
| 355 |
+
if COMPUTE_BIAS_GRAD and tile_n == 0:
|
| 356 |
+
bias_acc += tl.sum(
|
| 357 |
+
lhs, axis=1
|
| 358 |
+
) # Sum across M dimension [K, M] -> [K]
|
| 359 |
+
|
| 360 |
+
if TRANS_LHS:
|
| 361 |
+
lhs_ptrs += BLOCK_SIZE_M * K
|
| 362 |
+
else:
|
| 363 |
+
lhs_ptrs += BLOCK_SIZE_M
|
| 364 |
+
|
| 365 |
+
rhs_ptrs += BLOCK_SIZE_M * N
|
| 366 |
+
|
| 367 |
+
if not m_divisible_by_block_m:
|
| 368 |
+
offs_lhs_k = (
|
| 369 |
+
tile_k.to(tl.int64) * BLOCK_SIZE_K + tl.arange(0, BLOCK_SIZE_K)
|
| 370 |
+
) % K
|
| 371 |
+
offs_rhs_n = (
|
| 372 |
+
tile_n.to(tl.int64) * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
|
| 373 |
+
) % N
|
| 374 |
+
offs_m = loop_m.to(tl.int64) * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
|
| 375 |
+
lhs = tl.load(lhs_ptrs, mask=offs_m[None, :] < m, other=0)
|
| 376 |
+
rhs = tl.load(rhs_ptrs, mask=offs_m[:, None] < m, other=0)
|
| 377 |
+
acc = tl.dot(lhs, rhs, acc=acc)
|
| 378 |
+
|
| 379 |
+
# Accumulate last chunk for bias gradient
|
| 380 |
+
if COMPUTE_BIAS_GRAD and tile_n == 0:
|
| 381 |
+
bias_acc += tl.sum(lhs, axis=1)
|
| 382 |
+
|
| 383 |
+
acc = acc.to(out_ptr.type.element_ty)
|
| 384 |
+
|
| 385 |
+
offs_out_k = tile_k.to(tl.int64) * BLOCK_SIZE_K + tl.arange(0, BLOCK_SIZE_K)
|
| 386 |
+
offs_out_n = tile_n.to(tl.int64) * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
|
| 387 |
+
|
| 388 |
+
out_ptrs = (
|
| 389 |
+
out_ptr
|
| 390 |
+
+ g.to(tl.int64) * K * N
|
| 391 |
+
+ offs_out_k[:, None] * N
|
| 392 |
+
+ offs_out_n[None, :]
|
| 393 |
+
)
|
| 394 |
+
|
| 395 |
+
mask = (offs_out_k[:, None] < K) & (offs_out_n[None, :] < N)
|
| 396 |
+
if ACCUMULATE:
|
| 397 |
+
# Load existing values and add to them (like beta=1 in BLAS)
|
| 398 |
+
old_vals = tl.load(out_ptrs, mask=mask, other=0.0)
|
| 399 |
+
tl.store(out_ptrs, acc + old_vals, mask=mask)
|
| 400 |
+
else:
|
| 401 |
+
# Overwrite output (like beta=0 in BLAS)
|
| 402 |
+
tl.store(out_ptrs, acc, mask=mask)
|
| 403 |
+
|
| 404 |
+
# Store bias gradient (only for first N tile, sum across all M)
|
| 405 |
+
if COMPUTE_BIAS_GRAD and tile_n == 0:
|
| 406 |
+
# Keep as float32 for atomic_add (bf16 not supported for atomics)
|
| 407 |
+
bias_grad_ptrs = bias_grad_ptr + g.to(tl.int64) * K + offs_out_k
|
| 408 |
+
# Use atomic add since multiple K-tiles may write to same expert's bias
|
| 409 |
+
tl.atomic_add(
|
| 410 |
+
bias_grad_ptrs, bias_acc, mask=offs_out_k < K, sem="relaxed"
|
| 411 |
+
)
|
| 412 |
+
|
| 413 |
+
# Go to the next tile by advancing number of programs.
|
| 414 |
+
tile += GRID_DIM
|
| 415 |
+
tl.device_assert(tile > 0, "tile <= 0 (at update)")
|
| 416 |
+
|
| 417 |
+
# Get ready to go to the next MM problem.
|
| 418 |
+
|
| 419 |
+
last_mm_tile += num_tiles
|
| 420 |
+
# last_mm_tile can be zero if group 0 is skipped
|
| 421 |
+
tl.device_assert(last_mm_tile >= 0, "last_mm_tile < 0 (at update)")
|
| 422 |
+
|
| 423 |
+
last_m += m
|
| 424 |
+
# last_m can be zero if group 0 is skipped
|
| 425 |
+
tl.device_assert(last_m >= 0, "last_m < 0 (at update)")
|
| 426 |
+
tl.device_assert(last_m <= M, "last_m > M (at update)")
|
| 427 |
+
|
| 428 |
+
|
| 429 |
+
# Regular non-persistent TGMM kernel.
|
| 430 |
+
# ------------------------------------------------------------------------------
|
| 431 |
+
|
| 432 |
+
|
| 433 |
+
@triton.heuristics({"BLOCK_SIZE_G": lambda META: triton.next_power_of_2(META["G"])})
|
| 434 |
+
@triton.jit
|
| 435 |
+
def tgmm_non_persistent_kernel(
|
| 436 |
+
# Tensor pointers:
|
| 437 |
+
lhs_ptr,
|
| 438 |
+
rhs_ptr,
|
| 439 |
+
group_sizes_ptr,
|
| 440 |
+
out_ptr,
|
| 441 |
+
bias_grad_ptr,
|
| 442 |
+
# Tensor shapes:
|
| 443 |
+
M: int,
|
| 444 |
+
K: int,
|
| 445 |
+
N: int,
|
| 446 |
+
G: int,
|
| 447 |
+
# Meta-parameters:
|
| 448 |
+
TRANS_LHS: tl.constexpr,
|
| 449 |
+
BLOCK_SIZE_G: tl.constexpr,
|
| 450 |
+
BLOCK_SIZE_M: tl.constexpr,
|
| 451 |
+
BLOCK_SIZE_K: tl.constexpr,
|
| 452 |
+
BLOCK_SIZE_N: tl.constexpr,
|
| 453 |
+
GROUP_SIZE: tl.constexpr,
|
| 454 |
+
COMPUTE_BIAS_GRAD: tl.constexpr,
|
| 455 |
+
ACCUMULATE: tl.constexpr,
|
| 456 |
+
):
|
| 457 |
+
tl.assume(M > 0)
|
| 458 |
+
tl.assume(K > 0)
|
| 459 |
+
tl.assume(N > 0)
|
| 460 |
+
tl.assume(G > 0)
|
| 461 |
+
|
| 462 |
+
# Get group ID from grid.
|
| 463 |
+
g = tl.program_id(0)
|
| 464 |
+
tl.device_assert(g >= 0, "g < 0")
|
| 465 |
+
tl.device_assert(g < G, "g >= G")
|
| 466 |
+
|
| 467 |
+
# Get m dimension of current MM group.
|
| 468 |
+
m = tl.load(group_sizes_ptr + g)
|
| 469 |
+
# m can be zero if group is empty.
|
| 470 |
+
tl.device_assert(m >= 0, "m < 0")
|
| 471 |
+
|
| 472 |
+
# Skip empty groups.
|
| 473 |
+
if m == 0:
|
| 474 |
+
return
|
| 475 |
+
|
| 476 |
+
# Compute sum(group_sizes) until current group g.
|
| 477 |
+
# It's the starting column of lhs and starting row of rhs.
|
| 478 |
+
offs_g = tl.arange(0, BLOCK_SIZE_G)
|
| 479 |
+
group_sizes = tl.load(group_sizes_ptr + offs_g, mask=offs_g < g, other=0)
|
| 480 |
+
start_m = tl.sum(group_sizes)
|
| 481 |
+
|
| 482 |
+
num_k_tiles = tl.cdiv(K, BLOCK_SIZE_K)
|
| 483 |
+
tl.device_assert(num_k_tiles > 0, "num_k_tiles <= 0")
|
| 484 |
+
|
| 485 |
+
num_n_tiles = tl.cdiv(N, BLOCK_SIZE_N)
|
| 486 |
+
tl.device_assert(num_n_tiles > 0, "num_n_tiles <= 0")
|
| 487 |
+
|
| 488 |
+
# Get MM tile from grid.
|
| 489 |
+
tile_in_mm = tl.program_id(1)
|
| 490 |
+
tl.device_assert(tile_in_mm >= 0, "tile_in_mm < 0")
|
| 491 |
+
|
| 492 |
+
tile_k, tile_n = _remap_xcd_tile_grid(
|
| 493 |
+
tile_in_mm, num_k_tiles, num_n_tiles, GROUP_SIZE=GROUP_SIZE
|
| 494 |
+
)
|
| 495 |
+
|
| 496 |
+
tl.device_assert(tile_k * BLOCK_SIZE_K >= 0, "tile_k * BLOCK_SIZE_K < 0")
|
| 497 |
+
tl.device_assert(tile_n * BLOCK_SIZE_N >= 0, "tile_n * BLOCK_SIZE_N < 0")
|
| 498 |
+
|
| 499 |
+
offs_lhs_k = (tile_k.to(tl.int64) * BLOCK_SIZE_K + tl.arange(0, BLOCK_SIZE_K)) % K
|
| 500 |
+
offs_rhs_n = (tile_n.to(tl.int64) * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
|
| 501 |
+
offs_m = tl.arange(0, BLOCK_SIZE_M).to(tl.int64)
|
| 502 |
+
|
| 503 |
+
if TRANS_LHS:
|
| 504 |
+
lhs_ptrs = lhs_ptr + offs_lhs_k[:, None] + (start_m + offs_m[None, :]) * K
|
| 505 |
+
else:
|
| 506 |
+
lhs_ptrs = lhs_ptr + offs_lhs_k[:, None] * M + (start_m + offs_m[None, :])
|
| 507 |
+
|
| 508 |
+
rhs_ptrs = rhs_ptr + (start_m + offs_m[:, None]) * N + offs_rhs_n[None, :]
|
| 509 |
+
|
| 510 |
+
loop_m = tl.cdiv(m, BLOCK_SIZE_M)
|
| 511 |
+
m_divisible_by_block_m = m % BLOCK_SIZE_M == 0
|
| 512 |
+
if not m_divisible_by_block_m:
|
| 513 |
+
loop_m -= 1
|
| 514 |
+
|
| 515 |
+
acc = tl.zeros((BLOCK_SIZE_K, BLOCK_SIZE_N), dtype=tl.float32)
|
| 516 |
+
# Initialize bias accumulator
|
| 517 |
+
bias_acc = tl.zeros((BLOCK_SIZE_K,), dtype=tl.float32)
|
| 518 |
+
|
| 519 |
+
for _ in range(0, loop_m):
|
| 520 |
+
lhs = tl.load(lhs_ptrs)
|
| 521 |
+
rhs = tl.load(rhs_ptrs)
|
| 522 |
+
|
| 523 |
+
acc = tl.dot(lhs, rhs, acc=acc)
|
| 524 |
+
|
| 525 |
+
# Accumulate for bias gradient: sum lhs across M dimension
|
| 526 |
+
if COMPUTE_BIAS_GRAD and tile_n == 0:
|
| 527 |
+
bias_acc += tl.sum(lhs, axis=1) # [K, M] -> [K]
|
| 528 |
+
|
| 529 |
+
if TRANS_LHS:
|
| 530 |
+
lhs_ptrs += BLOCK_SIZE_M * K
|
| 531 |
+
else:
|
| 532 |
+
lhs_ptrs += BLOCK_SIZE_M
|
| 533 |
+
|
| 534 |
+
rhs_ptrs += BLOCK_SIZE_M * N
|
| 535 |
+
|
| 536 |
+
if not m_divisible_by_block_m:
|
| 537 |
+
offs_lhs_k = (
|
| 538 |
+
tile_k.to(tl.int64) * BLOCK_SIZE_K + tl.arange(0, BLOCK_SIZE_K)
|
| 539 |
+
) % K
|
| 540 |
+
offs_rhs_n = (
|
| 541 |
+
tile_n.to(tl.int64) * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
|
| 542 |
+
) % N
|
| 543 |
+
offs_m = loop_m.to(tl.int64) * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
|
| 544 |
+
lhs = tl.load(lhs_ptrs, mask=offs_m[None, :] < m, other=0)
|
| 545 |
+
rhs = tl.load(rhs_ptrs, mask=offs_m[:, None] < m, other=0)
|
| 546 |
+
acc = tl.dot(lhs, rhs, acc=acc)
|
| 547 |
+
# Accumulate last chunk for bias gradient
|
| 548 |
+
if COMPUTE_BIAS_GRAD and tile_n == 0:
|
| 549 |
+
bias_acc += tl.sum(lhs, axis=1)
|
| 550 |
+
|
| 551 |
+
acc = acc.to(out_ptr.type.element_ty)
|
| 552 |
+
|
| 553 |
+
offs_out_k = tile_k.to(tl.int64) * BLOCK_SIZE_K + tl.arange(0, BLOCK_SIZE_K)
|
| 554 |
+
offs_out_n = tile_n.to(tl.int64) * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
|
| 555 |
+
|
| 556 |
+
out_ptrs = (
|
| 557 |
+
out_ptr + g.to(tl.int64) * K * N + offs_out_k[:, None] * N + offs_out_n[None, :]
|
| 558 |
+
)
|
| 559 |
+
|
| 560 |
+
mask = (offs_out_k[:, None] < K) & (offs_out_n[None, :] < N)
|
| 561 |
+
if ACCUMULATE:
|
| 562 |
+
# Load existing values and add to them (like beta=1 in BLAS)
|
| 563 |
+
old_vals = tl.load(out_ptrs, mask=mask, other=0.0)
|
| 564 |
+
tl.store(out_ptrs, acc + old_vals, mask=mask)
|
| 565 |
+
else:
|
| 566 |
+
# Overwrite output (like beta=0 in BLAS)
|
| 567 |
+
tl.store(out_ptrs, acc, mask=mask)
|
| 568 |
+
|
| 569 |
+
# Store bias gradient (only for first N tile, sum across all M)
|
| 570 |
+
if COMPUTE_BIAS_GRAD and tile_n == 0:
|
| 571 |
+
# Keep as float32 for atomic_add (bf16/fp16 not supported for atomics)
|
| 572 |
+
bias_grad_ptrs = bias_grad_ptr + g.to(tl.int64) * K + offs_out_k
|
| 573 |
+
# Use atomic add since multiple K-tiles may write to same expert's bias
|
| 574 |
+
tl.atomic_add(bias_grad_ptrs, bias_acc, mask=offs_out_k < K, sem="relaxed")
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/adapter.py
ADDED
|
@@ -0,0 +1,53 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# SPDX-License-Identifier: Apache-2.0
|
| 2 |
+
"""Adapt AITER's Triton grouped GEMM to MegaBlocks' ``gmm`` calling convention.
|
| 3 |
+
|
| 4 |
+
MegaBlocks (following tgale96/grouped_gemm) uses a single ``gmm`` entry point
|
| 5 |
+
with ``trans_a`` / ``trans_b`` flags:
|
| 6 |
+
|
| 7 |
+
* ``trans_a=False, trans_b=False``: a(M,K) @ b(G,K,N) -> c(M,N)
|
| 8 |
+
* ``trans_a=False, trans_b=True`` : a(M,K) @ b(G,N,K)^T -> c(M,N) (dgrad)
|
| 9 |
+
* ``trans_a=True`` : a(M,K)^T @ b(M,N) per group -> c(G,K,N) (wgrad)
|
| 10 |
+
|
| 11 |
+
AITER exposes these as two kernels: ``gmm`` ((M,K)@(G,K,N)->(M,N), transposition
|
| 12 |
+
of the 3D operand inferred from strides) and ``ptgmm`` ((K,M)@(M,N)->(G,K,N),
|
| 13 |
+
transposition of the 2D operand inferred from strides).
|
| 14 |
+
"""
|
| 15 |
+
|
| 16 |
+
import torch
|
| 17 |
+
|
| 18 |
+
from .gmm import gmm as _aiter_gmm
|
| 19 |
+
from .gmm import ptgmm as _aiter_ptgmm
|
| 20 |
+
|
| 21 |
+
|
| 22 |
+
def gmm(a, b, c, batch_sizes, trans_a=False, trans_b=False):
|
| 23 |
+
# AITER requires group sizes to be int32 and to live on the compute device.
|
| 24 |
+
group_sizes = batch_sizes.to(device=a.device, dtype=torch.int32)
|
| 25 |
+
|
| 26 |
+
# AITER asserts exact strides: gmm wants lhs/rhs row-major (a transposed
|
| 27 |
+
# 3D operand must be exactly column-major), tgmm wants rhs row-major and
|
| 28 |
+
# lhs row/column-major. Make operands contiguous first so the transposed
|
| 29 |
+
# views have the precise strides the kernels expect. `.contiguous()` is a
|
| 30 |
+
# no-op when the tensor is already contiguous.
|
| 31 |
+
if trans_a:
|
| 32 |
+
# Weight gradient: a(M,K), b(M,N) -> c(G,K,N).
|
| 33 |
+
# Pass a transposed so AITER sees lhs(K,M) column-major (TRANS_LHS).
|
| 34 |
+
_aiter_ptgmm(
|
| 35 |
+
a.contiguous().transpose(0, 1),
|
| 36 |
+
b.contiguous(),
|
| 37 |
+
group_sizes,
|
| 38 |
+
preferred_element_type=c.dtype,
|
| 39 |
+
existing_out=c,
|
| 40 |
+
)
|
| 41 |
+
else:
|
| 42 |
+
# trans_b contracts b's last dim: pass a column-major (G,K,N) view.
|
| 43 |
+
rhs = b.contiguous()
|
| 44 |
+
if trans_b:
|
| 45 |
+
rhs = rhs.transpose(1, 2)
|
| 46 |
+
_aiter_gmm(
|
| 47 |
+
a.contiguous(),
|
| 48 |
+
rhs,
|
| 49 |
+
group_sizes,
|
| 50 |
+
preferred_element_type=c.dtype,
|
| 51 |
+
existing_out=c,
|
| 52 |
+
)
|
| 53 |
+
return c
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/configs.py
ADDED
|
@@ -0,0 +1,5 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# SPDX-License-Identifier: MIT
|
| 2 |
+
# Tuned GMM configs vendored from ROCm/aiter (aiter/ops/triton/configs/).
|
| 3 |
+
# Inlined as a Python module so packaging always includes them.
|
| 4 |
+
|
| 5 |
+
CONFIGS = {'gfx1250': {'gmm': {'default': {'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 256, 'GROUP_SIZE': 1, 'GRID_DIM': 256, 'num_warps': 8, 'num_stages': 1}}, 'ptgmm': {'default': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N': 256, 'GROUP_SIZE': 1, 'GRID_DIM': 256, 'num_warps': 8, 'num_stages': 1}, 'accumulate': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N': 128, 'GROUP_SIZE': 1, 'GRID_DIM': 256, 'num_warps': 8, 'num_stages': 1}}, 'nptgmm': {'default': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N': 256, 'GROUP_SIZE': 1, 'num_warps': 8, 'num_stages': 1}, 'accumulate': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N': 128, 'GROUP_SIZE': 1, 'num_warps': 8, 'num_stages': 1}}}, 'gfx942': {'gmm': {'default': {'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 256, 'GROUP_SIZE': 1, 'GRID_DIM': 304, 'num_warps': 8, 'num_stages': 1}}, 'ptgmm': {'default': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N': 256, 'GROUP_SIZE': 1, 'GRID_DIM': 304, 'num_warps': 8, 'num_stages': 1}, 'accumulate': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N': 128, 'GROUP_SIZE': 1, 'GRID_DIM': 304, 'num_warps': 8, 'num_stages': 1}}, 'nptgmm': {'default': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N': 256, 'GROUP_SIZE': 1, 'num_warps': 8, 'num_stages': 1}, 'accumulate': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N': 128, 'GROUP_SIZE': 1, 'num_warps': 8, 'num_stages': 1}}}, 'gfx950': {'gmm': {'default': {'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 256, 'GROUP_SIZE': 1, 'GRID_DIM': 256, 'num_warps': 8, 'num_stages': 1}}, 'ptgmm': {'default': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N': 256, 'GROUP_SIZE': 1, 'GRID_DIM': 256, 'num_warps': 8, 'num_stages': 1}, 'accumulate': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N': 128, 'GROUP_SIZE': 1, 'GRID_DIM': 256, 'num_warps': 8, 'num_stages': 1}}, 'nptgmm': {'default': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N': 256, 'GROUP_SIZE': 1, 'num_warps': 8, 'num_stages': 1}, 'accumulate': {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N': 128, 'GROUP_SIZE': 1, 'num_warps': 8, 'num_stages': 1}}}}
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/gmm.py
ADDED
|
@@ -0,0 +1,567 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# SPDX-License-Identifier: MIT
|
| 2 |
+
# Copyright (C) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
| 3 |
+
|
| 4 |
+
|
| 5 |
+
# Imports.
|
| 6 |
+
# ------------------------------------------------------------------------------
|
| 7 |
+
|
| 8 |
+
# PyTorch
|
| 9 |
+
import torch
|
| 10 |
+
from torch import Tensor
|
| 11 |
+
|
| 12 |
+
# Triton
|
| 13 |
+
import triton
|
| 14 |
+
|
| 15 |
+
# AITER: GMM utility functions
|
| 16 |
+
from .utils.gmm_common import (
|
| 17 |
+
DTYPE,
|
| 18 |
+
is_power_of_2,
|
| 19 |
+
check_input_device_dtype,
|
| 20 |
+
check_bias_shape_stride,
|
| 21 |
+
get_gmm_shape,
|
| 22 |
+
get_gmm_output,
|
| 23 |
+
get_gmm_transposition,
|
| 24 |
+
get_tgmm_shape,
|
| 25 |
+
get_tgmm_output,
|
| 26 |
+
get_tgmm_bias_grad,
|
| 27 |
+
get_tgmm_transposition,
|
| 28 |
+
)
|
| 29 |
+
|
| 30 |
+
# AITER: GMM Triton kernels
|
| 31 |
+
from ._triton_kernels.gmm import (
|
| 32 |
+
gmm_kernel,
|
| 33 |
+
tgmm_persistent_kernel,
|
| 34 |
+
tgmm_non_persistent_kernel,
|
| 35 |
+
get_config,
|
| 36 |
+
)
|
| 37 |
+
|
| 38 |
+
# GMM PyTorch wrapper.
|
| 39 |
+
# ------------------------------------------------------------------------------
|
| 40 |
+
|
| 41 |
+
|
| 42 |
+
def _gmm_grid(
|
| 43 |
+
N: int,
|
| 44 |
+
block_size_m: int,
|
| 45 |
+
block_size_n: int,
|
| 46 |
+
group_sizes: Tensor,
|
| 47 |
+
grid_dim: int,
|
| 48 |
+
) -> tuple[int]:
|
| 49 |
+
assert N > 0, f"N must be positive, it's {N}."
|
| 50 |
+
assert is_power_of_2(
|
| 51 |
+
block_size_m
|
| 52 |
+
), f"M-dimension tile size must be a power of 2 (it's {block_size_m})."
|
| 53 |
+
assert is_power_of_2(
|
| 54 |
+
block_size_n
|
| 55 |
+
), f"N-dimension tile size must be a power of 2 (it's {block_size_n})."
|
| 56 |
+
assert torch.all(group_sizes >= 0).item(), "All group_sizes must be non-negative."
|
| 57 |
+
assert grid_dim > 0, f"Grid dimension must be positive (it's {grid_dim})."
|
| 58 |
+
num_m_tiles = (group_sizes + block_size_m - 1) // block_size_m
|
| 59 |
+
assert torch.all(num_m_tiles >= 0).item(), "All num_m_tiles must be non-negative."
|
| 60 |
+
num_n_tiles = triton.cdiv(N, block_size_n)
|
| 61 |
+
assert num_n_tiles > 0, f"num_n_tiles must be positive, it's {num_n_tiles}."
|
| 62 |
+
num_tiles = torch.sum(num_m_tiles * num_n_tiles).item()
|
| 63 |
+
assert num_tiles > 0, f"num_tiles must be positive, it's {num_tiles}."
|
| 64 |
+
num_programs = int(min(grid_dim, num_tiles))
|
| 65 |
+
assert num_programs > 0, f"num_programs must be positive, it's {num_programs}."
|
| 66 |
+
return (num_programs,)
|
| 67 |
+
|
| 68 |
+
|
| 69 |
+
def gmm(
|
| 70 |
+
lhs: Tensor,
|
| 71 |
+
rhs: Tensor,
|
| 72 |
+
group_sizes: Tensor,
|
| 73 |
+
preferred_element_type: torch.dtype = DTYPE,
|
| 74 |
+
existing_out: Tensor | None = None,
|
| 75 |
+
config: dict[str, int] | None = None,
|
| 76 |
+
bias: Tensor | None = None,
|
| 77 |
+
) -> Tensor:
|
| 78 |
+
"""
|
| 79 |
+
Perform Group Matrix Multiplication (GMM): out = lhs @ rhs + bias
|
| 80 |
+
|
| 81 |
+
lhs rows are divided into G groups. Each group of lhs rows is matrix multiplied with a plane of
|
| 82 |
+
rhs 3D tensor and then stored in a slice of out. In PyTorch parlance, it can be implemented as
|
| 83 |
+
follows for a given group g:
|
| 84 |
+
out[group_start:group_end, :] = lhs[group_start:group_end, :] @ rhs[g] + bias[g]
|
| 85 |
+
|
| 86 |
+
The size of each group, and their respective start and end positions are specified by
|
| 87 |
+
group_sizes tensor. For instance, suppose that group_sizes = [3, 2, 4, 1]. In this particular
|
| 88 |
+
case we have 4 groups. The 1st group starts at 0 and ends at 2, the second group starts at 3 and
|
| 89 |
+
ends at 4, the third group starts at 5 and ends at 8, and the fourth and final group consists of
|
| 90 |
+
just the 10th (last) row of lhs.
|
| 91 |
+
|
| 92 |
+
Parameters
|
| 93 |
+
----------
|
| 94 |
+
lhs : torch.Tensor
|
| 95 |
+
Left-hand side 2D input tensor. Shape: (M, K).
|
| 96 |
+
lhs data type must be torch.float16 or torch.bfloat16, and must match rhs data type.
|
| 97 |
+
lhs must be on the same device of rhs and group_sizes.
|
| 98 |
+
rhs : torch.Tensor
|
| 99 |
+
Right-hand side 3D input tensor. Shape: (G, K, N).
|
| 100 |
+
rhs data type must be torch.float16 or torch.bfloat16, and must match lhs data type.
|
| 101 |
+
rhs must be on the same device of lhs and group_sizes.
|
| 102 |
+
group_sizes : torch.Tensor
|
| 103 |
+
1D input tensor describing group sizes. Shape: (G,).
|
| 104 |
+
group_sizes data type must be torch.int32 and all its elements must be non-negative.
|
| 105 |
+
group_sizes must be on the same device of lhs and rhs.
|
| 106 |
+
preferred_element_type : torch.dtype, optional
|
| 107 |
+
Desired data type for output tensor. Default is torch.bfloat16.
|
| 108 |
+
Supported output types are torch.float16 and torch.bfloat16.
|
| 109 |
+
existing_out : torch.Tensor or None, optional
|
| 110 |
+
Preallocated output tensor. Default is None.
|
| 111 |
+
If provided, results are written into this tensor. Otherwise, a new output tensor is
|
| 112 |
+
allocated.
|
| 113 |
+
If provided then it must have shape (M, N), its data type must match preferred_element_type
|
| 114 |
+
and it must be on the same device of other input tensors.
|
| 115 |
+
config : dict[str, int] or None, optional
|
| 116 |
+
Optional dictionary with kernel metaparameters. If absent, config will be queried from
|
| 117 |
+
internal tuning database.
|
| 118 |
+
bias : torch.Tensor or None, optional
|
| 119 |
+
Optional bias tensor. Shape: (G, N).
|
| 120 |
+
If provided, bias data type must match lhs and rhs data type, and bias must be on the same
|
| 121 |
+
device as other input tensors. Each group g adds bias[g] to the output.
|
| 122 |
+
|
| 123 |
+
Returns
|
| 124 |
+
-------
|
| 125 |
+
torch.Tensor
|
| 126 |
+
The computed output 2D tensor. Shape: (M, N).
|
| 127 |
+
Output tensor data type is given by preferred_element_type.
|
| 128 |
+
If existing_out is provided then existing_out is also returned.
|
| 129 |
+
|
| 130 |
+
Implementation Notes
|
| 131 |
+
--------------------
|
| 132 |
+
- GMM is implemented with a persistent Triton kernel.
|
| 133 |
+
- lhs must be row-major (lhs.stride() == (K, 1)).
|
| 134 |
+
- rhs can be row-major (rhs.stride() == (K * N, N, 1)) or column-major (rhs.stride() ==
|
| 135 |
+
(K * N, 1, K)). If rhs is row-major then kernel parameter TRANS_RHS == False, this is useful
|
| 136 |
+
for implementing forward pass. If rhs is column-major then kernel parameter TRANS_RHS == True,
|
| 137 |
+
this is useful for computing the lhs derivative in the backward pass, while fusing the
|
| 138 |
+
transposition.
|
| 139 |
+
- out must be row-major (out.stride() == (N, 1)).
|
| 140 |
+
- bias must be row-major (bias.stride() == (N, 1)) if provided.
|
| 141 |
+
"""
|
| 142 |
+
use_bias = bias is not None
|
| 143 |
+
check_input_device_dtype(lhs, rhs, group_sizes, bias)
|
| 144 |
+
|
| 145 |
+
M, K, N, G = get_gmm_shape(lhs, rhs, group_sizes)
|
| 146 |
+
|
| 147 |
+
if use_bias:
|
| 148 |
+
check_bias_shape_stride(bias, G, N)
|
| 149 |
+
|
| 150 |
+
out = get_gmm_output(
|
| 151 |
+
M,
|
| 152 |
+
N,
|
| 153 |
+
device=lhs.device,
|
| 154 |
+
preferred_element_type=preferred_element_type,
|
| 155 |
+
existing_out=existing_out,
|
| 156 |
+
)
|
| 157 |
+
|
| 158 |
+
trans_rhs, _ = get_gmm_transposition(lhs, rhs, out)
|
| 159 |
+
|
| 160 |
+
if config is None:
|
| 161 |
+
config = get_config("gmm", M, K, N, G)
|
| 162 |
+
|
| 163 |
+
assert all(
|
| 164 |
+
key in config
|
| 165 |
+
and isinstance(config[key], int)
|
| 166 |
+
and (
|
| 167 |
+
is_power_of_2(config[key])
|
| 168 |
+
if key.startswith("BLOCK_SIZE_")
|
| 169 |
+
else config[key] > 0
|
| 170 |
+
)
|
| 171 |
+
for key in {
|
| 172 |
+
"BLOCK_SIZE_M",
|
| 173 |
+
"BLOCK_SIZE_K",
|
| 174 |
+
"BLOCK_SIZE_N",
|
| 175 |
+
"GROUP_SIZE",
|
| 176 |
+
"GRID_DIM",
|
| 177 |
+
}
|
| 178 |
+
), "Invalid GMM kernel config."
|
| 179 |
+
|
| 180 |
+
grid = _gmm_grid(
|
| 181 |
+
N,
|
| 182 |
+
config["BLOCK_SIZE_M"],
|
| 183 |
+
config["BLOCK_SIZE_N"],
|
| 184 |
+
group_sizes,
|
| 185 |
+
config["GRID_DIM"],
|
| 186 |
+
)
|
| 187 |
+
|
| 188 |
+
# fmt: off
|
| 189 |
+
gmm_kernel[grid](
|
| 190 |
+
# Tensor pointers:
|
| 191 |
+
lhs, rhs, group_sizes, out, bias,
|
| 192 |
+
# Tensor shapes:
|
| 193 |
+
M, K, N, G,
|
| 194 |
+
# Meta-parameters:
|
| 195 |
+
TRANS_RHS=trans_rhs,
|
| 196 |
+
USE_BIAS=use_bias,
|
| 197 |
+
**config,
|
| 198 |
+
)
|
| 199 |
+
# fmt: on
|
| 200 |
+
|
| 201 |
+
return out
|
| 202 |
+
|
| 203 |
+
|
| 204 |
+
# Persistent TGMM PyTorch wrapper.
|
| 205 |
+
# ------------------------------------------------------------------------------
|
| 206 |
+
|
| 207 |
+
|
| 208 |
+
def _ptgmm_grid(
|
| 209 |
+
K: int,
|
| 210 |
+
N: int,
|
| 211 |
+
G: int,
|
| 212 |
+
block_size_k: int,
|
| 213 |
+
block_size_n: int,
|
| 214 |
+
grid_dim: int,
|
| 215 |
+
) -> tuple[int]:
|
| 216 |
+
assert K > 0, f"K must be positive, it's {K}."
|
| 217 |
+
assert N > 0, f"N must be positive, it's {N}."
|
| 218 |
+
assert G > 0, f"G must be positive, it's {G}."
|
| 219 |
+
assert is_power_of_2(
|
| 220 |
+
block_size_k
|
| 221 |
+
), f"K-dimension tile size must be a power of 2 (it's {block_size_k})."
|
| 222 |
+
assert is_power_of_2(
|
| 223 |
+
block_size_n
|
| 224 |
+
), f"N-dimension tile size must be a power of 2 (it's {block_size_n})."
|
| 225 |
+
assert grid_dim > 0, f"Grid dimension must be positive (it's {grid_dim})."
|
| 226 |
+
num_k_tiles = triton.cdiv(K, block_size_k)
|
| 227 |
+
assert num_k_tiles > 0, f"num_k_tiles must be positive, it's {num_k_tiles}."
|
| 228 |
+
num_n_tiles = triton.cdiv(N, block_size_n)
|
| 229 |
+
assert num_n_tiles > 0, f"num_n_tiles must be positive, it's {num_n_tiles}."
|
| 230 |
+
num_tiles = G * num_k_tiles * num_n_tiles
|
| 231 |
+
assert num_tiles > 0, f"num_tiles must be positive, it's {num_tiles}."
|
| 232 |
+
num_programs = min(grid_dim, num_tiles)
|
| 233 |
+
assert num_programs > 0, f"num_programs must be positive, it's {num_programs}."
|
| 234 |
+
return (num_programs,)
|
| 235 |
+
|
| 236 |
+
|
| 237 |
+
def ptgmm(
|
| 238 |
+
lhs: Tensor,
|
| 239 |
+
rhs: Tensor,
|
| 240 |
+
group_sizes: Tensor,
|
| 241 |
+
preferred_element_type: torch.dtype = DTYPE,
|
| 242 |
+
existing_out: Tensor | None = None,
|
| 243 |
+
config: dict[str, int] | None = None,
|
| 244 |
+
bias_grad: Tensor | None = None,
|
| 245 |
+
accumulate: bool = False,
|
| 246 |
+
) -> Tensor:
|
| 247 |
+
"""
|
| 248 |
+
Perform a Group Matrix Multiplication (GMM) variant: out = lhs @ rhs
|
| 249 |
+
|
| 250 |
+
lhs columns and rhs rows are divided into G groups. Each group of lhs is matrix multiplied with
|
| 251 |
+
the respective group of rhs and then stored in a plane of the output 3D tensor. In PyTorch
|
| 252 |
+
parlance, it can be implemented as follows for a given group g:
|
| 253 |
+
out[g] = lhs[:, group_start:group_end] @ rhs[group_start:group_end, :]
|
| 254 |
+
|
| 255 |
+
The 't' in the operator name derives from MaxText implementation
|
| 256 |
+
(https://github.com/AI-Hypercomputer/maxtext/blob/main/src/MaxText/kernels/megablox/gmm.py),
|
| 257 |
+
which served as the initial inspiration for this one. TGMM differs from GMM in terms of tensor
|
| 258 |
+
shapes. GMM does (M, K) @ (G, K, N) = (M, N) while TGMM does (K, M) @ (M, N) = (G, K, N).
|
| 259 |
+
|
| 260 |
+
The 'p' in the operator name means that it is implemented with a persistent kernel. There is
|
| 261 |
+
also the non-persistent variation, which is implemented with a regular kernel. Please take a
|
| 262 |
+
look at nptgmm operator. Both ptgmm and nptgmm implement the same computation, choosing one or
|
| 263 |
+
the other is a matter of performance for the target workload.
|
| 264 |
+
|
| 265 |
+
Parameters
|
| 266 |
+
----------
|
| 267 |
+
lhs : torch.Tensor
|
| 268 |
+
Left-hand side 2D input tensor. Shape: (K, M).
|
| 269 |
+
lhs data type must be torch.float16 or torch.bfloat16, and must match rhs data type.
|
| 270 |
+
lhs must be on the same device of rhs and group_sizes.
|
| 271 |
+
rhs : torch.Tensor
|
| 272 |
+
Right-hand side 2D input tensor. Shape: (M, N).
|
| 273 |
+
rhs data type must be torch.float16 or torch.bfloat16, and must match lhs data type.
|
| 274 |
+
rhs must be on the same device of lhs and group_sizes.
|
| 275 |
+
group_sizes : torch.Tensor
|
| 276 |
+
1D input tensor describing group sizes. Shape: (G,).
|
| 277 |
+
group_sizes data type must be torch.int32 and all its elements must be non-negative.
|
| 278 |
+
group_sizes must be on the same device of lhs and rhs.
|
| 279 |
+
preferred_element_type : torch.dtype, optional
|
| 280 |
+
Desired data type for output tensor. Default is torch.bfloat16.
|
| 281 |
+
Supported output types are torch.float16 and torch.bfloat16.
|
| 282 |
+
existing_out : torch.Tensor or None, optional
|
| 283 |
+
Preallocated output tensor. Default is None.
|
| 284 |
+
If provided, results are written into this tensor. Otherwise, a new output tensor is
|
| 285 |
+
allocated.
|
| 286 |
+
If provided then it must have shape (G, K, N), its data type must match
|
| 287 |
+
preferred_element_type and it must be on the same device of other input tensors.
|
| 288 |
+
config : dict[str, int] or None, optional
|
| 289 |
+
Optional dictionary with kernel metaparameters. If absent, config will be queried from
|
| 290 |
+
internal tuning database.
|
| 291 |
+
bias_grad : torch.Tensor or None, optional
|
| 292 |
+
Optional bias gradient output tensor. Shape: (G, K).
|
| 293 |
+
If provided, the kernel will compute the bias gradient and write it to this tensor.
|
| 294 |
+
bias_grad must be torch.float32 (kernel uses atomic_add which requires float32),
|
| 295 |
+
accumulate : bool, optional
|
| 296 |
+
Whether to accumulate into existing output tensor values. Default is False.
|
| 297 |
+
If False, output will be overwritten with fresh computation.
|
| 298 |
+
If True, results will be added to existing output tensor values.
|
| 299 |
+
|
| 300 |
+
Returns
|
| 301 |
+
-------
|
| 302 |
+
torch.Tensor
|
| 303 |
+
The computed output 3D tensor. Shape: (G, K, N).
|
| 304 |
+
Output tensor data type is given by preferred_element_type.
|
| 305 |
+
If existing_out is provided then existing_out is also returned.
|
| 306 |
+
|
| 307 |
+
Implementation Notes
|
| 308 |
+
--------------------
|
| 309 |
+
- PTGMM is implemented with a persistent Triton kernel.
|
| 310 |
+
- lhs can be row-major (lhs.stride() == (M, 1)) or column-major (lhs.stride() == (1, K)). If lhs
|
| 311 |
+
is row-major then kernel parameter TRANS_LHS == False. If lhs is column-major then kernel
|
| 312 |
+
parameter TRANS_LHS == True, this is useful for computing the rhs derivative in the backward
|
| 313 |
+
pass, while fusing the transposition.
|
| 314 |
+
- rhs must be row-major (rhs.stride() == (N, 1)).
|
| 315 |
+
- out must be row-major (out.stride() == (K * N, N, 1)).
|
| 316 |
+
"""
|
| 317 |
+
check_input_device_dtype(lhs, rhs, group_sizes)
|
| 318 |
+
|
| 319 |
+
M, K, N, G = get_tgmm_shape(lhs, rhs, group_sizes)
|
| 320 |
+
|
| 321 |
+
out = get_tgmm_output(
|
| 322 |
+
K,
|
| 323 |
+
N,
|
| 324 |
+
G,
|
| 325 |
+
device=lhs.device,
|
| 326 |
+
preferred_element_type=preferred_element_type,
|
| 327 |
+
existing_out=existing_out,
|
| 328 |
+
)
|
| 329 |
+
|
| 330 |
+
trans_lhs, _ = get_tgmm_transposition(lhs, rhs, out)
|
| 331 |
+
|
| 332 |
+
if config is None:
|
| 333 |
+
config = get_config("ptgmm", M, K, N, G, accumulate)
|
| 334 |
+
|
| 335 |
+
assert all(
|
| 336 |
+
key in config
|
| 337 |
+
and isinstance(config[key], int)
|
| 338 |
+
and (
|
| 339 |
+
is_power_of_2(config[key])
|
| 340 |
+
if key.startswith("BLOCK_SIZE_")
|
| 341 |
+
else config[key] > 0
|
| 342 |
+
)
|
| 343 |
+
for key in {
|
| 344 |
+
"BLOCK_SIZE_M",
|
| 345 |
+
"BLOCK_SIZE_K",
|
| 346 |
+
"BLOCK_SIZE_N",
|
| 347 |
+
"GROUP_SIZE",
|
| 348 |
+
"GRID_DIM",
|
| 349 |
+
}
|
| 350 |
+
), "Invalid PTGMM kernel config."
|
| 351 |
+
|
| 352 |
+
# Bias gradient handling.
|
| 353 |
+
# -----------------------
|
| 354 |
+
# Get or validate bias gradient tensor.
|
| 355 |
+
compute_bias_grad = bias_grad is not None
|
| 356 |
+
bias_grad_ptr = get_tgmm_bias_grad(
|
| 357 |
+
K,
|
| 358 |
+
G,
|
| 359 |
+
device=lhs.device,
|
| 360 |
+
existing_bias_grad=bias_grad,
|
| 361 |
+
)
|
| 362 |
+
|
| 363 |
+
grid = _ptgmm_grid(
|
| 364 |
+
K,
|
| 365 |
+
N,
|
| 366 |
+
G,
|
| 367 |
+
config["BLOCK_SIZE_K"],
|
| 368 |
+
config["BLOCK_SIZE_N"],
|
| 369 |
+
config["GRID_DIM"],
|
| 370 |
+
)
|
| 371 |
+
|
| 372 |
+
# fmt: off
|
| 373 |
+
tgmm_persistent_kernel[grid](
|
| 374 |
+
# Tensor pointers:
|
| 375 |
+
lhs, rhs, group_sizes, out, bias_grad_ptr,
|
| 376 |
+
# Tensor shapes:
|
| 377 |
+
M, K, N, G,
|
| 378 |
+
# Meta-parameters:
|
| 379 |
+
TRANS_LHS=trans_lhs,
|
| 380 |
+
COMPUTE_BIAS_GRAD=compute_bias_grad,
|
| 381 |
+
ACCUMULATE=accumulate,
|
| 382 |
+
**config,
|
| 383 |
+
)
|
| 384 |
+
# fmt: on
|
| 385 |
+
|
| 386 |
+
return out
|
| 387 |
+
|
| 388 |
+
|
| 389 |
+
# Regular non-persistent TGMM PyTorch wrapper.
|
| 390 |
+
# ------------------------------------------------------------------------------
|
| 391 |
+
|
| 392 |
+
|
| 393 |
+
def _nptgmm_grid(
|
| 394 |
+
K: int,
|
| 395 |
+
N: int,
|
| 396 |
+
G: int,
|
| 397 |
+
block_size_k: int,
|
| 398 |
+
block_size_n: int,
|
| 399 |
+
) -> tuple[int, int]:
|
| 400 |
+
assert K > 0, f"K must be positive, it's {K}."
|
| 401 |
+
assert N > 0, f"N must be positive, it's {N}."
|
| 402 |
+
assert G > 0, f"G must be positive, it's {G}."
|
| 403 |
+
assert is_power_of_2(
|
| 404 |
+
block_size_k
|
| 405 |
+
), f"K-dimension tile size must be a power of 2 (it's {block_size_k})."
|
| 406 |
+
assert is_power_of_2(
|
| 407 |
+
block_size_n
|
| 408 |
+
), f"N-dimension tile size must be a power of 2 (it's {block_size_n})."
|
| 409 |
+
num_k_tiles = triton.cdiv(K, block_size_k)
|
| 410 |
+
assert num_k_tiles > 0, f"num_k_tiles must be positive, it's {num_k_tiles}."
|
| 411 |
+
num_n_tiles = triton.cdiv(N, block_size_n)
|
| 412 |
+
assert num_n_tiles > 0, f"num_n_tiles must be positive, it's {num_n_tiles}."
|
| 413 |
+
num_tiles_per_mm = num_k_tiles * num_n_tiles
|
| 414 |
+
assert (
|
| 415 |
+
num_tiles_per_mm > 0
|
| 416 |
+
), f"num_tiles_per_mm must be positive, it's {num_tiles_per_mm}."
|
| 417 |
+
return (G, num_tiles_per_mm)
|
| 418 |
+
|
| 419 |
+
|
| 420 |
+
def nptgmm(
|
| 421 |
+
lhs: Tensor,
|
| 422 |
+
rhs: Tensor,
|
| 423 |
+
group_sizes: Tensor,
|
| 424 |
+
preferred_element_type: torch.dtype = DTYPE,
|
| 425 |
+
existing_out: Tensor | None = None,
|
| 426 |
+
config: dict[str, int] | None = None,
|
| 427 |
+
bias_grad: Tensor | None = None,
|
| 428 |
+
accumulate: bool = False,
|
| 429 |
+
) -> Tensor:
|
| 430 |
+
"""
|
| 431 |
+
Perform a Group Matrix Multiplication (GMM) variant: out = lhs @ rhs
|
| 432 |
+
|
| 433 |
+
lhs columns and rhs rows are divided into G groups. Each group of lhs is matrix multiplied with
|
| 434 |
+
the respective group of rhs and then stored in a plane of the output 3D tensor. In PyTorch
|
| 435 |
+
parlance, it can be implemented as follows for a given group g:
|
| 436 |
+
out[g] = lhs[:, group_start:group_end] @ rhs[group_start:group_end, :]
|
| 437 |
+
|
| 438 |
+
The 't' in the operator name derives from MaxText implementation
|
| 439 |
+
(https://github.com/AI-Hypercomputer/maxtext/blob/main/src/MaxText/kernels/megablox/gmm.py),
|
| 440 |
+
which served as the initial inspiration for this one. TGMM differs from GMM in terms of tensor
|
| 441 |
+
shapes. GMM does (M, K) @ (G, K, N) = (M, N) while TGMM does (K, M) @ (M, N) = (G, K, N).
|
| 442 |
+
|
| 443 |
+
The 'np' in the operator name means that it is implemented with a non-persistent, i.e. regular
|
| 444 |
+
kernel. There is also the persistent variation, which is implemented with a persistent kernel.
|
| 445 |
+
Please take a look at ptgmm operator. Both nptgmm and ptgmm implement the same computation,
|
| 446 |
+
choosing one or the other is a matter of performance for the target workload.
|
| 447 |
+
|
| 448 |
+
Parameters
|
| 449 |
+
----------
|
| 450 |
+
lhs : torch.Tensor
|
| 451 |
+
Left-hand side 2D input tensor. Shape: (K, M).
|
| 452 |
+
lhs data type must be torch.float16 or torch.bfloat16, and must match rhs data type.
|
| 453 |
+
lhs must be on the same device of rhs and group_sizes.
|
| 454 |
+
rhs : torch.Tensor
|
| 455 |
+
Right-hand side 2D input tensor. Shape: (M, N).
|
| 456 |
+
rhs data type must be torch.float16 or torch.bfloat16, and must match lhs data type.
|
| 457 |
+
rhs must be on the same device of lhs and group_sizes.
|
| 458 |
+
group_sizes : torch.Tensor
|
| 459 |
+
1D input tensor describing group sizes. Shape: (G,).
|
| 460 |
+
group_sizes data type must be torch.int32 and all its elements must be non-negative.
|
| 461 |
+
group_sizes must be on the same device of lhs and rhs.
|
| 462 |
+
preferred_element_type : torch.dtype, optional
|
| 463 |
+
Desired data type for output tensor. Default is torch.bfloat16.
|
| 464 |
+
Supported output types are torch.float16 and torch.bfloat16.
|
| 465 |
+
existing_out : torch.Tensor or None, optional
|
| 466 |
+
Preallocated output tensor. Default is None.
|
| 467 |
+
If provided, results are written into this tensor. Otherwise, a new output tensor is
|
| 468 |
+
allocated.
|
| 469 |
+
If provided then it must have shape (G, K, N), its data type must match
|
| 470 |
+
preferred_element_type and it must be on the same device of other input tensors.
|
| 471 |
+
config : dict[str, int] or None, optional
|
| 472 |
+
Optional dictionary with kernel metaparameters. If absent, config will be queried from
|
| 473 |
+
internal tuning database.
|
| 474 |
+
bias_grad : torch.Tensor or None, optional
|
| 475 |
+
Optional bias gradient output tensor. Shape: (G, K).
|
| 476 |
+
If provided, the kernel will compute the bias gradient and write it to this tensor.
|
| 477 |
+
bias_grad must be torch.float32 (kernel uses atomic_add which requires float32),
|
| 478 |
+
accumulate : bool, optional
|
| 479 |
+
Whether to accumulate into existing output tensor values. Default is False.
|
| 480 |
+
If False, output will be overwritten with fresh computation.
|
| 481 |
+
If True, results will be added to existing output tensor values.
|
| 482 |
+
|
| 483 |
+
Returns
|
| 484 |
+
-------
|
| 485 |
+
torch.Tensor
|
| 486 |
+
The computed output 3D tensor. Shape: (G, K, N).
|
| 487 |
+
Output tensor data type is given by preferred_element_type.
|
| 488 |
+
If existing_out is provided then existing_out is also returned.
|
| 489 |
+
|
| 490 |
+
Implementation Notes
|
| 491 |
+
--------------------
|
| 492 |
+
- NPTGMM is implemented with a non-persistent regular Triton kernel.
|
| 493 |
+
- lhs can be row-major (lhs.stride() == (M, 1)) or column-major (lhs.stride() == (1, K)). If lhs
|
| 494 |
+
is row-major then kernel parameter TRANS_LHS == False. If lhs is column-major then kernel
|
| 495 |
+
parameter TRANS_LHS == True, this is useful for computing the rhs derivative in the backward
|
| 496 |
+
pass, while fusing the transposition.
|
| 497 |
+
- rhs must be row-major (rhs.stride() == (N, 1)).
|
| 498 |
+
- out must be row-major (out.stride() == (K * N, N, 1)).
|
| 499 |
+
"""
|
| 500 |
+
check_input_device_dtype(lhs, rhs, group_sizes)
|
| 501 |
+
|
| 502 |
+
M, K, N, G = get_tgmm_shape(lhs, rhs, group_sizes)
|
| 503 |
+
|
| 504 |
+
out = get_tgmm_output(
|
| 505 |
+
K,
|
| 506 |
+
N,
|
| 507 |
+
G,
|
| 508 |
+
device=lhs.device,
|
| 509 |
+
preferred_element_type=preferred_element_type,
|
| 510 |
+
existing_out=existing_out,
|
| 511 |
+
)
|
| 512 |
+
|
| 513 |
+
trans_lhs, _ = get_tgmm_transposition(lhs, rhs, out)
|
| 514 |
+
|
| 515 |
+
# Bias gradient handling.
|
| 516 |
+
# -----------------------
|
| 517 |
+
# Get or validate bias gradient tensor.
|
| 518 |
+
compute_bias_grad = bias_grad is not None
|
| 519 |
+
bias_grad_ptr = get_tgmm_bias_grad(
|
| 520 |
+
K,
|
| 521 |
+
G,
|
| 522 |
+
device=lhs.device,
|
| 523 |
+
existing_bias_grad=bias_grad,
|
| 524 |
+
)
|
| 525 |
+
|
| 526 |
+
if config is None:
|
| 527 |
+
config = get_config("nptgmm", M, K, N, G, accumulate)
|
| 528 |
+
|
| 529 |
+
assert all(
|
| 530 |
+
key in config
|
| 531 |
+
and isinstance(config[key], int)
|
| 532 |
+
and (
|
| 533 |
+
is_power_of_2(config[key])
|
| 534 |
+
if key.startswith("BLOCK_SIZE_")
|
| 535 |
+
else config[key] > 0
|
| 536 |
+
)
|
| 537 |
+
for key in {
|
| 538 |
+
"BLOCK_SIZE_M",
|
| 539 |
+
"BLOCK_SIZE_K",
|
| 540 |
+
"BLOCK_SIZE_N",
|
| 541 |
+
"GROUP_SIZE",
|
| 542 |
+
}
|
| 543 |
+
), "Invalid NPTGMM kernel config."
|
| 544 |
+
|
| 545 |
+
grid = _nptgmm_grid(
|
| 546 |
+
K,
|
| 547 |
+
N,
|
| 548 |
+
G,
|
| 549 |
+
config["BLOCK_SIZE_K"],
|
| 550 |
+
config["BLOCK_SIZE_N"],
|
| 551 |
+
)
|
| 552 |
+
|
| 553 |
+
# fmt: off
|
| 554 |
+
tgmm_non_persistent_kernel[grid](
|
| 555 |
+
# Tensor pointers:
|
| 556 |
+
lhs, rhs, group_sizes, out, bias_grad_ptr,
|
| 557 |
+
# Tensor shapes:
|
| 558 |
+
M, K, N, G,
|
| 559 |
+
# Meta-parameters:
|
| 560 |
+
TRANS_LHS=trans_lhs,
|
| 561 |
+
COMPUTE_BIAS_GRAD=compute_bias_grad,
|
| 562 |
+
ACCUMULATE=accumulate,
|
| 563 |
+
**config,
|
| 564 |
+
)
|
| 565 |
+
# fmt: on
|
| 566 |
+
|
| 567 |
+
return out
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/__init__.py
ADDED
|
File without changes
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/_triton/__init__.py
ADDED
|
File without changes
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/_triton/arch_info.py
ADDED
|
@@ -0,0 +1,46 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
import triton
|
| 2 |
+
|
| 3 |
+
# Detect the GPU arch lazily: querying the triton driver at import time fails
|
| 4 |
+
# in headless environments (e.g. the kernel-builder ABI check sandbox has no
|
| 5 |
+
# GPU), and the original JAX fallback pulled in an unrelated runtime dep. The
|
| 6 |
+
# arch is only actually needed when a GMM kernel is dispatched, so resolve and
|
| 7 |
+
# cache on first call.
|
| 8 |
+
_CACHED_ARCH = None
|
| 9 |
+
|
| 10 |
+
|
| 11 |
+
def get_arch():
|
| 12 |
+
global _CACHED_ARCH
|
| 13 |
+
if _CACHED_ARCH is not None:
|
| 14 |
+
return _CACHED_ARCH
|
| 15 |
+
try:
|
| 16 |
+
_CACHED_ARCH = triton.runtime.driver.active.get_current_target().arch
|
| 17 |
+
except RuntimeError:
|
| 18 |
+
try:
|
| 19 |
+
from jax._src.lib import gpu_triton as triton_kernel_call_lib
|
| 20 |
+
_CACHED_ARCH = triton_kernel_call_lib.get_arch_details("0").split(":")[0]
|
| 21 |
+
except ImportError as e:
|
| 22 |
+
raise RuntimeError(
|
| 23 |
+
"Cannot determine GPU arch: triton driver is inactive and "
|
| 24 |
+
"JAX is not available. A GPU is required for grouped GEMM."
|
| 25 |
+
) from e
|
| 26 |
+
return _CACHED_ARCH
|
| 27 |
+
|
| 28 |
+
|
| 29 |
+
def is_gluon_avail():
|
| 30 |
+
return get_arch() in ("gfx950", "gfx1250")
|
| 31 |
+
|
| 32 |
+
|
| 33 |
+
def is_fp4_avail():
|
| 34 |
+
return get_arch() in ("gfx950", "gfx1250")
|
| 35 |
+
|
| 36 |
+
|
| 37 |
+
def is_fp8_avail():
|
| 38 |
+
return get_arch() in ("gfx942", "gfx950", "gfx1250", "gfx1200", "gfx1201")
|
| 39 |
+
|
| 40 |
+
|
| 41 |
+
def is_mx_scale_preshuffling_avail():
|
| 42 |
+
return get_arch() in ("gfx950", "gfx1250")
|
| 43 |
+
|
| 44 |
+
|
| 45 |
+
def is_tdm_avail():
|
| 46 |
+
return get_arch() in ("gfx1250",)
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/_triton/pid_preprocessing.py
ADDED
|
@@ -0,0 +1,100 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# SPDX-License-Identifier: MIT
|
| 2 |
+
|
| 3 |
+
# Copyright (C) 2024-2026, Advanced Micro Devices, Inc. All rights reserved.
|
| 4 |
+
|
| 5 |
+
import triton
|
| 6 |
+
import triton.language as tl
|
| 7 |
+
|
| 8 |
+
|
| 9 |
+
@triton.jit
|
| 10 |
+
def remap_xcd_chunked(
|
| 11 |
+
pid, GRID_MN, NUM_XCDS: tl.constexpr = 8, CHUNK_SIZE: tl.constexpr = 2
|
| 12 |
+
):
|
| 13 |
+
# Compute current XCD and local PID
|
| 14 |
+
xcd = pid % NUM_XCDS
|
| 15 |
+
# distribute the modulo pids in round robin
|
| 16 |
+
if pid > (GRID_MN // (NUM_XCDS * CHUNK_SIZE)) * (NUM_XCDS * CHUNK_SIZE):
|
| 17 |
+
return pid
|
| 18 |
+
local_pid = pid // NUM_XCDS
|
| 19 |
+
# Calculate chunk index and position within chunk
|
| 20 |
+
chunk_idx = local_pid // CHUNK_SIZE
|
| 21 |
+
pos_in_chunk = local_pid % CHUNK_SIZE
|
| 22 |
+
# Calculate new PID
|
| 23 |
+
new_pid = chunk_idx * NUM_XCDS * CHUNK_SIZE + xcd * CHUNK_SIZE + pos_in_chunk
|
| 24 |
+
return new_pid
|
| 25 |
+
|
| 26 |
+
|
| 27 |
+
@triton.jit
|
| 28 |
+
def remap_xcd(pid, GRID_MN, NUM_XCDS: tl.constexpr = 8):
|
| 29 |
+
## pid remapping on xcds
|
| 30 |
+
# Number of pids per XCD in the new arrangement
|
| 31 |
+
pids_per_xcd = (GRID_MN + NUM_XCDS - 1) // NUM_XCDS
|
| 32 |
+
# When GRID_MN cannot divide NUM_XCDS, some xcds will have
|
| 33 |
+
# pids_per_xcd pids, the other will have pids_per_xcd - 1 pids.
|
| 34 |
+
# We calculate the number of xcds that have pids_per_xcd pids as
|
| 35 |
+
# tall_xcds
|
| 36 |
+
tall_xcds = GRID_MN % NUM_XCDS
|
| 37 |
+
tall_xcds = NUM_XCDS if tall_xcds == 0 else tall_xcds
|
| 38 |
+
# Compute current XCD and local pid within the XCD
|
| 39 |
+
xcd = pid % NUM_XCDS
|
| 40 |
+
local_pid = pid // NUM_XCDS
|
| 41 |
+
# Calculate new pid based on the new grouping
|
| 42 |
+
# Note that we need to consider the following two cases:
|
| 43 |
+
# 1. the current pid is on a tall xcd
|
| 44 |
+
# 2. the current pid is on a short xcd
|
| 45 |
+
if xcd < tall_xcds:
|
| 46 |
+
pid = xcd * pids_per_xcd + local_pid
|
| 47 |
+
else:
|
| 48 |
+
pid = (
|
| 49 |
+
tall_xcds * pids_per_xcd
|
| 50 |
+
+ (xcd - tall_xcds) * (pids_per_xcd - 1)
|
| 51 |
+
+ local_pid
|
| 52 |
+
)
|
| 53 |
+
|
| 54 |
+
return pid
|
| 55 |
+
|
| 56 |
+
|
| 57 |
+
@triton.jit
|
| 58 |
+
def pid_grid(pid: int, num_pid_m: int, num_pid_n: int, GROUP_SIZE_M: tl.constexpr = 1):
|
| 59 |
+
"""
|
| 60 |
+
Maps 1D pid to 2D grid coords (pid_m, pid_n).
|
| 61 |
+
|
| 62 |
+
Args:
|
| 63 |
+
- pid: 1D pid
|
| 64 |
+
- num_pid_m: grid m size
|
| 65 |
+
- num_pid_n: grid n size
|
| 66 |
+
- GROUP_SIZE_M: tl.constexpr: default is 1
|
| 67 |
+
"""
|
| 68 |
+
if GROUP_SIZE_M == 1:
|
| 69 |
+
pid_m = pid // num_pid_n
|
| 70 |
+
pid_n = pid % num_pid_n
|
| 71 |
+
else:
|
| 72 |
+
num_pid_in_group = GROUP_SIZE_M * num_pid_n
|
| 73 |
+
group_id = pid // num_pid_in_group
|
| 74 |
+
first_pid_m = group_id * GROUP_SIZE_M
|
| 75 |
+
group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
|
| 76 |
+
tl.assume(group_size_m >= 0)
|
| 77 |
+
pid_m = first_pid_m + (pid % group_size_m)
|
| 78 |
+
pid_n = (pid % num_pid_in_group) // group_size_m
|
| 79 |
+
|
| 80 |
+
return pid_m, pid_n
|
| 81 |
+
|
| 82 |
+
|
| 83 |
+
@triton.jit
|
| 84 |
+
def pid_grid_3d(pid: int, num_pid_m: int, num_pid_n: int, num_pid_k):
|
| 85 |
+
"""
|
| 86 |
+
Maps 1D pid to 3D grid coords (pid_m, pid_n, pid_k).
|
| 87 |
+
Args:
|
| 88 |
+
- pid: 1D pid
|
| 89 |
+
- num_pid_m: grid m size
|
| 90 |
+
- num_pid_n: grid n size
|
| 91 |
+
- num_pid_k: grid k size
|
| 92 |
+
|
| 93 |
+
Returns:
|
| 94 |
+
- pid_m, pid_n, pid_k: 3D grid coordinates
|
| 95 |
+
"""
|
| 96 |
+
pid_m = pid % num_pid_m
|
| 97 |
+
pid_n = (pid // num_pid_m) % num_pid_n
|
| 98 |
+
pid_k = pid // (num_pid_m * num_pid_n) % num_pid_k
|
| 99 |
+
|
| 100 |
+
return pid_m, pid_n, pid_k
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/gmm_common.py
ADDED
|
@@ -0,0 +1,752 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# SPDX-License-Identifier: MIT
|
| 2 |
+
# Copyright (C) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
| 3 |
+
|
| 4 |
+
# Imports.
|
| 5 |
+
# ------------------------------------------------------------------------------
|
| 6 |
+
|
| 7 |
+
# PyTorch
|
| 8 |
+
import torch
|
| 9 |
+
from torch import Tensor
|
| 10 |
+
|
| 11 |
+
# AITER: logging
|
| 12 |
+
from .logger import AiterTritonLogger
|
| 13 |
+
|
| 14 |
+
_LOGGER: AiterTritonLogger = AiterTritonLogger()
|
| 15 |
+
|
| 16 |
+
|
| 17 |
+
# Supported data types.
|
| 18 |
+
# ------------------------------------------------------------------------------
|
| 19 |
+
|
| 20 |
+
# Supported data types, as strings.
|
| 21 |
+
SUPPORTED_DTYPES_STR: set[str] = {"fp16", "bf16"}
|
| 22 |
+
|
| 23 |
+
|
| 24 |
+
# Convert string data type to PyTorch data type.
|
| 25 |
+
def dtype_from_str(dtype_str: str) -> torch.dtype:
|
| 26 |
+
dtype_str = dtype_str.strip().lower()
|
| 27 |
+
dtype_str = dtype_str[1:] if dtype_str[0] in {"i", "o"} else dtype_str
|
| 28 |
+
assert (
|
| 29 |
+
dtype_str in SUPPORTED_DTYPES_STR
|
| 30 |
+
), "String data type isn't in set of supported string data types."
|
| 31 |
+
return {"fp16": torch.float16, "bf16": torch.bfloat16}[dtype_str]
|
| 32 |
+
|
| 33 |
+
|
| 34 |
+
# Supported data types, as PyTorch types.
|
| 35 |
+
SUPPORTED_DTYPES: set[torch.dtype] = {
|
| 36 |
+
dtype_from_str(dtype_str) for dtype_str in SUPPORTED_DTYPES_STR
|
| 37 |
+
}
|
| 38 |
+
|
| 39 |
+
|
| 40 |
+
# Convert PyTorch data type to string data type.
|
| 41 |
+
def str_from_dtype(dtype: torch.dtype) -> str:
|
| 42 |
+
assert (
|
| 43 |
+
dtype in SUPPORTED_DTYPES
|
| 44 |
+
), "PyTorch data type isn't in set of supported PyTorch data types."
|
| 45 |
+
return {torch.float16: "fp16", torch.bfloat16: "bf16"}[dtype]
|
| 46 |
+
|
| 47 |
+
|
| 48 |
+
# Default data type, as string.
|
| 49 |
+
DTYPE_STR: str = "bf16"
|
| 50 |
+
assert (
|
| 51 |
+
DTYPE_STR in SUPPORTED_DTYPES_STR
|
| 52 |
+
), "Default string data type isn't in set of supported string data types."
|
| 53 |
+
|
| 54 |
+
|
| 55 |
+
# Default data type, as PyTorch type.
|
| 56 |
+
DTYPE: torch.dtype = dtype_from_str(DTYPE_STR)
|
| 57 |
+
|
| 58 |
+
|
| 59 |
+
# Other defaults.
|
| 60 |
+
# ------------------------------------------------------------------------------
|
| 61 |
+
|
| 62 |
+
# Default device.
|
| 63 |
+
DEVICE: torch.device | str = "cuda"
|
| 64 |
+
|
| 65 |
+
# Default RNG seed for input generation.
|
| 66 |
+
RNG_SEED: int = 0
|
| 67 |
+
|
| 68 |
+
# Default number of group sizes.
|
| 69 |
+
NUM_GROUP_SIZES: int = 1
|
| 70 |
+
|
| 71 |
+
# Default transposition (NN).
|
| 72 |
+
TRANS_LHS: bool = False
|
| 73 |
+
TRANS_RHS: bool = False
|
| 74 |
+
|
| 75 |
+
|
| 76 |
+
# Parameter checking functions.
|
| 77 |
+
# ------------------------------------------------------------------------------
|
| 78 |
+
|
| 79 |
+
|
| 80 |
+
def is_power_of_2(x: int) -> bool:
|
| 81 |
+
return (x > 0) and (x & (x - 1) == 0)
|
| 82 |
+
|
| 83 |
+
|
| 84 |
+
def check_input_device_dtype(
|
| 85 |
+
lhs: Tensor, rhs: Tensor, group_sizes: Tensor, bias: Tensor | None = None
|
| 86 |
+
) -> None:
|
| 87 |
+
assert (
|
| 88 |
+
lhs.device == rhs.device == group_sizes.device
|
| 89 |
+
), f"All input tensors must be in the same device (lhs = {lhs.device}, rhs = {rhs.device}, group_sizes = {group_sizes.device})."
|
| 90 |
+
assert (
|
| 91 |
+
lhs.dtype == rhs.dtype
|
| 92 |
+
), f"lhs and rhs types must match (lhs = {lhs.dtype}, rhs = {rhs.dtype})."
|
| 93 |
+
assert group_sizes.dtype == torch.int32, "group_sizes type must be int32."
|
| 94 |
+
|
| 95 |
+
if bias is not None:
|
| 96 |
+
assert (
|
| 97 |
+
bias.device == lhs.device
|
| 98 |
+
), f"bias must be on the same device as lhs (bias = {bias.device}, lhs = {lhs.device})."
|
| 99 |
+
assert (
|
| 100 |
+
bias.dtype == lhs.dtype
|
| 101 |
+
), f"bias dtype must match lhs dtype (bias = {bias.dtype}, lhs = {lhs.dtype})."
|
| 102 |
+
|
| 103 |
+
|
| 104 |
+
def check_bias_shape_stride(bias: Tensor, G: int, N: int) -> None:
|
| 105 |
+
assert bias.shape == (
|
| 106 |
+
G,
|
| 107 |
+
N,
|
| 108 |
+
), f"bias must have shape (G, N) = ({G}, {N}), got {bias.shape}."
|
| 109 |
+
assert bias.stride() == (N, 1), "bias must be row-major (bias.stride() == (N, 1))."
|
| 110 |
+
|
| 111 |
+
|
| 112 |
+
# Generation of group sizes.
|
| 113 |
+
# ------------------------------------------------------------------------------
|
| 114 |
+
|
| 115 |
+
|
| 116 |
+
# Probabilities for generating random group sizes.
|
| 117 |
+
UNUSED_TOKENS_PROB: float = 0.0
|
| 118 |
+
UNUSED_EXPERTS_PROB: float = 0.1
|
| 119 |
+
|
| 120 |
+
|
| 121 |
+
def gen_uniform_group_sizes(
|
| 122 |
+
M: int,
|
| 123 |
+
G: int,
|
| 124 |
+
device: torch.device | str = DEVICE,
|
| 125 |
+
) -> Tensor:
|
| 126 |
+
assert M >= 0, f"Number of tokens M must be non-negative (it's {M})."
|
| 127 |
+
assert G > 0, f"Number of experts G must be positive (it's {G})."
|
| 128 |
+
|
| 129 |
+
base = M // G
|
| 130 |
+
remainder = M % G
|
| 131 |
+
group_sizes = torch.full((G,), base, dtype=torch.int32, device=device)
|
| 132 |
+
if remainder > 0:
|
| 133 |
+
group_sizes[:remainder] += 1
|
| 134 |
+
|
| 135 |
+
assert (
|
| 136 |
+
len(group_sizes) == G
|
| 137 |
+
), f"Group sizes don't have {G} elements (it's {len(group_sizes)})."
|
| 138 |
+
assert torch.all(group_sizes >= 0).item(), "All group sizes must be non-negative."
|
| 139 |
+
assert (
|
| 140 |
+
torch.sum(group_sizes).item() == M
|
| 141 |
+
), f"Group sizes don't add up to total tokens {M}."
|
| 142 |
+
assert group_sizes.dtype == torch.int32, "Group sizes must be int32."
|
| 143 |
+
|
| 144 |
+
return group_sizes
|
| 145 |
+
|
| 146 |
+
|
| 147 |
+
def gen_group_sizes(
|
| 148 |
+
M: int,
|
| 149 |
+
G: int,
|
| 150 |
+
device: torch.device | str = DEVICE,
|
| 151 |
+
rng_seed: int | None = RNG_SEED,
|
| 152 |
+
unused_tokens_prob: float = UNUSED_TOKENS_PROB,
|
| 153 |
+
unused_experts_prob: float = UNUSED_EXPERTS_PROB,
|
| 154 |
+
) -> Tensor:
|
| 155 |
+
assert M >= 0, f"Number of tokens M must be non-negative (it's {M})."
|
| 156 |
+
assert G > 0, f"Number of experts G must be positive (it's {G})."
|
| 157 |
+
assert (
|
| 158 |
+
0 <= unused_tokens_prob <= 1
|
| 159 |
+
), f"Probability of unused tokens must be in [0, 1] interval (it's {unused_tokens_prob})."
|
| 160 |
+
assert (
|
| 161 |
+
0 <= unused_experts_prob <= 1
|
| 162 |
+
), f"Probability of unused experts must be in [0, 1] interval (it's {unused_experts_prob})."
|
| 163 |
+
|
| 164 |
+
if rng_seed is not None:
|
| 165 |
+
torch.manual_seed(rng_seed)
|
| 166 |
+
|
| 167 |
+
if unused_tokens_prob > 0:
|
| 168 |
+
# Optionally drop tokens to simulate routing sparsity, some tokens may not be routed.
|
| 169 |
+
num_unused_tokens = M
|
| 170 |
+
while num_unused_tokens == M:
|
| 171 |
+
num_unused_tokens = int(
|
| 172 |
+
torch.binomial(
|
| 173 |
+
torch.tensor(float(M), device=device),
|
| 174 |
+
torch.tensor(unused_tokens_prob, device=device),
|
| 175 |
+
).item()
|
| 176 |
+
)
|
| 177 |
+
else:
|
| 178 |
+
num_unused_tokens = 0
|
| 179 |
+
num_used_tokens = M - num_unused_tokens
|
| 180 |
+
assert (
|
| 181 |
+
num_unused_tokens >= 0
|
| 182 |
+
), f"Number of unused tokens must be non-negative (it's {num_unused_tokens})."
|
| 183 |
+
assert (
|
| 184 |
+
num_used_tokens > 0
|
| 185 |
+
), f"Number of used tokens must be positive (it's {num_used_tokens})."
|
| 186 |
+
assert (
|
| 187 |
+
num_used_tokens + num_unused_tokens == M
|
| 188 |
+
), f"Unused + used tokens don't add up total tokens ({num_used_tokens} + {num_unused_tokens} != {M})."
|
| 189 |
+
|
| 190 |
+
if num_unused_tokens > 0:
|
| 191 |
+
_LOGGER.debug(
|
| 192 |
+
f"Group sizes generation: dropped {num_unused_tokens} token{'s' if num_unused_tokens > 1 else ''}.",
|
| 193 |
+
)
|
| 194 |
+
|
| 195 |
+
if unused_experts_prob > 0:
|
| 196 |
+
# Some experts may have zero tokens assigned to them.
|
| 197 |
+
num_used_experts = 0
|
| 198 |
+
while num_used_experts == 0:
|
| 199 |
+
used_experts = torch.nonzero(
|
| 200 |
+
torch.rand((G,), device=device) >= unused_experts_prob
|
| 201 |
+
).squeeze()
|
| 202 |
+
num_used_experts = used_experts.numel()
|
| 203 |
+
else:
|
| 204 |
+
used_experts = torch.arange(0, G, device=device)
|
| 205 |
+
num_used_experts = G
|
| 206 |
+
num_unused_experts = G - num_used_experts
|
| 207 |
+
assert (
|
| 208 |
+
num_unused_experts >= 0
|
| 209 |
+
), f"Number of unused experts must be non-negative (it's {num_unused_experts})."
|
| 210 |
+
assert (
|
| 211 |
+
num_used_experts >= 1
|
| 212 |
+
), f"At least one expert must be used (it's {num_used_experts})."
|
| 213 |
+
assert (
|
| 214 |
+
num_unused_experts + num_used_experts == G
|
| 215 |
+
), f"Unused + used experts don't add up total experts ({num_unused_experts} + {num_used_experts} != {G})."
|
| 216 |
+
|
| 217 |
+
if num_unused_experts > 0:
|
| 218 |
+
_LOGGER.debug(
|
| 219 |
+
f"Group sizes generation: dropped {num_unused_experts} expert{'s' if num_unused_experts > 1 else ''}.",
|
| 220 |
+
)
|
| 221 |
+
|
| 222 |
+
group_sizes = torch.bincount(
|
| 223 |
+
used_experts[
|
| 224 |
+
torch.randint(low=0, high=num_used_experts, size=(num_used_tokens,))
|
| 225 |
+
],
|
| 226 |
+
minlength=G,
|
| 227 |
+
).to(torch.int32)
|
| 228 |
+
|
| 229 |
+
assert (
|
| 230 |
+
len(group_sizes) == G
|
| 231 |
+
), f"Group sizes don't have {G} elements (it's {len(group_sizes)})."
|
| 232 |
+
assert torch.all(group_sizes >= 0).item(), "All group sizes must be non-negative."
|
| 233 |
+
assert (
|
| 234 |
+
torch.sum(group_sizes).item() == num_used_tokens
|
| 235 |
+
), f"Group sizes don't add up to used tokens {num_used_tokens}."
|
| 236 |
+
assert group_sizes.dtype == torch.int32, "Group sizes must be int32."
|
| 237 |
+
|
| 238 |
+
return group_sizes
|
| 239 |
+
|
| 240 |
+
|
| 241 |
+
def gen_multiple_group_sizes(
|
| 242 |
+
num_group_sizes: int,
|
| 243 |
+
M: int,
|
| 244 |
+
G: int,
|
| 245 |
+
device: torch.device | str = DEVICE,
|
| 246 |
+
rng_seed: int | None = RNG_SEED,
|
| 247 |
+
unused_tokens_prob: float = UNUSED_TOKENS_PROB,
|
| 248 |
+
unused_experts_prob: float = UNUSED_EXPERTS_PROB,
|
| 249 |
+
group_sizes_0: Tensor | None = None,
|
| 250 |
+
) -> list[Tensor]:
|
| 251 |
+
assert (
|
| 252 |
+
num_group_sizes > 0
|
| 253 |
+
), f"Number of group sizes to be generated must be positive, it's {num_group_sizes}."
|
| 254 |
+
multiple_group_sizes = [
|
| 255 |
+
gen_group_sizes(
|
| 256 |
+
M,
|
| 257 |
+
G,
|
| 258 |
+
device=device,
|
| 259 |
+
rng_seed=rng_seed if g == 0 else None,
|
| 260 |
+
unused_tokens_prob=unused_tokens_prob,
|
| 261 |
+
unused_experts_prob=unused_experts_prob,
|
| 262 |
+
)
|
| 263 |
+
for g in range(
|
| 264 |
+
num_group_sizes if group_sizes_0 is None else num_group_sizes - 1
|
| 265 |
+
)
|
| 266 |
+
]
|
| 267 |
+
if group_sizes_0 is not None:
|
| 268 |
+
multiple_group_sizes.insert(0, group_sizes_0)
|
| 269 |
+
assert (
|
| 270 |
+
len(multiple_group_sizes) == num_group_sizes
|
| 271 |
+
), f"Expecting {num_group_sizes} distinct group sizes (it's {len(multiple_group_sizes)})."
|
| 272 |
+
return multiple_group_sizes
|
| 273 |
+
|
| 274 |
+
|
| 275 |
+
# GMM helpers: tensor generation.
|
| 276 |
+
# ------------------------------------------------------------------------------
|
| 277 |
+
|
| 278 |
+
|
| 279 |
+
def gen_gmm_input(
|
| 280 |
+
M: int,
|
| 281 |
+
K: int,
|
| 282 |
+
N: int,
|
| 283 |
+
G: int,
|
| 284 |
+
device: torch.device | str = DEVICE,
|
| 285 |
+
preferred_element_type: torch.dtype = DTYPE,
|
| 286 |
+
trans_rhs: bool = TRANS_RHS,
|
| 287 |
+
rng_seed: int | None = RNG_SEED,
|
| 288 |
+
unif_group_sizes: bool = False,
|
| 289 |
+
) -> tuple[Tensor, Tensor, Tensor]:
|
| 290 |
+
assert M > 0, f"Number of lhs rows M must be positive (M = {M})."
|
| 291 |
+
assert K > 0, f"Number of lhs columns / rhs rows K must be positive (K = {K})."
|
| 292 |
+
assert N > 0, f"Number of rhs columns N must be positive (N = {N})."
|
| 293 |
+
assert G > 0, f"Number of groups G must be positive (G = {G})."
|
| 294 |
+
|
| 295 |
+
if rng_seed is not None:
|
| 296 |
+
torch.manual_seed(rng_seed)
|
| 297 |
+
|
| 298 |
+
lhs = torch.randn((M, K), dtype=torch.float32, device=device)
|
| 299 |
+
lhs = lhs.to(preferred_element_type)
|
| 300 |
+
|
| 301 |
+
if trans_rhs:
|
| 302 |
+
rhs = torch.randn((G, N, K), dtype=torch.float32, device=device).permute(
|
| 303 |
+
0, 2, 1
|
| 304 |
+
)
|
| 305 |
+
else:
|
| 306 |
+
rhs = torch.randn((G, K, N), dtype=torch.float32, device=device)
|
| 307 |
+
rhs = rhs.to(preferred_element_type)
|
| 308 |
+
|
| 309 |
+
group_sizes = (
|
| 310 |
+
gen_uniform_group_sizes(M, G, device=device)
|
| 311 |
+
if unif_group_sizes
|
| 312 |
+
else gen_group_sizes(M, G, device=device, rng_seed=None)
|
| 313 |
+
)
|
| 314 |
+
|
| 315 |
+
return lhs, rhs, group_sizes
|
| 316 |
+
|
| 317 |
+
|
| 318 |
+
def gen_gmm_output(
|
| 319 |
+
M: int,
|
| 320 |
+
N: int,
|
| 321 |
+
device: torch.device | str = DEVICE,
|
| 322 |
+
preferred_element_type: torch.dtype = DTYPE,
|
| 323 |
+
) -> Tensor:
|
| 324 |
+
assert M > 0, f"Number of out rows M must be positive (M = {M})."
|
| 325 |
+
assert N > 0, f"Number of out columns N must be positive (N = {N})."
|
| 326 |
+
|
| 327 |
+
out = torch.empty((M, N), dtype=preferred_element_type, device=device)
|
| 328 |
+
|
| 329 |
+
return out
|
| 330 |
+
|
| 331 |
+
|
| 332 |
+
def gen_gmm_tensors(
|
| 333 |
+
M: int,
|
| 334 |
+
K: int,
|
| 335 |
+
N: int,
|
| 336 |
+
G: int,
|
| 337 |
+
num_group_sizes: int,
|
| 338 |
+
device: torch.device | str = DEVICE,
|
| 339 |
+
input_type: torch.dtype = DTYPE,
|
| 340 |
+
output_type: torch.dtype = DTYPE,
|
| 341 |
+
trans_lhs: bool = False,
|
| 342 |
+
trans_rhs: bool = TRANS_RHS,
|
| 343 |
+
rng_seed: int | None = RNG_SEED,
|
| 344 |
+
unif_group_sizes: bool = False,
|
| 345 |
+
use_bias: bool = False,
|
| 346 |
+
) -> tuple[Tensor, Tensor, list[Tensor], Tensor, Tensor | None]:
|
| 347 |
+
lhs, rhs, group_sizes_0 = gen_gmm_input(
|
| 348 |
+
M,
|
| 349 |
+
K,
|
| 350 |
+
N,
|
| 351 |
+
G,
|
| 352 |
+
device=device,
|
| 353 |
+
preferred_element_type=input_type,
|
| 354 |
+
trans_rhs=trans_rhs,
|
| 355 |
+
rng_seed=rng_seed,
|
| 356 |
+
unif_group_sizes=unif_group_sizes,
|
| 357 |
+
)
|
| 358 |
+
multiple_group_sizes = gen_multiple_group_sizes(
|
| 359 |
+
num_group_sizes, M, G, device=device, rng_seed=None, group_sizes_0=group_sizes_0
|
| 360 |
+
)
|
| 361 |
+
out = gen_gmm_output(M, N, device=device, preferred_element_type=output_type)
|
| 362 |
+
bias = None
|
| 363 |
+
if use_bias:
|
| 364 |
+
torch.manual_seed(rng_seed + 1000) # Different seed for bias
|
| 365 |
+
bias = torch.randn(G, N, dtype=input_type, device=device)
|
| 366 |
+
|
| 367 |
+
return lhs, rhs, multiple_group_sizes, out, bias
|
| 368 |
+
|
| 369 |
+
|
| 370 |
+
# GMM helpers: get information from tensors.
|
| 371 |
+
# ------------------------------------------------------------------------------
|
| 372 |
+
|
| 373 |
+
|
| 374 |
+
def get_gmm_shape(
|
| 375 |
+
lhs: Tensor, rhs: Tensor, group_sizes: Tensor
|
| 376 |
+
) -> tuple[int, int, int, int]:
|
| 377 |
+
assert lhs.dim() == 2, f"lhs must have 2 dimensions (it's {lhs.dim()})."
|
| 378 |
+
assert rhs.dim() == 3, f"rhs must have 3 dimensions (it's {rhs.dim()})."
|
| 379 |
+
assert (
|
| 380 |
+
group_sizes.dim() == 1
|
| 381 |
+
), f"group_sizes must have 1 dimension (it's {group_sizes.dim()})."
|
| 382 |
+
|
| 383 |
+
M, lhs_k = lhs.shape
|
| 384 |
+
rhs_g, rhs_k, N = rhs.shape
|
| 385 |
+
group_sizes_g = group_sizes.shape[0]
|
| 386 |
+
|
| 387 |
+
assert (
|
| 388 |
+
lhs_k == rhs_k
|
| 389 |
+
), f"K dimension of lhs and rhs don't match (lhs = {lhs_k}, rhs = {rhs_k})."
|
| 390 |
+
K = lhs_k
|
| 391 |
+
assert (
|
| 392 |
+
rhs_g == group_sizes_g
|
| 393 |
+
), f"G dimension of rhs and group_sizes don't match (rhs = {rhs_g}, group_sizes = {group_sizes_g})."
|
| 394 |
+
G = rhs_g
|
| 395 |
+
|
| 396 |
+
assert M > 0, f"M must be positive, it's {M}."
|
| 397 |
+
assert K > 0, f"K must be positive, it's {K}."
|
| 398 |
+
assert N > 0, f"N must be positive, it's {N}"
|
| 399 |
+
assert G > 0, f"G must be positive, it's {G}"
|
| 400 |
+
|
| 401 |
+
return M, K, N, G
|
| 402 |
+
|
| 403 |
+
|
| 404 |
+
def get_gmm_output(
|
| 405 |
+
M: int,
|
| 406 |
+
N: int,
|
| 407 |
+
device: torch.device | str = DEVICE,
|
| 408 |
+
preferred_element_type: torch.dtype = DTYPE,
|
| 409 |
+
existing_out: Tensor | None = None,
|
| 410 |
+
) -> Tensor:
|
| 411 |
+
assert M > 0, f"Number of out rows M must be positive (M = {M})."
|
| 412 |
+
assert N > 0, f"Number of out columns N must be positive (N = {N})."
|
| 413 |
+
|
| 414 |
+
if existing_out is not None:
|
| 415 |
+
assert (
|
| 416 |
+
existing_out.device == device
|
| 417 |
+
), f"Existing output device and provided device don't match (existing = {existing_out.device}, provided = {device})."
|
| 418 |
+
assert (
|
| 419 |
+
existing_out.dtype == preferred_element_type
|
| 420 |
+
), f"Existing output type and preferred output type don't match (existing = {existing_out.dtype}, preferred = {preferred_element_type})."
|
| 421 |
+
assert existing_out.shape == (
|
| 422 |
+
M,
|
| 423 |
+
N,
|
| 424 |
+
), f"Existing output shape and GMM shape don't match (existing = {tuple(existing_out.shape)}, provided = {(M, N)})."
|
| 425 |
+
return existing_out
|
| 426 |
+
|
| 427 |
+
return gen_gmm_output(
|
| 428 |
+
M,
|
| 429 |
+
N,
|
| 430 |
+
device=device,
|
| 431 |
+
preferred_element_type=preferred_element_type,
|
| 432 |
+
)
|
| 433 |
+
|
| 434 |
+
|
| 435 |
+
def get_gmm_transposition(lhs: Tensor, rhs: Tensor, out: Tensor) -> tuple[bool, int]:
|
| 436 |
+
assert lhs.dim() == 2, f"lhs must have 2 dimensions (it's {lhs.dim()})."
|
| 437 |
+
assert rhs.dim() == 3, f"rhs must have 3 dimensions (it's {rhs.dim()})."
|
| 438 |
+
assert out.dim() == 2, f"out must have 2 dimensions (it's {out.dim()})."
|
| 439 |
+
|
| 440 |
+
lhs_m, lhs_k = lhs.shape
|
| 441 |
+
G, rhs_k, rhs_n = rhs.shape
|
| 442 |
+
out_m, out_n = out.shape
|
| 443 |
+
|
| 444 |
+
assert (
|
| 445 |
+
lhs_m == out_m
|
| 446 |
+
), f"M dimension of lhs and out don't match (lhs = {lhs_m}, rhs = {out_m})."
|
| 447 |
+
M = lhs_m
|
| 448 |
+
assert (
|
| 449 |
+
lhs_k == rhs_k
|
| 450 |
+
), f"K dimension of lhs and rhs don't match (lhs = {lhs_k}, rhs = {rhs_k})."
|
| 451 |
+
K = lhs_k
|
| 452 |
+
assert (
|
| 453 |
+
rhs_n == out_n
|
| 454 |
+
), f"N dimension of rhs and out don't match (lhs = {rhs_n}, rhs = {out_n})."
|
| 455 |
+
N = rhs_n
|
| 456 |
+
|
| 457 |
+
assert M > 0, f"M must be positive, it's {M}."
|
| 458 |
+
assert K > 0, f"K must be positive, it's {K}."
|
| 459 |
+
assert N > 0, f"N must be positive, it's {N}"
|
| 460 |
+
assert G > 0, f"G must be positive, it's {G}"
|
| 461 |
+
|
| 462 |
+
is_lhs_row_major = lhs.stride() == (K, 1)
|
| 463 |
+
assert is_lhs_row_major, "lhs must be row-major."
|
| 464 |
+
is_rhs_row_major = rhs.stride() == (K * N, N, 1)
|
| 465 |
+
is_rhs_col_major = rhs.stride() == (K * N, 1, K)
|
| 466 |
+
assert (
|
| 467 |
+
is_rhs_row_major != is_rhs_col_major
|
| 468 |
+
), "rhs must be row-major or column-major."
|
| 469 |
+
is_out_row_major = out.stride() == (N, 1)
|
| 470 |
+
assert is_out_row_major, "out must be row-major."
|
| 471 |
+
|
| 472 |
+
# Get rhs leading dimension according to transposition configuration.
|
| 473 |
+
ld_rhs = N if is_rhs_row_major else K
|
| 474 |
+
|
| 475 |
+
return is_rhs_col_major, ld_rhs
|
| 476 |
+
|
| 477 |
+
|
| 478 |
+
# TGMM helpers: tensor generation.
|
| 479 |
+
# ------------------------------------------------------------------------------
|
| 480 |
+
|
| 481 |
+
|
| 482 |
+
def gen_tgmm_input(
|
| 483 |
+
M: int,
|
| 484 |
+
K: int,
|
| 485 |
+
N: int,
|
| 486 |
+
G: int,
|
| 487 |
+
device: torch.device | str = DEVICE,
|
| 488 |
+
preferred_element_type: torch.dtype = DTYPE,
|
| 489 |
+
trans_lhs: bool = TRANS_LHS,
|
| 490 |
+
rng_seed: int | None = RNG_SEED,
|
| 491 |
+
unif_group_sizes: bool = False,
|
| 492 |
+
) -> tuple[Tensor, Tensor, Tensor]:
|
| 493 |
+
assert K > 0, f"Number of lhs rows K must be positive (M = {K})."
|
| 494 |
+
assert M > 0, f"Number of lhs columns / rhs rows M must be positive (K = {M})."
|
| 495 |
+
assert N > 0, f"Number of rhs columns N must be positive (N = {N})."
|
| 496 |
+
assert G > 0, f"Number of groups G must be positive (G = {G})."
|
| 497 |
+
|
| 498 |
+
if rng_seed is not None:
|
| 499 |
+
torch.manual_seed(rng_seed)
|
| 500 |
+
|
| 501 |
+
if trans_lhs:
|
| 502 |
+
lhs = torch.randn((M, K), dtype=torch.float32, device=device).T
|
| 503 |
+
else:
|
| 504 |
+
lhs = torch.randn((K, M), dtype=torch.float32, device=device)
|
| 505 |
+
lhs = lhs.to(preferred_element_type)
|
| 506 |
+
|
| 507 |
+
rhs = torch.randn((M, N), dtype=torch.float32, device=device)
|
| 508 |
+
rhs = rhs.to(preferred_element_type)
|
| 509 |
+
|
| 510 |
+
group_sizes = (
|
| 511 |
+
gen_uniform_group_sizes(M, G, device=device)
|
| 512 |
+
if unif_group_sizes
|
| 513 |
+
else gen_group_sizes(M, G, device=device, rng_seed=None)
|
| 514 |
+
)
|
| 515 |
+
|
| 516 |
+
return lhs, rhs, group_sizes
|
| 517 |
+
|
| 518 |
+
|
| 519 |
+
def gen_tgmm_output(
|
| 520 |
+
K: int,
|
| 521 |
+
N: int,
|
| 522 |
+
G: int,
|
| 523 |
+
device: torch.device | str = DEVICE,
|
| 524 |
+
preferred_element_type: torch.dtype = DTYPE,
|
| 525 |
+
) -> Tensor:
|
| 526 |
+
assert K > 0, f"Number of out rows K must be positive (K = {K})."
|
| 527 |
+
assert N > 0, f"Number of out columns N must be positive (N = {N})."
|
| 528 |
+
assert G > 0, f"Number of groups G must be positive (G = {G})."
|
| 529 |
+
|
| 530 |
+
out = torch.empty((G, K, N), dtype=preferred_element_type, device=device)
|
| 531 |
+
|
| 532 |
+
return out
|
| 533 |
+
|
| 534 |
+
|
| 535 |
+
def gen_tgmm_bias_grad(
|
| 536 |
+
K: int,
|
| 537 |
+
G: int,
|
| 538 |
+
device: torch.device | str = DEVICE,
|
| 539 |
+
with_bias_grad: bool = False,
|
| 540 |
+
) -> Tensor:
|
| 541 |
+
if with_bias_grad:
|
| 542 |
+
assert K > 0, f"Number of bias_grad rows K must be positive (K = {K})."
|
| 543 |
+
assert G > 0, f"Number of groups G must be positive (G = {G})."
|
| 544 |
+
return torch.empty((G, K), device=device, dtype=torch.float32)
|
| 545 |
+
else:
|
| 546 |
+
# Return dummy pointer when bias_grad is not needed.
|
| 547 |
+
# Must be float32 because atomic_add does not support bf16/fp16,
|
| 548 |
+
# and Triton validates the pointer dtype even in dead branches.
|
| 549 |
+
return torch.tensor([], device=device, dtype=torch.float32)
|
| 550 |
+
|
| 551 |
+
|
| 552 |
+
def gen_tgmm_tensors(
|
| 553 |
+
M: int,
|
| 554 |
+
K: int,
|
| 555 |
+
N: int,
|
| 556 |
+
G: int,
|
| 557 |
+
num_group_sizes: int,
|
| 558 |
+
device: torch.device | str = DEVICE,
|
| 559 |
+
input_type: torch.dtype = DTYPE,
|
| 560 |
+
output_type: torch.dtype = DTYPE,
|
| 561 |
+
trans_lhs: bool = TRANS_LHS,
|
| 562 |
+
trans_rhs: bool = False,
|
| 563 |
+
rng_seed: int | None = RNG_SEED,
|
| 564 |
+
unif_group_sizes: bool = False,
|
| 565 |
+
use_bias: bool = False,
|
| 566 |
+
) -> tuple[Tensor, Tensor, list[Tensor], Tensor, Tensor | None]:
|
| 567 |
+
lhs, rhs, group_sizes_0 = gen_tgmm_input(
|
| 568 |
+
M,
|
| 569 |
+
K,
|
| 570 |
+
N,
|
| 571 |
+
G,
|
| 572 |
+
device=device,
|
| 573 |
+
preferred_element_type=input_type,
|
| 574 |
+
trans_lhs=trans_lhs,
|
| 575 |
+
rng_seed=rng_seed,
|
| 576 |
+
unif_group_sizes=unif_group_sizes,
|
| 577 |
+
)
|
| 578 |
+
multiple_group_sizes = gen_multiple_group_sizes(
|
| 579 |
+
num_group_sizes, M, G, device=device, rng_seed=None, group_sizes_0=group_sizes_0
|
| 580 |
+
)
|
| 581 |
+
out = gen_tgmm_output(K, N, G, device=device, preferred_element_type=output_type)
|
| 582 |
+
if use_bias:
|
| 583 |
+
bias_grad = gen_tgmm_bias_grad(K, G, device=device, with_bias_grad=True)
|
| 584 |
+
else:
|
| 585 |
+
bias_grad = None
|
| 586 |
+
return lhs, rhs, multiple_group_sizes, out, bias_grad
|
| 587 |
+
|
| 588 |
+
|
| 589 |
+
# TGMM helpers: get information from tensors.
|
| 590 |
+
# ------------------------------------------------------------------------------
|
| 591 |
+
|
| 592 |
+
|
| 593 |
+
def get_tgmm_shape(
|
| 594 |
+
lhs: Tensor, rhs: Tensor, group_sizes: Tensor
|
| 595 |
+
) -> tuple[int, int, int, int]:
|
| 596 |
+
assert lhs.dim() == 2, f"lhs must have 2 dimensions (it's {lhs.dim()})."
|
| 597 |
+
assert rhs.dim() == 2, f"rhs must have 2 dimensions (it's {rhs.dim()})."
|
| 598 |
+
assert (
|
| 599 |
+
group_sizes.dim() == 1
|
| 600 |
+
), f"group_sizes must have 1 dimension (it's {group_sizes.dim()})."
|
| 601 |
+
|
| 602 |
+
K, lhs_m = lhs.shape
|
| 603 |
+
rhs_m, N = rhs.shape
|
| 604 |
+
G = group_sizes.shape[0]
|
| 605 |
+
|
| 606 |
+
assert (
|
| 607 |
+
lhs_m == rhs_m
|
| 608 |
+
), f"M dimension of lhs and rhs don't match (lhs = {lhs_m}, rhs = {rhs_m})."
|
| 609 |
+
M = lhs_m
|
| 610 |
+
|
| 611 |
+
assert M > 0, f"M must be positive, it's {M}."
|
| 612 |
+
assert K > 0, f"K must be positive, it's {K}."
|
| 613 |
+
assert N > 0, f"N must be positive, it's {N}"
|
| 614 |
+
assert G > 0, f"G must be positive, it's {G}"
|
| 615 |
+
|
| 616 |
+
return M, K, N, G
|
| 617 |
+
|
| 618 |
+
|
| 619 |
+
def get_tgmm_output(
|
| 620 |
+
K: int,
|
| 621 |
+
N: int,
|
| 622 |
+
G: int,
|
| 623 |
+
device: torch.device | str = DEVICE,
|
| 624 |
+
preferred_element_type: torch.dtype = DTYPE,
|
| 625 |
+
existing_out: Tensor | None = None,
|
| 626 |
+
) -> Tensor:
|
| 627 |
+
assert K > 0, f"Number of out rows K must be positive (K = {K})."
|
| 628 |
+
assert N > 0, f"Number of out columns N must be positive (N = {N})."
|
| 629 |
+
assert G > 0, f"Number of groups G must be positive (G = {G})."
|
| 630 |
+
|
| 631 |
+
if existing_out is not None:
|
| 632 |
+
assert (
|
| 633 |
+
existing_out.device == device
|
| 634 |
+
), f"Existing output device and provided device don't match (existing = {existing_out.device}, provided = {device})."
|
| 635 |
+
assert (
|
| 636 |
+
existing_out.dtype == preferred_element_type
|
| 637 |
+
), f"Existing output type and preferred output type don't match (existing = {existing_out.dtype}, preferred = {preferred_element_type})."
|
| 638 |
+
assert existing_out.shape == (
|
| 639 |
+
G,
|
| 640 |
+
K,
|
| 641 |
+
N,
|
| 642 |
+
), f"Existing output shape and GMM shape don't match (existing = {tuple(existing_out.shape)}, provided = {(G, K, N)})."
|
| 643 |
+
return existing_out
|
| 644 |
+
|
| 645 |
+
return gen_tgmm_output(
|
| 646 |
+
K,
|
| 647 |
+
N,
|
| 648 |
+
G,
|
| 649 |
+
device=device,
|
| 650 |
+
preferred_element_type=preferred_element_type,
|
| 651 |
+
)
|
| 652 |
+
|
| 653 |
+
|
| 654 |
+
def get_tgmm_bias_grad(
|
| 655 |
+
K: int,
|
| 656 |
+
G: int,
|
| 657 |
+
device: torch.device | str = DEVICE,
|
| 658 |
+
existing_bias_grad: Tensor | None = None,
|
| 659 |
+
) -> Tensor:
|
| 660 |
+
"""
|
| 661 |
+
Get or validate bias gradient tensor for TGMM.
|
| 662 |
+
|
| 663 |
+
If existing_bias_grad is provided, validates its shape, device, dtype, and stride,
|
| 664 |
+
and always zeros it before returning (since the kernel uses atomic_add).
|
| 665 |
+
If existing_bias_grad is None, returns a dummy tensor (for use when COMPUTE_BIAS_GRAD=False).
|
| 666 |
+
Parameters
|
| 667 |
+
----------
|
| 668 |
+
K : int
|
| 669 |
+
Number of rows in the bias gradient tensor.
|
| 670 |
+
G : int
|
| 671 |
+
Number of groups.
|
| 672 |
+
device : torch.device or str
|
| 673 |
+
Device for the tensor.
|
| 674 |
+
existing_bias_grad : torch.Tensor or None
|
| 675 |
+
Existing bias gradient tensor to validate and use.
|
| 676 |
+
Returns
|
| 677 |
+
-------
|
| 678 |
+
torch.Tensor
|
| 679 |
+
Valid bias gradient tensor or dummy tensor.
|
| 680 |
+
"""
|
| 681 |
+
assert K > 0, f"Number of bias_grad rows K must be positive (K = {K})."
|
| 682 |
+
assert G > 0, f"Number of groups G must be positive (G = {G})."
|
| 683 |
+
|
| 684 |
+
if existing_bias_grad is not None:
|
| 685 |
+
# Validate existing bias_grad tensor.
|
| 686 |
+
expected_shape = (G, K)
|
| 687 |
+
assert (
|
| 688 |
+
tuple(existing_bias_grad.shape) == expected_shape
|
| 689 |
+
), f"bias_grad must have shape {expected_shape}, got {tuple(existing_bias_grad.shape)}."
|
| 690 |
+
assert (
|
| 691 |
+
existing_bias_grad.device == device
|
| 692 |
+
), f"bias_grad must be on the same device (bias_grad = {existing_bias_grad.device}, device = {device})."
|
| 693 |
+
assert (
|
| 694 |
+
existing_bias_grad.dtype == torch.float32
|
| 695 |
+
), f"bias_grad must be torch.float32 (kernel uses atomic_add which requires float32), got {existing_bias_grad.dtype}."
|
| 696 |
+
assert existing_bias_grad.stride() == (
|
| 697 |
+
K,
|
| 698 |
+
1,
|
| 699 |
+
), f"bias_grad must be row-major with stride (K, 1) = ({K}, 1), got {existing_bias_grad.stride()}."
|
| 700 |
+
|
| 701 |
+
# Always zero the tensor since bias_grad represents gradients for the current
|
| 702 |
+
# computation and should start fresh. The kernel uses atomic_add which adds to
|
| 703 |
+
# existing values, so we must zero before the kernel runs.
|
| 704 |
+
existing_bias_grad.zero_()
|
| 705 |
+
|
| 706 |
+
return existing_bias_grad
|
| 707 |
+
|
| 708 |
+
else:
|
| 709 |
+
return gen_tgmm_bias_grad(K, G, device=device, with_bias_grad=False)
|
| 710 |
+
|
| 711 |
+
|
| 712 |
+
def get_tgmm_transposition(lhs: Tensor, rhs: Tensor, out: Tensor) -> tuple[bool, int]:
|
| 713 |
+
assert lhs.dim() == 2, f"lhs must have 2 dimensions (it's {lhs.dim()})."
|
| 714 |
+
assert rhs.dim() == 2, f"rhs must have 2 dimensions (it's {rhs.dim()})."
|
| 715 |
+
assert out.dim() == 3, f"out must have 3 dimensions (it's {out.dim()})."
|
| 716 |
+
|
| 717 |
+
lhs_k, lhs_m = lhs.shape
|
| 718 |
+
rhs_m, rhs_n = rhs.shape
|
| 719 |
+
G, out_k, out_n = out.shape
|
| 720 |
+
|
| 721 |
+
assert (
|
| 722 |
+
lhs_m == rhs_m
|
| 723 |
+
), f"M dimension of lhs and rhs don't match (lhs = {lhs_m}, rhs = {rhs_m})."
|
| 724 |
+
M = lhs_m
|
| 725 |
+
assert (
|
| 726 |
+
lhs_k == out_k
|
| 727 |
+
), f"K dimension of lhs and out don't match (lhs = {lhs_k}, rhs = {out_k})."
|
| 728 |
+
K = lhs_k
|
| 729 |
+
assert (
|
| 730 |
+
rhs_n == out_n
|
| 731 |
+
), f"N dimension of rhs and out don't match (lhs = {rhs_n}, rhs = {out_n})."
|
| 732 |
+
N = rhs_n
|
| 733 |
+
|
| 734 |
+
assert M > 0, f"M must be positive, it's {M}."
|
| 735 |
+
assert K > 0, f"K must be positive, it's {K}."
|
| 736 |
+
assert N > 0, f"N must be positive, it's {N}"
|
| 737 |
+
assert G > 0, f"G must be positive, it's {G}"
|
| 738 |
+
|
| 739 |
+
is_lhs_row_major = lhs.stride() == (M, 1)
|
| 740 |
+
is_lhs_col_major = lhs.stride() == (1, K)
|
| 741 |
+
assert (
|
| 742 |
+
is_lhs_row_major != is_lhs_col_major
|
| 743 |
+
), "lhs must be row-major or column-major."
|
| 744 |
+
is_rhs_row_major = rhs.stride() == (N, 1)
|
| 745 |
+
assert is_rhs_row_major, "rhs must be row-major."
|
| 746 |
+
is_out_row_major = out.stride() == (K * N, N, 1)
|
| 747 |
+
assert is_out_row_major, "out must be row-major."
|
| 748 |
+
|
| 749 |
+
# Get lhs leading dimension according to transposition configuration.
|
| 750 |
+
ld_lhs = M if is_lhs_row_major else K
|
| 751 |
+
|
| 752 |
+
return is_lhs_col_major, ld_lhs
|
build/torch214-cxx11-cu126-x86_64-linux/_grouped_gemm_triton/utils/logger.py
ADDED
|
@@ -0,0 +1,47 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
import os
|
| 2 |
+
import logging
|
| 3 |
+
|
| 4 |
+
|
| 5 |
+
# AITER Triton Logger which is singleton object around python logging.
|
| 6 |
+
# Note: Python logging is also a singleton object, but we want to read the
|
| 7 |
+
# env var AITER_LOG_LEVEL once at the beginning. Another alternative is to do
|
| 8 |
+
# this in __init__.py. In fact, that's how CK logger is setup. We can look at
|
| 9 |
+
# switching to that at some point
|
| 10 |
+
#
|
| 11 |
+
# AITER_LOG_LEVEL follows python logging levels
|
| 12 |
+
# DEBUG
|
| 13 |
+
# INFO
|
| 14 |
+
# WARNING
|
| 15 |
+
# ERROR
|
| 16 |
+
# CRITICAL
|
| 17 |
+
#
|
| 18 |
+
class AiterTritonLogger(object):
|
| 19 |
+
_instance = None
|
| 20 |
+
|
| 21 |
+
def __new__(cls):
|
| 22 |
+
if cls._instance is None:
|
| 23 |
+
cls._instance = super(AiterTritonLogger, cls).__new__(cls)
|
| 24 |
+
log_level_str = os.getenv("AITER_TRITON_LOG_LEVEL", "WARNING").upper()
|
| 25 |
+
numeric_level = getattr(logging, log_level_str, logging.WARNING)
|
| 26 |
+
cls._instance._logger = logging.getLogger("AITER_TRITON")
|
| 27 |
+
cls._instance._logger.setLevel(numeric_level)
|
| 28 |
+
|
| 29 |
+
return cls._instance
|
| 30 |
+
|
| 31 |
+
def get_logger(self):
|
| 32 |
+
return self._logger
|
| 33 |
+
|
| 34 |
+
def debug(self, msg):
|
| 35 |
+
self._logger.debug(msg)
|
| 36 |
+
|
| 37 |
+
def info(self, msg):
|
| 38 |
+
self._logger.info(msg)
|
| 39 |
+
|
| 40 |
+
def warning(self, msg):
|
| 41 |
+
self._logger.warning(msg)
|
| 42 |
+
|
| 43 |
+
def error(self, msg):
|
| 44 |
+
self._logger.error(msg)
|
| 45 |
+
|
| 46 |
+
def critical(self, msg):
|
| 47 |
+
self._logger.critical(msg)
|
build/torch214-cxx11-cu126-x86_64-linux/_layers/__init__.py
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2024 Databricks
|
| 2 |
+
# SPDX-License-Identifier: Apache-2.0
|
| 3 |
+
|
| 4 |
+
# from megablocks.layers.dmoe import dMoE
|
| 5 |
+
from .moe import MoE
|
| 6 |
+
|
| 7 |
+
__all__ = [
|
| 8 |
+
'MoE',
|
| 9 |
+
# 'dMoE',
|
| 10 |
+
]
|
build/torch214-cxx11-cu126-x86_64-linux/_layers/activation_fn.py
ADDED
|
@@ -0,0 +1,33 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2024 Databricks
|
| 2 |
+
# SPDX-License-Identifier: Apache-2.0
|
| 3 |
+
|
| 4 |
+
from typing import Any, Callable, Union
|
| 5 |
+
|
| 6 |
+
import torch
|
| 7 |
+
from ..stk import Matrix
|
| 8 |
+
|
| 9 |
+
|
| 10 |
+
def act_fn(
|
| 11 |
+
x: Matrix,
|
| 12 |
+
function: Callable,
|
| 13 |
+
return_grad_fn: bool = False,
|
| 14 |
+
**kwargs,
|
| 15 |
+
) -> Union[tuple[Matrix, Any] | Matrix]:
|
| 16 |
+
assert isinstance(x, Matrix)
|
| 17 |
+
with torch.set_grad_enabled(torch.is_grad_enabled() or return_grad_fn):
|
| 18 |
+
if return_grad_fn:
|
| 19 |
+
x.data.requires_grad = True
|
| 20 |
+
out = function(x.data, **kwargs)
|
| 21 |
+
y = Matrix(
|
| 22 |
+
x.size(),
|
| 23 |
+
out,
|
| 24 |
+
x.row_indices,
|
| 25 |
+
x.column_indices,
|
| 26 |
+
x.offsets,
|
| 27 |
+
x.column_indices_t,
|
| 28 |
+
x.offsets_t,
|
| 29 |
+
x.block_offsets_t,
|
| 30 |
+
)
|
| 31 |
+
if return_grad_fn:
|
| 32 |
+
return y, out.backward
|
| 33 |
+
return y
|
build/torch214-cxx11-cu126-x86_64-linux/_layers/all_to_all.py
ADDED
|
@@ -0,0 +1,54 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2024 Databricks
|
| 2 |
+
# SPDX-License-Identifier: Apache-2.0
|
| 3 |
+
|
| 4 |
+
import torch
|
| 5 |
+
import torch.distributed as dist
|
| 6 |
+
|
| 7 |
+
|
| 8 |
+
class AllToAllOp(torch.autograd.Function):
|
| 9 |
+
|
| 10 |
+
@staticmethod
|
| 11 |
+
def forward(ctx, x, output_split_sizes, input_split_sizes, group, async_op):
|
| 12 |
+
out = torch.empty((sum(output_split_sizes),) + x.shape[1:], device=x.device, dtype=x.dtype)
|
| 13 |
+
|
| 14 |
+
ctx.input_shape = x.shape
|
| 15 |
+
ctx.output_split_sizes = output_split_sizes
|
| 16 |
+
ctx.input_split_sizes = input_split_sizes
|
| 17 |
+
ctx.group = group
|
| 18 |
+
handle = dist.all_to_all_single(
|
| 19 |
+
out,
|
| 20 |
+
x,
|
| 21 |
+
output_split_sizes=output_split_sizes,
|
| 22 |
+
input_split_sizes=input_split_sizes,
|
| 23 |
+
group=group,
|
| 24 |
+
async_op=async_op,
|
| 25 |
+
)
|
| 26 |
+
return out, handle
|
| 27 |
+
|
| 28 |
+
@staticmethod
|
| 29 |
+
def backward(ctx, grad, _):
|
| 30 |
+
if ctx.needs_input_grad[0]:
|
| 31 |
+
out = torch.empty(
|
| 32 |
+
ctx.input_shape,
|
| 33 |
+
device=grad.device,
|
| 34 |
+
dtype=grad.dtype,
|
| 35 |
+
)
|
| 36 |
+
dist.all_to_all_single(
|
| 37 |
+
out,
|
| 38 |
+
grad,
|
| 39 |
+
output_split_sizes=ctx.input_split_sizes,
|
| 40 |
+
input_split_sizes=ctx.output_split_sizes,
|
| 41 |
+
group=ctx.group,
|
| 42 |
+
)
|
| 43 |
+
return out, None, None, None, None
|
| 44 |
+
return None, None, None, None, None
|
| 45 |
+
|
| 46 |
+
|
| 47 |
+
def all_to_all(x, output_split_sizes, input_split_sizes, group, async_op=False):
|
| 48 |
+
return AllToAllOp.apply(
|
| 49 |
+
x,
|
| 50 |
+
output_split_sizes,
|
| 51 |
+
input_split_sizes,
|
| 52 |
+
group,
|
| 53 |
+
async_op,
|
| 54 |
+
)
|
build/torch214-cxx11-cu126-x86_64-linux/_layers/arguments.py
ADDED
|
@@ -0,0 +1,101 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2024 Databricks
|
| 2 |
+
# SPDX-License-Identifier: Apache-2.0
|
| 3 |
+
|
| 4 |
+
import dataclasses
|
| 5 |
+
from functools import partial
|
| 6 |
+
from typing import Any, Callable, Optional, Union
|
| 7 |
+
|
| 8 |
+
import torch
|
| 9 |
+
import torch.distributed as dist
|
| 10 |
+
import torch.nn.functional as F
|
| 11 |
+
|
| 12 |
+
# import megablocks.grouped_gemm_util as grouped_gemm
|
| 13 |
+
from .. import grouped_gemm_util as grouped_gemm
|
| 14 |
+
|
| 15 |
+
# Type annotation for in-place Tensor initialization function.
|
| 16 |
+
InitFn = Union[Callable[[torch.Tensor], None], partial[torch.Tensor]]
|
| 17 |
+
|
| 18 |
+
_ALLOWED_BITWIDTHS = (-1, 4, 8)
|
| 19 |
+
|
| 20 |
+
DEFAULT_ACTIVATION_FN = partial(F.gelu, approximate='tanh')
|
| 21 |
+
|
| 22 |
+
|
| 23 |
+
@dataclasses.dataclass
|
| 24 |
+
class Arguments:
|
| 25 |
+
# Model arguments.
|
| 26 |
+
hidden_size: int = 1024
|
| 27 |
+
ffn_hidden_size: int = 4096
|
| 28 |
+
num_layers: int = 1
|
| 29 |
+
bias: bool = True
|
| 30 |
+
return_bias: bool = True
|
| 31 |
+
activation_fn: Optional[Callable] = DEFAULT_ACTIVATION_FN
|
| 32 |
+
|
| 33 |
+
# MoE arguments.
|
| 34 |
+
moe_num_experts: int = 1
|
| 35 |
+
moe_top_k: int = 1
|
| 36 |
+
moe_capacity_factor: int = 1
|
| 37 |
+
moe_normalize_expert_weights: Optional[Union[int, float]] = None
|
| 38 |
+
moe_loss_weight: float = 0.1
|
| 39 |
+
moe_jitter_eps: Optional[float] = None
|
| 40 |
+
moe_lbl_in_fp32: bool = False
|
| 41 |
+
|
| 42 |
+
# Parallelism arguments.
|
| 43 |
+
moe_expert_model_parallelism: bool = False
|
| 44 |
+
expert_parallel_group: Optional[dist.ProcessGroup] = None
|
| 45 |
+
pipeline_model_parallel_size: int = 1
|
| 46 |
+
num_layers_per_virtual_pipeline_stage: Optional[int] = None
|
| 47 |
+
|
| 48 |
+
# Compute arguments.
|
| 49 |
+
memory_optimized_mlp: bool = False
|
| 50 |
+
mlp_type: str = 'mlp'
|
| 51 |
+
mlp_impl: str = 'sparse'
|
| 52 |
+
|
| 53 |
+
# Initialization arguments.
|
| 54 |
+
fp16: bool = True
|
| 55 |
+
bf16: bool = False
|
| 56 |
+
device: Union[int, torch.device] = dataclasses.field(default_factory=torch.cuda.current_device)
|
| 57 |
+
init_method: InitFn = partial(torch.nn.init.normal_, mean=0.0, std=0.02)
|
| 58 |
+
output_layer_init_method: InitFn = init_method
|
| 59 |
+
|
| 60 |
+
# Benchmarking arguments.
|
| 61 |
+
uniform_expert_assignment: bool = False
|
| 62 |
+
|
| 63 |
+
# shared expert arguments
|
| 64 |
+
shared_expert: bool = False # enable using shared expert
|
| 65 |
+
fc_cls: Any = torch.nn.Linear # class of the fully connected layer in shared expert (purpose: to allow using custom FC layer eg te.Linear (for FP8))
|
| 66 |
+
fc_kwargs: dict[str, Any] = dataclasses.field(default_factory=dict,) # kwargs for custom fc layers
|
| 67 |
+
remat_act_fn: bool = True # enable act fn to be rematerialized instead of stored
|
| 68 |
+
shared_expert_hidden_size: Optional[
|
| 69 |
+
int] = None # hidden size of the shared expert IF we want to set it to something different from hidden_size
|
| 70 |
+
shared_expert_weighted_sum: bool = False # enable using weighted sum for shared expert output (wieghted by number of experts used)
|
| 71 |
+
|
| 72 |
+
# Router Z-loss arguments
|
| 73 |
+
moe_zloss_weight: float = 0 # 1e-3 is a reasonable value
|
| 74 |
+
moe_zloss_in_fp32: bool = False
|
| 75 |
+
|
| 76 |
+
def __post_init__(self):
|
| 77 |
+
# Sparse MLP is not supported with triton >=3.2.0
|
| 78 |
+
# TODO: Remove this once sparse is supported with triton >=3.2.0
|
| 79 |
+
if self.__getattribute__('mlp_impl') == 'sparse':
|
| 80 |
+
try:
|
| 81 |
+
import triton
|
| 82 |
+
if triton.__version__ >= '3.2.0':
|
| 83 |
+
raise ValueError(
|
| 84 |
+
'Sparse MLP is not supported with triton >=3.2.0. Please use mlp_impl="grouped" instead.',
|
| 85 |
+
)
|
| 86 |
+
except ImportError:
|
| 87 |
+
raise ImportError('Triton is required for sparse MLP implementation')
|
| 88 |
+
|
| 89 |
+
if self.__getattribute__('mlp_impl') == 'grouped':
|
| 90 |
+
grouped_gemm.assert_grouped_gemm_is_available()
|
| 91 |
+
|
| 92 |
+
if self.shared_expert_hidden_size is None:
|
| 93 |
+
self.shared_expert_hidden_size = self.ffn_hidden_size
|
| 94 |
+
|
| 95 |
+
|
| 96 |
+
def from_megatron(megatron_args: Any):
|
| 97 |
+
args = Arguments()
|
| 98 |
+
for field in dataclasses.fields(args):
|
| 99 |
+
if hasattr(megatron_args, field.name):
|
| 100 |
+
setattr(args, field.name, getattr(megatron_args, field.name))
|
| 101 |
+
return args
|
build/torch214-cxx11-cu126-x86_64-linux/_layers/common.py
ADDED
|
@@ -0,0 +1,26 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2024 Databricks
|
| 2 |
+
# SPDX-License-Identifier: Apache-2.0
|
| 3 |
+
|
| 4 |
+
import torch
|
| 5 |
+
|
| 6 |
+
from .arguments import Arguments
|
| 7 |
+
|
| 8 |
+
|
| 9 |
+
def dtype(args: Arguments):
|
| 10 |
+
if args.fp16:
|
| 11 |
+
return torch.float16
|
| 12 |
+
elif args.bf16:
|
| 13 |
+
return torch.bfloat16
|
| 14 |
+
return None
|
| 15 |
+
|
| 16 |
+
|
| 17 |
+
def cast_if_autocast_enabled(tensor):
|
| 18 |
+
if torch.is_autocast_enabled():
|
| 19 |
+
if tensor.device.type == 'cuda':
|
| 20 |
+
dtype = torch.get_autocast_gpu_dtype()
|
| 21 |
+
elif tensor.device.type == 'cpu':
|
| 22 |
+
dtype = torch.get_autocast_cpu_dtype()
|
| 23 |
+
else:
|
| 24 |
+
raise NotImplementedError()
|
| 25 |
+
return tensor.to(dtype=dtype)
|
| 26 |
+
return tensor
|