diff --git a/device_lib/amd-gpu-gfx942.json b/device_lib/amd-gpu-gfx942.json index c3f5b8ab..74e8849c 100644 --- a/device_lib/amd-gpu-gfx942.json +++ b/device_lib/amd-gpu-gfx942.json @@ -1,1059 +1,402 @@ { - "name": "CDNA 3", - "vendor": "AMD", - "generation": 3, - "releaseYear": 2023, - "fabricationProcess": { - "processNode": 5, - "manufacturer": "TSMC", - "technology": "N5" + "Name": "CDNA 3", + "Vendor": "AMD", + "Architecture": "gfx942", + "ReleaseYear": 2023, + "FabricationProcess": { + "ProcessNode": 5, + "Manufacturer": "TSMC", + "Technology": "N5 for the accelerator complex dies (XCDs); the four I/O dies (IODs) are built on N6 and the package uses 3.5D die stacking with hybrid bonding on a CoWoS interposer" }, - "coreSubsystem": { - "name": "Compute Unit (CU)", - "maxPerUnit": 304, - "subUnits": [ - { - "type": "SIMDCore", - "countPerUnit": 4, - "size": 64, - "description": "64-wide SIMD vector processor" + "CoreSubsystem": { + "Name": "Compute Unit (CU)", + "ChipType": "Chip", + "CoreType": "Compute Unit", + "UnitTypes": { + "Chip": { + "Count": 1, + "Size": 304, + "Description": "Primary subject of this description is the AMD Instinct MI300X discrete GPU (OAM module): 8 XCDs stacked over 4 IODs with 8 HBM3 stacks, 304 active compute units, 1216 Matrix Cores, 19,456 stream processors, 2100 MHz peak engine clock, 750 W maximum TBP, and roughly 153 billion transistors. The MI300A APU shares the same gfx942 ISA but substitutes 3 'Zen 4' CCDs (24 x86 cores) for 2 of the XCDs, leaving 6 XCDs / 228 active CUs / 912 Matrix Cores / 14,592 stream processors, 128 GB of HBM3 unified coherently with the CPU, and a 550 W TDP (760 W in liquid-cooled configurations). Both parts expose 256 MB of AMD Infinity Cache and 5.3 TB/s of peak HBM3 bandwidth", + "Memory": ["HBM3 (Device Memory)", "AMD Infinity Cache (MALL, per package)"], + "Subunits": ["XCD (Accelerator Complex Die)", "IOD (I/O Die)", "HBM3 Stack"] }, - { - "type": "MatrixCore", - "countPerUnit": 4, - "size": 32, - "description": "Dedicated matrix multiplication unit for AI acceleration" + "XCD (Accelerator Complex Die)": { + "Count": 8, + "Size": 38, + "Description": "Accelerator complex die on TSMC N5. Each XCD physically contains 40 CUs of which 38 are active (2 disabled for yield), grouped under 4 Asynchronous Compute Engines and a hardware scheduler, all sharing one 4 MB L2 cache. MI300X uses 8 XCDs (304 CUs); MI300A uses 6 (228 CUs). XCDs are the unit of spatial partitioning: MI300X supports SPX/DPX/QPX/CPX modes down to one partition per XCD, MI300A supports up to 3 partitions of 2 XCDs", + "Memory": ["L2 Cache (per XCD)"], + "Subunits": ["Compute Unit", "Asynchronous Compute Engine (ACE)", "Hardware Scheduler (HWS)"] + }, + "Compute Unit": { + "Count": 38, + "Size": 64, + "Description": "CDNA 3 compute unit: instruction fetch and scheduling, a scalar unit, 4 SIMD units totalling 64 stream processors, 4 Matrix Cores, load/store pipelines, 64 KB of Local Data Share and a 32 KB L1 vector data cache. Executes wave64 wavefronts; up to 16 wavefronts (1024 work-items) per work-group", + "Memory": ["Vector Register File (VGPR, per SIMD)", "Scalar Register File (SGPR)", "LDS (Local Data Share, per CU)", "L1 Vector Data Cache (per CU)", "Instruction Cache (shared by 2 CUs)"], + "Subunits": ["SIMD Unit", "Matrix Core"] + }, + "SIMD Unit": { + "Count": 4, + "Size": 16, + "Description": "16-lane vector ALU; a 64 work-item wavefront is issued across 4 cycles. CDNA 3 doubles vector FP32 throughput over CDNA 2 with packed FP32 operations (packed-fp32-ops), reaching 256 FLOPs/clock/CU for vector FP32 and 128 FLOPs/clock/CU for vector FP64" + }, + "Matrix Core": { + "Count": 4, + "Size": 1, + "Description": "Matrix multiply-accumulate unit driven by the MFMA/SMFMAC instruction families. Per CU per clock: 256 FLOPs matrix FP64, 256 matrix FP32, 1024 TF32 (XF32), 2048 FP16, 2048 BF16, 4096 FP8 (E4M3 and E5M2), 4096 INT8. Structured sparsity (at least 2 zeros in every group of 4 values) doubles matrix INT8/FP8/FP16/BF16 throughput to 8192 operations per clock per CU" + }, + "Asynchronous Compute Engine (ACE)": { + "Count": 4, + "Size": 1, + "Description": "Compute accelerator that dispatches compute-shader workgroups to the CUs of its XCD; each ACE is associated with 10 of the 40 physical CUs and drives 8 hardware queues" + }, + "Hardware Scheduler (HWS)": { + "Count": 1, + "Size": 1, + "Description": "Per-XCD global resource that schedules work across the four ACEs and their hardware queues" + }, + "IOD (I/O Die)": { + "Count": 4, + "Size": 1, + "Description": "Active I/O die on TSMC N6, vertically stacked beneath a pair of XCDs. Hosts the AMD Infinity Cache slices, the HBM3 memory controllers for 2 stacks, the AMD Infinity Fabric network, and the media engines", + "Memory": ["AMD Infinity Cache (MALL, per package)"], + "Subunits": ["Infinity Fabric Link", "Video Decoder", "HBM3 Memory Controller"] + }, + "Infinity Fabric Link": { + "Count": 2, + "Size": 1, + "Description": "Two 16-lane bidirectional inter-package AMD Infinity Fabric links per IOD (8 per package), signalling at up to 32 Gbps for 128 GB/s bidirectional each. One link per IOD is multi-purpose and can instead be configured as x16 PCIe Gen 5. MI300X uses seven as scale-up Infinity Fabric links (896 GB/s aggregate, fully connecting 8 GPUs on an OCP UBB 2.0 board) plus one x16 PCIe Gen 5 host link; MI300A dedicates 4 to Infinity Fabric and leaves 4 assignable to Infinity Fabric or PCIe Gen 5" + }, + "HBM3 Memory Controller": { + "Count": 2, + "Size": 1, + "Description": "Each IOD fans out through the package to 2 HBM3 stacks; the controllers drive the bus at 5.2 Gbps" + }, + "Video Decoder": { + "Count": 1, + "Size": 1, + "Description": "Decode-only media engine for HEVC/H.265, AVC/H.264, VP9 or AV1. MI300X has 4 such decoder groups, MI300A has 3. No video encoders are present on CDNA 3", + "Subunits": ["JPEG/MJPEG Codec Core"] + }, + "JPEG/MJPEG Codec Core": { + "Count": 8, + "Size": 1, + "Description": "8 JPEG/MJPEG codec cores accompany each decoder group, for 32 cores on MI300X and 24 on MI300A" + }, + "HBM3 Stack": { + "Count": 8, + "Size": 1, + "Description": "8 HBM3 stacks per package, 24 GB each on MI300X (192 GB total) and 16 GB each on MI300A (128 GB total). Each stack is associated with 16 Infinity Cache channels" } - ] - }, - "shaderModel": { - "directX": "N/A", - "vulkan": "N/A", - "openGL": "N/A", - "openCL": "3.0", - "metal": "N/A" + } }, - "memorySubsystem": { - "supportedMemoryTypes": ["HBM3"], - "maxMemoryBandwidth": 5300, - "maxMemorySize": 192, - "maxBusWidth": 8192, - "cacheHierarchy": [ + "MemorySubsystem": { + "SupportedMemoryTypes": ["HBM3"], + "MemoryTypes": [ + { + "Type": "Vector Register File (VGPR, per SIMD)", + "Size": 128, + "description": "A wavefront may allocate up to 512 VGPRs, split flexibly between at most 256 architectural VGPRs and at most 256 accumulation VGPRs (AGPRs) used by the matrix VALU instructions; 64-bit operands and AGPR tuples require even alignment (LLVM vgpr-align2). 512 registers x 64 lanes x 4 bytes gives the 128 KB per-SIMD figure recorded here (512 KB per CU); AMD did not publish a change to the physical register file capacity between CDNA 2 and CDNA 3" + }, + { + "Type": "Scalar Register File (SGPR)", + "description": "A wavefront is allocated 16 to 102 SGPRs in units of 16 dwords, with VCC stored in the top two and 16 further SGPRs reserved when a trap handler is present. AMD does not publish the physical SGPR file capacity for CDNA 3, so no size is recorded" + }, + { + "Type": "LDS (Local Data Share, per CU)", + "Size": 64, + "BankCount": 32, + "description": "64 KB of software-managed on-chip memory per compute unit, organized as 32 banks of 512 dwords with 32 integer atomic units. Shared by the work-items of a work-group; a work-group may request the full 64 KB. Reads across a wavefront are dispatched over four cycles. Unchanged from CDNA 2" + }, + { + "Type": "L1 Vector Data Cache (per CU)", + "Size": 32, + "description": "Per-CU L1 vector data cache. CDNA 3 doubles the line size to 128 B and the capacity to 32 KB versus CDNA 2, and widens both the request bus to the core and the fill path from L2 to match, doubling bandwidth to the core. Relaxed coherency model requiring explicit synchronization" + }, { - "level": "L1 Vector Cache", - "sizePerUnit": 16, - "totalSize": 3648, - "description": "L1 cache for vector operations per CU" + "Type": "Instruction Cache (shared by 2 CUs)", + "Size": 64, + "description": "64 KB 8-way set-associative instruction cache shared between two compute units, doubling both capacity and associativity versus CDNA 2" }, { - "level": "L1 Scalar Cache", - "sizePerUnit": 16, - "totalSize": 3648, - "description": "L1 cache for scalar operations per CU" + "Type": "L2 Cache (per XCD)", + "Size": 4096, + "BankCount": 16, + "MaxMemoryBandwidth": 34400, + "description": "4 MB 16-way set-associative writeback, write-allocate cache shared by the 38 active CUs of an XCD, built from 16 channels of 256 KB. Each channel reads out a 128 B line and the cache sustains four requests from different CUs per cycle for 2 KB/clock per XCD; writes are half-line (64 B) per channel. Lowest level at which coherency is maintained in hardware. The recorded bandwidth is the whitepaper's 34.4 TB/s aggregate read bandwidth across all 8 XCD L2 instances (its memory-hierarchy figure legend instead labels the L2-to-XCD path 51.6 TB/s)" }, { - "level": "L2 Cache", - "sizePerUnit": 4096, - "totalSize": 98304, - "description": "Shared L2 cache across multiple CUs" + "Type": "AMD Infinity Cache (MALL, per package)", + "Size": 262144, + "BankCount": 128, + "MaxMemoryBandwidth": 17200, + "description": "256 MB 16-way set-associative memory-side last level cache distributed across the 4 IODs: 16 channels per HBM3 stack, 128 channels total, each 64 B wide and backed by 2 MB of banked data array, for 17.2 TB/s peak bandwidth. Being memory-side it holds no dirty data evicted from lower levels, does not participate in coherency or absorb snoop traffic, and can cache nominally uncacheable memory; it does contain a snoop filter covering the XCD L2 caches. On MI300A it is shared between the XCDs and the 'Zen 4' CPU dies" }, { - "level": "L3 Cache", - "sizePerUnit": 0, - "totalSize": 131072, - "description": "Shared last-level cache" + "Type": "HBM3 (Device Memory)", + "Size": 201326592, + "BankCount": 8, + "MaxMemoryBandwidth": 5300, + "maxBusWidth": 8192, + "description": "192 GB of HBM3 on MI300X across 8 stacks of 24 GB, on an 8192-bit interface clocked at 5.2 Gbps for 5.3 TB/s peak theoretical bandwidth. MI300A carries 128 GB (8 stacks of 16 GB) at the same 5.3 TB/s, unified and cache-coherent between the CPU and GPU dies in a single physical and virtual address space. Full-chip ECC with page retirement and page avoidance" } ] }, - "specializedHardware": { - "rayTracingAccelerators": { - "present": false, - "name": null, - "countPerComputeUnit": 0, - "performance": { - "raysPerSecond": 0 - } - }, - "aiAccelerators": { - "present": true, - "name": "Matrix Core", - "countPerComputeUnit": 4, - "supportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "INT4"], - "performance": { - "fp16TopsPerGPU": 1638, - "int8TopsPerGPU": 3276 - } - }, - "videoCodecs": { - "encoders": [], - "decoders": [] - } - }, - "powerEfficiency": { - "maxTDP": 750, - "powerStates": [ + "KernelModel": { + "LLVMTarget": "amdgcn", + "LLVMTriple": "amdgcn-amd-amdhsa", + "LLVMFeatures": [ { - "name": "Performance Mode", - "description": "Maximum performance for compute-intensive workloads" + "Name": "gfx9-insts", + "Description": "Additional instructions for GFX9+; gfx942 is ISA version 9.4.2 in the GCN GFX9 encoding family" }, { - "name": "Balanced Mode", - "description": "Optimized balance between performance and power consumption" + "Name": "gfx90a-insts", + "Description": "Additional instructions introduced with GFX90A (CDNA 2)" }, { - "name": "Power Saving Mode", - "description": "Reduced performance for power-constrained environments" - } - ], - "clockGating": true, - "dynamicVoltageFrequencyScaling": true - }, - "displayOutputs": { - "maxDisplays": 0, - "maxResolution": "N/A", - "maxRefreshRate": 0, - "interfaces": [], - "hdr": false - }, - "pciExpress": { - "version": "5.0", - "lanes": 128, - "bandwidth": 256 - }, - "multiGpuSupport": { - "technologies": ["Infinity Fabric"], - "maxGpus": 8, - "interconnectBandwidth": 800 - }, - "softwareFeatures": { - "upscalingTechnologies": [], - "meshShading": false, - "variableRateShading": false, - "samplerFeedback": false - }, - "additionalFeatures": { - "chipletDesign": true, - "xgmiLinks": 12, - "cpuIntegration": { - "present": true, - "description": "Integrated with Zen 4 CPU cores in APU design" - }, - "memoryCoherence": true, - "ecc": true, - "secureBoot": true, - "isa": "gfx942", - "instructionExtensions": ["MFMA", "WMMA", "DPAS"], - "roiAccelerators": true, - "virtualization": { - "present": true, - "technologies": ["SR-IOV", "MxGPU"] - }, - "fp64Performance": { - "peakTeraflops": 102.4, - "ratiotoFP32": 0.5 - } - }, - "target": { - "properties": { - "target_name": "gfx942", - "architecture": "amdgcn", - "generation": "GFX9", - "family": "CDNA2", - "llvm_target": "AMDGPU", - "isa_version": { - "major": 9, - "minor": 4, - "stepping": 2 - }, - "elf_machine_type": "EF_AMDGPU_MACH_AMDGCN_GFX942", - "elf_machine_code": "0x04c" - }, - "compute_capabilities": { - "wavefront_size": 64, - "max_workgroup_size": 1024, - "local_memory_size": 65536, - "supports_generic_address_space": true, - "address_space_pointer_size": 64, - "supports_flat_instructions": true, - "supports_architected_flat_scratch": true, - "supports_architected_sgprs": true - }, - "target_features": { - "16_bit_insts": true, - "atomic_buffer_global_pk_add_f16_insts": true, - "atomic_ds_pk_add_16_insts": true, - "atomic_fadd_rtn_insts": true, - "atomic_flat_pk_add_16_insts": true, - "atomic_global_pk_add_bf16_inst": true, - "atomic_buffer_pk_add_bf16_inst": true, - "ci_insts": true, - "dl_insts": true, - "dot1_insts": true, - "dot2_insts": true, - "dot3_insts": true, - "dot4_insts": true, - "dot5_insts": true, - "dot6_insts": true, - "dot7_insts": true, - "dot8_insts": true, - "dot9_insts": true, - "dot10_insts": true, - "dpp": true, - "fp8_insts": true, - "gfx8_insts": true, - "gfx9_insts": true, - "gfx90a_insts": true, - "gfx940_insts": true, - "mai_insts": true, - "s_memrealtime": true, - "s_memtime_inst": true, - "wavefrontsize64": true, - "packed_tid": true, - "image_insts": true, - "extended_image_insts": true, - "fp8_conversion_insts": true, - "vcmpx_permlane_hazard": true, - "salu_float_insts": true, - "pseudo_scalar_trans": true, - "has_restricted_soffset": true, - "scalar_dwordx3_loads": true, - "dpp_src1_sgpr": true, - "max_hard_clause_length_32": true, - "1_5x_vgprs": true, - "memory_atomic_fadd_f32_denormal_support": true, - "bvh_dual_and_bvh8_insts": true - }, - "instruction_sets": { - "ds_instructions": { - "supported": true, - "atomic_operations": true, - "floating_point_atomics": true, - "packed_operations": true, - "examples": [ - "ds_add_f32", - "ds_add_f64", - "ds_pk_add_bf16", - "ds_pk_add_f16", - "ds_read_b32", - "ds_write_b32", - "ds_cmpst_f32", - "ds_max_f32" - ] - }, - "flat_instructions": { - "supported": true, - "global_memory": true, - "scratch_memory": true, - "atomic_operations": true, - "examples": [ - "flat_load_dword", - "flat_store_dword", - "flat_atomic_add", - "flat_atomic_add_f32", - "flat_atomic_pk_add_f16", - "global_load_dword", - "scratch_load_dword" - ] - }, - "buffer_instructions": { - "supported": true, - "typed_buffer_loads": true, - "format_conversion": true, - "atomic_operations": true, - "examples": [ - "buffer_load_dword", - "buffer_store_dword", - "buffer_atomic_add", - "buffer_atomic_add_f32", - "buffer_load_format_x" - ] - }, - "scalar_memory_instructions": { - "supported": true, - "cache_control": true, - "atomic_operations": true, - "examples": [ - "s_load_dword", - "s_store_dword", - "s_atomic_add", - "s_memtime", - "s_memrealtime" - ] - }, - "vector_alu_instructions": { - "vop1": { - "supported": true, - "floating_point": true, - "integer": true, - "conversion": true, - "examples": [ - "v_mov_b32", - "v_cvt_f32_i32", - "v_sqrt_f32", - "v_rcp_f32", - "v_cvt_f32_bf8", - "v_cvt_f32_fp8" - ] - }, - "vop2": { - "supported": true, - "arithmetic": true, - "comparison": true, - "packed_operations": true, - "examples": [ - "v_add_f32", - "v_mul_f32", - "v_dot2c_f32_f16", - "v_dot4c_i32_i8", - "v_pk_fmac_f16" - ] - }, - "vop3": { - "supported": true, - "three_operand": true, - "modifiers": true, - "comparison": true, - "examples": [ - "v_mad_f32", - "v_fma_f32", - "v_cmp_lt_f32", - "v_div_fixup_f32", - "v_dot2_f32_f16" - ] - }, - "vop3p": { - "supported": true, - "packed_operations": true, - "matrix_operations": true, - "mixed_precision": true, - "examples": [ - "v_pk_add_f16", - "v_pk_mul_f16", - "v_mfma_f32_16x16x4_f32", - "v_mfma_f32_32x32x8_f16", - "v_dot2_f32_f16" - ] - } + "Name": "gfx940-insts", + "Description": "Additional instructions for GFX940+ (CDNA 3)" }, - "scalar_alu_instructions": { - "sop1": { - "supported": true, - "examples": [ - "s_mov_b32", - "s_not_b32", - "s_brev_b32" - ] - }, - "sop2": { - "supported": true, - "examples": [ - "s_add_u32", - "s_and_b32", - "s_lshl_b32" - ] - }, - "sopc": { - "supported": true, - "examples": [ - "s_cmp_eq_i32", - "s_cmp_gt_i32" - ] - }, - "sopk": { - "supported": true, - "examples": [ - "s_movk_i32", - "s_addk_i32" - ] - }, - "sopp": { - "supported": true, - "examples": [ - "s_nop", - "s_endpgm", - "s_barrier" - ] - } - } - }, - "matrix_operations": { - "mfma_instructions": { - "supported": true, - "encoding": "VOP3P-MAI", - "wavefront_execution": true, - "precision_support": { - "f32": true, - "f16": true, - "bf16": true, - "i8": true, - "f64": true, - "fp8": true, - "bf8": true - }, - "performance_characteristics": { - "preferred_size": "16x16", - "power_efficiency": "mfma_16x16 > mfma_32x32", - "clock_variation_across_xcds": "3-10%", - "hotspotting_avoidance": "avoid matrix stride multiples of 512 bytes" - }, - "detailed_instructions": { - "v_mfma_f32_4x4x1f32": { - "matrix_dimensions": {"M": 4, "N": 4, "K": 1, "blocks": 16}, - "flops": 512, - "execution_cycles": 8, - "flops_per_cu_per_cycle": 256, - "can_coexecute_with_valu": true, - "valu_coexecution_cycles": 4, - "register_usage": { - "gprs_required_a": 1, - "gprs_required_b": 1, - "gprs_required_c": 4, - "gprs_required_d": 4, - "gpr_alignment": "8_bytes" - }, - "issue_latency_cycles": 1, - "data_latency_cycles": 8, - "throughput_cycles": 8 - }, - "v_mfma_f32_16x16x1f32": { - "matrix_dimensions": {"M": 16, "N": 16, "K": 1, "blocks": 4}, - "flops": 1024, - "execution_cycles": 16, - "flops_per_cu_per_cycle": 256, - "can_coexecute_with_valu": true, - "valu_coexecution_cycles": 8, - "register_usage": { - "gprs_required_a": 1, - "gprs_required_b": 1, - "gprs_required_c": 16, - "gprs_required_d": 16, - "gpr_alignment": "8_bytes" - }, - "issue_latency_cycles": 1, - "data_latency_cycles": 16, - "throughput_cycles": 16 - }, - "v_mfma_f32_16x16x4f32": { - "matrix_dimensions": {"M": 16, "N": 16, "K": 4, "blocks": 1}, - "flops": 2048, - "execution_cycles": 32, - "flops_per_cu_per_cycle": 256, - "can_coexecute_with_valu": true, - "valu_coexecution_cycles": 16, - "register_usage": { - "gprs_required_a": 4, - "gprs_required_b": 4, - "gprs_required_c": 16, - "gprs_required_d": 16, - "gpr_alignment": "8_bytes" - }, - "issue_latency_cycles": 1, - "data_latency_cycles": 32, - "throughput_cycles": 32 - }, - "v_mfma_f32_16x16x16f16": { - "matrix_dimensions": {"M": 16, "N": 16, "K": 16, "blocks": 1}, - "flops": 8192, - "execution_cycles": 32, - "flops_per_cu_per_cycle": 1024, - "can_coexecute_with_valu": true, - "valu_coexecution_cycles": 16, - "register_usage": { - "gprs_required_a": 8, - "gprs_required_b": 8, - "gprs_required_c": 16, - "gprs_required_d": 16, - "gpr_alignment": "8_bytes" - }, - "issue_latency_cycles": 1, - "data_latency_cycles": 32, - "throughput_cycles": 32, - "preferred": true - }, - "v_mfma_f32_32x32x1f32": { - "matrix_dimensions": {"M": 32, "N": 32, "K": 1, "blocks": 1}, - "flops": 2048, - "execution_cycles": 64, - "flops_per_cu_per_cycle": 128, - "can_coexecute_with_valu": true, - "valu_coexecution_cycles": 32, - "register_usage": { - "gprs_required_a": 1, - "gprs_required_b": 1, - "gprs_required_c": 32, - "gprs_required_d": 32, - "gpr_alignment": "8_bytes" - }, - "issue_latency_cycles": 1, - "data_latency_cycles": 64, - "throughput_cycles": 64 - }, - "v_mfma_f32_32x32x8f16": { - "matrix_dimensions": {"M": 32, "N": 32, "K": 8, "blocks": 1}, - "flops": 16384, - "execution_cycles": 64, - "flops_per_cu_per_cycle": 1024, - "can_coexecute_with_valu": true, - "valu_coexecution_cycles": 32, - "register_usage": { - "gprs_required_a": 8, - "gprs_required_b": 8, - "gprs_required_c": 32, - "gprs_required_d": 32, - "gpr_alignment": "8_bytes" - }, - "issue_latency_cycles": 1, - "data_latency_cycles": 64, - "throughput_cycles": 64 - }, - "v_mfma_i32_16x16x32i8": { - "matrix_dimensions": {"M": 16, "N": 16, "K": 32, "blocks": 1}, - "ops": 16384, - "execution_cycles": 32, - "ops_per_cu_per_cycle": 2048, - "can_coexecute_with_valu": true, - "valu_coexecution_cycles": 16, - "register_usage": { - "gprs_required_a": 8, - "gprs_required_b": 8, - "gprs_required_c": 16, - "gprs_required_d": 16, - "gpr_alignment": "8_bytes" - }, - "issue_latency_cycles": 1, - "data_latency_cycles": 32, - "throughput_cycles": 32 - }, - "v_mfma_f64_16x16x4f64": { - "matrix_dimensions": {"M": 16, "N": 16, "K": 4, "blocks": 1}, - "flops": 8192, - "execution_cycles": 64, - "flops_per_cu_per_cycle": 512, - "can_coexecute_with_valu": false, - "register_usage": { - "gprs_required_a": 16, - "gprs_required_b": 16, - "gprs_required_c": 64, - "gprs_required_d": 64, - "gpr_alignment": "8_bytes" - }, - "issue_latency_cycles": 1, - "data_latency_cycles": 64, - "throughput_cycles": 64 - } - }, - "control_modifiers": { - "cbsz": { - "description": "Control Broadcast Size modifier", - "function": "broadcasts input values to neighboring blocks", - "supported_values": [0, 1, 2, 3], - "default": 0 - }, - "abid": { - "description": "A Block ID modifier", - "function": "selects which input block to broadcast", - "works_with": "cbsz", - "supported_values": [0, 1, 2, 3] - }, - "blgp": { - "description": "B Lane Group Pattern modifier", - "function": "controls data layout pattern for B matrix", - "supported_values": [0, 1, 2, 3] - } - } + { + "Name": "mai-insts", + "Description": "Matrix Arithmetic Instructions: the V_MFMA family executed on the Matrix Cores" }, - "smfmac_instructions": { - "supported": true, - "sparse_matrix": true, - "sparsity_support": "2x improvement in math efficiency", - "encoding": "VOP3P", - "examples": { - "v_smfmac_f32_16x16x32_f16": { - "matrix_dimensions": {"M": 16, "N": 16, "K": 32}, - "sparsity_pattern": "structured_sparse", - "issue_latency_cycles": 1, - "data_latency_cycles": 32, - "throughput_cycles": 32 - }, - "v_smfmac_f32_32x32x16_bf16": { - "matrix_dimensions": {"M": 32, "N": 32, "K": 16}, - "sparsity_pattern": "structured_sparse", - "issue_latency_cycles": 1, - "data_latency_cycles": 64, - "throughput_cycles": 64 - }, - "v_smfmac_i32_16x16x64_i8": { - "matrix_dimensions": {"M": 16, "N": 16, "K": 64}, - "sparsity_pattern": "structured_sparse", - "issue_latency_cycles": 1, - "data_latency_cycles": 32, - "throughput_cycles": 32 - } - } - } - }, - "memory_operations": { - "load_store_instructions": { - "buffer_operations": { - "buffer_load_dword": { - "description": "Load 32-bit word from buffer", - "issue_latency_cycles": 1, - "data_latency_cycles": 150, - "throughput_cycles": 1, - "cache_behavior": "cached_by_default", - "fifo_depth": 32, - "supports_modifiers": ["sc0", "nt", "sc1", "offset12"] - }, - "buffer_load_dwordx2": { - "description": "Load 64-bit double word from buffer", - "issue_latency_cycles": 1, - "data_latency_cycles": 150, - "throughput_cycles": 1, - "cache_behavior": "cached_by_default", - "fifo_depth": 32 - }, - "buffer_load_dwordx4": { - "description": "Load 128-bit quad word from buffer", - "issue_latency_cycles": 1, - "data_latency_cycles": 150, - "throughput_cycles": 1, - "cache_behavior": "cached_by_default", - "fifo_depth": 32 - }, - "buffer_store_dword": { - "description": "Store 32-bit word to buffer", - "issue_latency_cycles": 1, - "data_latency_cycles": 4, - "throughput_cycles": 1, - "cache_behavior": "write_through", - "fifo_depth": 32 - }, - "buffer_store_dwordx2": { - "description": "Store 64-bit double word to buffer", - "issue_latency_cycles": 1, - "data_latency_cycles": 4, - "throughput_cycles": 1, - "cache_behavior": "write_through", - "fifo_depth": 32 - }, - "buffer_store_dwordx4": { - "description": "Store 128-bit quad word to buffer", - "issue_latency_cycles": 1, - "data_latency_cycles": 4, - "throughput_cycles": 1, - "cache_behavior": "write_through", - "fifo_depth": 32 - }, - "buffer_atomic_add": { - "description": "Atomic add operation on buffer", - "issue_latency_cycles": 1, - "data_latency_cycles": 200, - "throughput_cycles": 1, - "cache_behavior": "bypass_cache", - "fifo_depth": 16 - }, - "buffer_atomic_add_f32": { - "description": "Atomic float32 add operation on buffer", - "issue_latency_cycles": 1, - "data_latency_cycles": 200, - "throughput_cycles": 1, - "cache_behavior": "bypass_cache", - "fifo_depth": 16 - }, - "buffer_load_format_x": { - "description": "Load formatted data from buffer", - "issue_latency_cycles": 1, - "data_latency_cycles": 160, - "throughput_cycles": 1, - "cache_behavior": "cached_by_default", - "fifo_depth": 32, - "supports_lds_load": true - } - }, - "flat_operations": { - "flat_load_dword": { - "description": "Load 32-bit word via flat addressing", - "issue_latency_cycles": 1, - "data_latency_cycles": 150, - "throughput_cycles": 1, - "cache_behavior": "cached_by_default", - "fifo_depth": 32, - "address_space": "generic" - }, - "flat_store_dword": { - "description": "Store 32-bit word via flat addressing", - "issue_latency_cycles": 1, - "data_latency_cycles": 4, - "throughput_cycles": 1, - "cache_behavior": "write_through", - "fifo_depth": 32, - "address_space": "generic" - }, - "global_load_dword": { - "description": "Load from global memory", - "issue_latency_cycles": 1, - "data_latency_cycles": 150, - "throughput_cycles": 1, - "cache_behavior": "cached_by_default", - "fifo_depth": 32, - "address_space": "global" - }, - "global_store_dword": { - "description": "Store to global memory", - "issue_latency_cycles": 1, - "data_latency_cycles": 4, - "throughput_cycles": 1, - "cache_behavior": "write_through", - "fifo_depth": 32, - "address_space": "global" - }, - "scratch_load_dword": { - "description": "Load from scratch memory", - "issue_latency_cycles": 1, - "data_latency_cycles": 120, - "throughput_cycles": 1, - "cache_behavior": "cached_by_default", - "fifo_depth": 32, - "address_space": "private" - }, - "scratch_store_dword": { - "description": "Store to scratch memory", - "issue_latency_cycles": 1, - "data_latency_cycles": 4, - "throughput_cycles": 1, - "cache_behavior": "write_through", - "fifo_depth": 32, - "address_space": "private" - } - }, - "ds_operations": { - "ds_read_b32": { - "description": "Read 32-bit from Local Data Share", - "issue_latency_cycles": 1, - "data_latency_cycles": 20, - "throughput_cycles": 1, - "cache_behavior": "lds_direct", - "fifo_depth": 64, - "address_space": "local" - }, - "ds_write_b32": { - "description": "Write 32-bit to Local Data Share", - "issue_latency_cycles": 1, - "data_latency_cycles": 2, - "throughput_cycles": 1, - "cache_behavior": "lds_direct", - "fifo_depth": 64, - "address_space": "local" - }, - "ds_add_rtn_u32": { - "description": "Atomic add with return from LDS", - "issue_latency_cycles": 1, - "data_latency_cycles": 25, - "throughput_cycles": 1, - "cache_behavior": "lds_direct", - "fifo_depth": 32, - "address_space": "local" - } - }, - "scalar_memory_operations": { - "s_load_dword": { - "description": "Scalar load 32-bit word", - "issue_latency_cycles": 1, - "data_latency_cycles": 100, - "throughput_cycles": 1, - "cache_behavior": "cached_by_default", - "fifo_depth": 16 - }, - "s_store_dword": { - "description": "Scalar store 32-bit word", - "issue_latency_cycles": 1, - "data_latency_cycles": 4, - "throughput_cycles": 1, - "cache_behavior": "write_through", - "fifo_depth": 16 - }, - "s_atomic_add": { - "description": "Scalar atomic add", - "issue_latency_cycles": 1, - "data_latency_cycles": 150, - "throughput_cycles": 1, - "cache_behavior": "bypass_cache", - "fifo_depth": 8 - } - } + { + "Name": "fp8-insts", + "Description": "Has fp8 (E4M3) and bf8 (E5M2) instructions, including the FP8 MFMA variants" }, - "fifo_specifications": { - "vector_memory_fifo": { - "depth": 32, - "width": 128, - "description": "Main vector memory operation queue", - "supports_reordering": true, - "max_outstanding_requests": 32 - }, - "scalar_memory_fifo": { - "depth": 16, - "width": 64, - "description": "Scalar memory operation queue", - "supports_reordering": false, - "max_outstanding_requests": 16 - }, - "lds_fifo": { - "depth": 64, - "width": 32, - "description": "Local Data Share operation queue", - "supports_reordering": true, - "max_outstanding_requests": 64 - }, - "atomic_fifo": { - "depth": 16, - "width": 64, - "description": "Atomic operation queue", - "supports_reordering": false, - "max_outstanding_requests": 16 - }, - "texture_fifo": { - "depth": 32, - "width": 128, - "description": "Texture/image operation queue", - "supports_reordering": true, - "max_outstanding_requests": 32 - } + { + "Name": "fp8-conversion-insts", + "Description": "Has fp8 and bf8 conversion instructions such as v_cvt_f32_fp8 and v_cvt_f32_bf8" }, - "cache_specifications": { - "l1_vector_cache": { - "size": "32KB", - "associativity": 8, - "line_size": 128, - "latency_cycles": 4, - "bandwidth_per_cycle": "128B", - "replacement_policy": "LRU" - }, - "l1_scalar_cache": { - "size": "16KB", - "associativity": 4, - "line_size": 64, - "latency_cycles": 4, - "bandwidth_per_cycle": "64B", - "replacement_policy": "LRU" - }, - "l2_cache": { - "size": "4MB", - "associativity": 16, - "line_size": 128, - "latency_cycles": 50, - "bandwidth_per_cycle": "256B", - "replacement_policy": "LRU" - }, - "infinity_cache": { - "size": "256MB", - "associativity": 16, - "line_size": 128, - "latency_cycles": 100, - "bandwidth_per_cycle": "1024B", - "replacement_policy": "LRU" - } + { + "Name": "cvt-fp8-sdwa-src-sel", + "Description": "FP8/BF8 conversions take the byte index from the SDWA src0_sel or op_sel field" }, - "special_features": { - "cache_swizzle_support": { - "introduced_in": "gfx942", - "description": "Enhanced cache swizzling for buffer fat pointers", - "improves": "memory_access_patterns" - }, - "buffer_fat_pointers": { - "supported": true, - "address_space": 7, - "descriptor_size": 128, - "offset_size": 32, - "total_pointer_size": 160 - } - } - }, - "address_spaces": { - "global": { - "id": 1, - "cacheable": true, - "coherent": true - }, - "local": { - "id": 3, - "shared_memory": true, - "size": 65536 - }, - "private": { - "id": 5, - "stack_memory": true - }, - "constant": { - "id": 4, - "read_only": true, - "cacheable": true - }, - "generic": { - "id": 0, - "flat_addressing": true + { + "Name": "xf32-insts", + "Description": "Has v_mfma_f32_16x16x8_xf32 and v_mfma_f32_32x32x4_xf32, the TF32 (XF32) matrix instructions new in CDNA 3" + }, + { + "Name": "packed-fp32-ops", + "Description": "Support packed fp32 instructions (v_pk_fma_f32, v_pk_mul_f32, v_pk_add_f32), which double vector FP32 throughput per CU versus CDNA 2" + }, + { + "Name": "full-rate-64-ops", + "Description": "Most fp64 instructions are full rate" + }, + { + "Name": "fmacf64-inst", + "Description": "Has v_fmac_f64 instruction" + }, + { + "Name": "pk-fmac-f16-inst", + "Description": "Has v_pk_fmac_f16 instruction" + }, + { + "Name": "fma-mix-insts", + "Description": "Has v_fma_mix_f32, v_fma_mixlo_f16 and v_fma_mixhi_f16 instructions" + }, + { + "Name": "dl-insts", + "Description": "Has v_fmac_f32 and v_xnor_b32 deep-learning instructions" + }, + { + "Name": "dot1-insts", + "Description": "Has v_dot4_i32_i8 and v_dot8_i32_i4 instructions" + }, + { + "Name": "dot2-insts", + "Description": "Has v_dot2_i32_i16 and v_dot2_u32_u16 instructions" + }, + { + "Name": "dot3-insts", + "Description": "Has v_dot8c_i32_i4 instruction" + }, + { + "Name": "dot4-insts", + "Description": "Has v_dot2c_i32_i16 instruction" + }, + { + "Name": "dot5-insts", + "Description": "Has v_dot2c_f32_f16 instruction" + }, + { + "Name": "dot6-insts", + "Description": "Has v_dot4c_i32_i8 instruction" + }, + { + "Name": "dot7-insts", + "Description": "Has v_dot4_u32_u8 and v_dot8_u32_u4 instructions" + }, + { + "Name": "dot10-insts", + "Description": "Has v_dot2_f32_f16 instruction" + }, + { + "Name": "dpp-64bit", + "Description": "Support DPP (Data Parallel Primitives) extension in the double-precision ALU" + }, + { + "Name": "atomic-fadd-rtn-insts", + "Description": "Has buffer/global/flat floating-point atomic add instructions that return the original value" + }, + { + "Name": "atomic-fadd-no-rtn-insts", + "Description": "Has floating-point atomic add instructions that do not return the original value" + }, + { + "Name": "flat-atomic-fadd-f32-inst", + "Description": "Has flat_atomic_add_f32 instruction" + }, + { + "Name": "flat-buffer-global-fadd-f64-inst", + "Description": "Has flat, buffer and global instructions for f64 atomic fadd" + }, + { + "Name": "memory-atomic-fadd-f32-denormal-support", + "Description": "global/flat/buffer atomic fadd for float supports denormal handling" + }, + { + "Name": "atomic-ds-pk-add-16-insts", + "Description": "Has ds_pk_add_f16 and ds_pk_add_rtn_f16 packed LDS atomic add instructions" + }, + { + "Name": "atomic-flat-pk-add-16-insts", + "Description": "Has flat_atomic_pk_add_f16 and flat_atomic_pk_add_bf16 instructions" + }, + { + "Name": "atomic-buffer-global-pk-add-f16-insts", + "Description": "Has buffer and global packed f16 atomic add instructions that can return the original value" + }, + { + "Name": "atomic-global-pk-add-bf16-inst", + "Description": "Has global_atomic_pk_add_bf16 instruction" + }, + { + "Name": "lds-atomic-add-f64", + "Description": "Has ds_add_f64 and ds_add_rtn_f64 instructions" + }, + { + "Name": "atomic-fmin-fmax-global-f64", + "Description": "Has global/buffer instructions for atomicrmw fmin/fmax on double-precision floats" + }, + { + "Name": "agent-scope-fine-grained-remote-memory-atomics", + "Description": "Agent-scope fine-grained atomics work on remote (non-local) device memory" + }, + { + "Name": "v-mov-b64-inst", + "Description": "Has v_mov_b64 instruction" + }, + { + "Name": "lshl-add-u64-inst", + "Description": "Has v_lshl_add_u64 instruction" + }, + { + "Name": "vgpr-align2", + "Description": "VGPR and AGPR tuple operands require even alignment" + }, + { + "Name": "architected-flat-scratch", + "Description": "Flat Scratch register is a read-only SPI-initialized architected register" + }, + { + "Name": "packed-tid", + "Description": "Work-item IDs are packed into v0 at kernel launch" + }, + { + "Name": "kernarg-preload", + "Description": "Hardware supports preloading of kernel arguments into user SGPRs" + }, + { + "Name": "back-off-barrier", + "Description": "Hardware supports backing off s_barrier if an exception occurs" + }, + { + "Name": "tgsplit-support", + "Description": "Hardware support for threadgroup split execution mode, letting a work-group's wavefronts run on different CUs" + }, + { + "Name": "xnack", + "Description": "Enable XNACK support (memory page-fault replay). Part of the gfx942 target ID as xnack- / xnack+" + }, + { + "Name": "sramecc", + "Description": "Enable SRAM ECC. Part of the gfx942 target ID as sramecc- / sramecc+" + }, + { + "Name": "wavefrontsize64", + "Description": "Wavefronts are 64 work-items wide; gfx942 is wave64 only" } - }, - "memory_ordering": { - "acquire": true, - "release": true, - "acq_rel": true, - "seq_cst": true - }, - "memory_scopes": [ - "workitem", - "subgroup", - "workgroup", - "device", - "system" ], - "register_file": { - "scalar_registers": { - "count": 512, - "width": 32, - "addressing": "s[0:511]" - }, - "vector_registers": { - "count": 512, - "width": 32, - "addressing": "v[0:511]", - "accumulator_registers": { - "count": 512, - "width": 32, - "addressing": "a[0:511]", - "alignment_requirement": "even_pairs" - } + "SubUnits": [ + { + "Type": "HIP", + "Version": "6.0" + }, + { + "Type": "OpenCL", + "Version": "2.0" } + ] + }, + "SpecializedHardware": { + "RayTracingAccelerators": { + "Present": false }, - "code_object": { - "versions_supported": [4, 5, 6], - "default_version": 5, - "elf_format": true, - "metadata_format": "yaml", - "relocatable": true + "AiAccelerators": { + "Present": true, + "Name": "AMD Matrix Core", + "Count": 4, + "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "Other"], + "Performance": { + "fp16TopsPerGPU": 1307.4, + "int8TopsPerGPU": 2614.9 + } }, - "compilation_options": { - "target_cpu": "gfx942", - "target_features": { - "xnack": { - "supported": true, - "default": false, - "description": "Enable XNACK replay for page fault handling" + "videoCodecs": { + "encoders": [], + "decoders": [ + { + "codec": "H.265/HEVC" + }, + { + "codec": "H.264" }, - "sramecc": { - "supported": true, - "default": false, - "description": "Enable SRAM ECC" + { + "codec": "VP9" }, - "cumode": { - "supported": true, - "default": false, - "description": "Enable CU mode" + { + "codec": "AV1" } - } - }, - "operating_systems": { - "amdhsa": { - "supported": true, - "description": "AMD HSA runtime" - }, - "amdpal": { - "supported": true, - "description": "AMD PAL runtime" - }, - "mesa3d": { - "supported": true, - "description": "Mesa 3D runtime" - } - }, - "llvm_integration": { - "target_machine": "AMDGPUTargetMachine", - "instruction_selection": "AMDGPUDAGToDAGISel", - "register_info": "SIRegisterInfo", - "instruction_info": "SIInstrInfo", - "subtarget": "GCNSubtarget", - "calling_conventions": [ - "AMDGPU_CS", - "AMDGPU_KERNEL", - "AMDGPU_PS", - "AMDGPU_VS", - "AMDGPU_GS", - "AMDGPU_HS", - "AMDGPU_DS" ] - }, - "optimization_features": { - "scheduler": "GCNMaxOccupancySchedStrategy", - "register_pressure_tracking": true, - "instruction_bundling": true, - "memory_clause_formation": true, - "wwm_register_allocation": true, - "exec_mask_optimization": true - }, - "debugging_support": { - "dwarf_support": true, - "trap_handler": true, - "printf_support": true, - "debugger_abi": "AMDGPU" - }, - "compatibility": { - "backward_compatible_with": ["gfx940", "gfx941"], - "forward_compatible_with": [], - "related_targets": ["gfx940", "gfx941", "gfx943"] - }, - "notes": { - "description": "GFX942 is part of the CDNA2 architecture family, likely targeting AMD Instinct MI300 series accelerators. It shares most features with GFX940 but may have minor variations or bug fixes.", - "typical_use_cases": [ - "High-performance computing", - "Machine learning training", - "AI inference", - "Scientific computing", - "Data analytics" - ], - "llvm_commit": "9d0572797233857397f3fdc35fffcfb490354f56", - "added_in_llvm_version": "17.0", - "status": "Active development target" } }, - "marketSegments": ["HPC", "AI Training", "AI Inference", "Data Centers"], - "variants": [ - { - "name": "MI300A", - "type": "APU", - "gpuCUs": 228, - "cpuCores": 24, - "memorySize": 128, - "description": "Combined CPU+GPU design with unified memory" - }, - { - "name": "MI300X", - "type": "GPU-only", - "gpuCUs": 228, - "cpuCores": 0, - "memorySize": 192, - "description": "GPU-only variant with expanded HBM3 memory capacity" - } - ], - "software": { - "drivers": ["AMDGPU-PRO", "ROCm"], - "sdks": ["ROCm SDK", "HIP", "OpenMP"], - "frameworks": [ - "TensorFlow", - "PyTorch", - "ONNX Runtime", - "MXNet", - "OpenVINO" + "PowerEfficiency": { + "BoostClock": 2100, + "MaxTDP": 750, + "PowerStates": [ + { + "Name": "Active", + "Description": "Full power state with all 8 XCDs, the Infinity Cache and all 8 HBM3 stacks active, up to the 750 W maximum TBP of the MI300X OAM module. MI300A is rated 550 W for air and liquid cooling and 760 W for liquid cooling only. AMD publishes only the 2100 MHz peak engine clock, not a base clock, so no BaseClock is recorded" + }, + { + "Name": "Idle", + "Description": "Reduced power state when no workloads are running" + } ], - "compilers": ["LLVM/Clang", "HIP Clang", "OpenMP"] + "ClockGating": true, + "DynamicVoltageFrequencyScaling": true }, - "inferencePerformance": { - "resnet50ImagesPerSecond": 200000, - "bert": { - "throughput": 30000, - "unit": "samples/second" - }, - "dlrm": { - "throughput": 50000, - "unit": "samples/second" - } + "PciExpress": { + "Version": "5.0", + "Lanes": 16, + "Bandwidth": 128 + }, + "MultiGpuSupport": { + "Technologies": ["Infinity Fabric"], + "MaxGpus": 8, + "InterconnectBandwidth": 128 } } diff --git a/device_lib/amd-gpu-gfx950.json b/device_lib/amd-gpu-gfx950.json index c9a6e03c..d66f8a2c 100644 --- a/device_lib/amd-gpu-gfx950.json +++ b/device_lib/amd-gpu-gfx950.json @@ -1,238 +1,260 @@ { - "name": "MI350", - "vendor": "AMD", - "generation": 0, - "releaseYear": 0, - "fabricationProcess": { - "processNode": 0.0, - "manufacturer": "ABCDEFGHIJKL", - "technology": "ABCDEFGHIJKLMNOPQRSTUVWXYZ" + "Name": "CDNA 4", + "Vendor": "AMD", + "Architecture": "gfx950", + "ReleaseYear": 2025, + "FabricationProcess": { + "ProcessNode": 3, + "Manufacturer": "TSMC", + "Technology": "N3P for the 8 accelerator complex dies (XCDs), N6 for the 2 I/O dies (IODs)" }, - "computeUnits": { - "name": "ABCDEFGHIJKLMNOPQRS", - "maxPerGPU": 0, - "subUnits": [ - { - "name": "ABCDEFGHIJKLMN", - "countPerComputeUnit": 0, - "description": "ABCDEF" + "CoreSubsystem": { + "Name": "CU", + "ChipType": "Chip", + "CoreType": "Compute Unit", + "UnitTypes": { + "Chip": { + "Count": 1, + "Size": 8, + "Description": "One Instinct MI350-series OAM package: 8 XCDs 3D-stacked in pairs of 4 on top of 2 IODs, plus 8 HBM3E stacks, 185 billion transistors total. 256 active compute units, 16384 stream processors and 1024 Matrix Cores. Primary device modelled here is the liquid-cooled MI355X (up to 2.4 GHz, 1400 W TBP); the air-cooled MI350X is the same silicon at up to 2.2 GHz and 1000 W TBP, giving proportionally ~9% lower peak throughput at identical memory capacity and bandwidth. Eight packages form a fully connected Infinity Fabric platform", + "Memory": ["HBM3E (Device DRAM)", "Infinity Cache (LLC, per IOD)"], + "Subunits": ["XCD", "IOD", "HBM3E Stack", "Infinity Fabric Link", "PCIe Controller"] }, - { - "name": "ABCDEFGHIJKLMNOP", - "countPerComputeUnit": 0, - "description": "ABCDEFGH" + "XCD": { + "Count": 8, + "Size": 32, + "Description": "Accelerator Complex Die built on TSMC N3P. Each XCD holds 36 physical compute units of which 32 are active (4 disabled for yield), a shared command processor and workgroup dispatcher, and 4 MB of L2 cache. Down from 38 active CUs per XCD on CDNA 3 / MI300X, offset by higher clocks", + "Memory": ["L2 (per XCD)"], + "Subunits": ["Compute Unit"] + }, + "Compute Unit": { + "Count": 32, + "Size": 4, + "Description": "CDNA 4 compute unit: 4 SIMD16 vector units, 4 Matrix Cores, 1 scalar unit, 160 KB of LDS and 32 KB of L1 vector cache. Executes wavefronts of 64 work-items, up to 8 waves per SIMD (32 per CU). Delivers 128 FP32 lanes worth of throughput per clock via packed FP32 dual-issue, and half that for FP64", + "Memory": ["LDS (per CU)", "L1 Vector Cache (per CU)", "Vector Register File (per SIMD)"], + "Subunits": ["SIMD", "Matrix Core", "Scalar Unit"] + }, + "SIMD": { + "Count": 4, + "Size": 16, + "Description": "16-lane vector ALU; a wave64 is issued over 4 cycles. Backed by a 128 KB vector register file (512 32-bit VGPRs per lane) that is unified between architected VGPRs and accumulation AGPRs", + "Memory": ["Vector Register File (per SIMD)"] + }, + "Matrix Core": { + "Count": 4, + "Size": 1, + "Description": "Matrix multiply engine executing the MFMA/SMFMAC instruction family. CDNA 4 roughly doubles per-CU matrix throughput versus CDNA 3 for low precision and adds the OCP block-scaled MXFP6 and MXFP4 formats alongside OCP FP8 (E4M3/E5M2), BF16, FP16, FP32, FP64 and INT8. New gfx950 shapes include v_mfma_f32_16x16x128_f8f6f4 and v_mfma_f32_32x32x64_f8f6f4 (with v_mfma_scale_* and v_mfma_ld_scale_b32 for per-block scales), plus wider FP16/BF16/INT8 shapes; SMFMAC provides 2:4 structured sparsity acceleration. Notably FP6 runs at the same rate as FP4 on CDNA 4. The XF32/TF32 MFMA instructions of CDNA 3 (gfx942) are NOT present on gfx950" + }, + "Scalar Unit": { + "Count": 1, + "Size": 1, + "Description": "Per-CU scalar ALU handling wave-uniform arithmetic, control flow and scalar memory (constant) fetches" + }, + "IOD": { + "Count": 2, + "Size": 4, + "Description": "I/O die on TSMC N6. Each IOD hosts 4 of the 8 XCDs stacked above it and 4 of the 8 HBM3E stacks, and carries 128 MB of Infinity Cache plus the Infinity Fabric switching. IOD-to-IOD bisection bandwidth is 5.5 TB/s", + "Memory": ["Infinity Cache (LLC, per IOD)"] + }, + "HBM3E Stack": { + "Count": 8, + "Size": 12, + "Description": "12-high HBM3E stack, 36 GB each, on a 1024-bit interface at 8 Gbps per pin. 8 stacks give 288 GB at 8 TB/s aggregate", + "Memory": ["HBM3E (Device DRAM)"] + }, + "Infinity Fabric Link": { + "Count": 7, + "Size": 1, + "Description": "One of 7 4th-generation Infinity Fabric peer-to-peer links, each 16 bits wide at 38.4 Gbps, for 1075.2 GB/s of aggregate GPU-to-GPU bandwidth in an 8-way fully connected platform" }, + "PCIe Controller": { + "Count": 1, + "Size": 1, + "Description": "PCIe 5.0 x16 host interface controller on the OAM module" + } + } + }, + "MemorySubsystem": { + "SupportedMemoryTypes": ["Other"], + "MemoryTypes": [ { - "name": "ABCDEFGHIJKLMNOPQRSTUVWXYZA", - "countPerComputeUnit": 0, - "description": "ABCDEFGHIJKLMNOPQRSTUVWXY" + "Type": "Vector Register File (per SIMD)", + "Size": 128, + "description": "512 32-bit vector registers per lane across 64 lanes = 128 KB per SIMD, 512 KB per CU. The file is unified: a wave can use up to 512 architected VGPRs, or split the space between architected VGPRs and accumulation AGPRs (up to 256 of each)" }, { - "name": "ABCDE", - "countPerComputeUnit": 0, - "description": "ABCDEFGHIJKLMNOPQRS" + "Type": "LDS (per CU)", + "Size": 160, + "description": "160 KB software-managed Local Data Share per compute unit, up from 64 KB on CDNA 3 — the single largest occupancy-relevant change in CDNA 4. Read bandwidth doubled to 256 bytes per clock, GLOBAL_LOAD_LDS widened to 128 bits per lane, and new read-with-transpose instructions feed the Matrix Cores directly. Total 40 MB across 256 CUs" }, { - "name": "ABCDEFGHIJKLMN", - "countPerComputeUnit": 0, - "description": "ABCDE" + "Type": "L1 Vector Cache (per CU)", + "Size": 32, + "description": "32 KB per-CU vector L1 data cache, unchanged from CDNA 3" }, { - "name": "ABCDEFGHIJKLMNOPQRSTUVW", - "countPerComputeUnit": 0, - "description": "ABCDEFGHIJKLMNOPQRSTUV" + "Type": "L2 (per XCD)", + "Size": 4096, + "description": "4 MB L2 cache per XCD, coherent across all 8 XCDs; 32 MB total per package. CDNA 4 improves L2 writeback handling of dirty data" }, { - "name": "ABCDEFGHIJKLMNOPQRSTUVWXY", - "countPerComputeUnit": 0, - "description": "ABCDEFGHIJKLMN" + "Type": "Infinity Cache (LLC, per IOD)", + "Size": 131072, + "description": "128 MB of memory-attached last-level Infinity Cache per IOD, 256 MB total, sitting in front of HBM3E and shared by all XCDs" }, { - "name": "ABCDEFGHIJKLMNOPQRSTUVWXYZAB", - "countPerComputeUnit": 0, - "description": "ABCDEFGHIJKLMNOPQ" + "Type": "HBM3E (Device DRAM)", + "Size": 301989888, + "BankCount": 8, + "MaxMemoryBandwidth": 8000, + "maxBusWidth": 8192, + "description": "288 GB of HBM3E across 8 12-high stacks at 8 Gbps per pin, 8 TB/s peak bandwidth over an aggregate 8192-bit interface (1024 bits per stack). Identical on MI350X and MI355X" } ] }, - "shaderModel": { - "directX": "ABCDEFGHIJKLMNOPQ", - "vulkan": "ABCDEFGHIJKLMNOPQ", - "openGL": "ABCDEFGHIJKLMNOPQRSTUV", - "openCL": "ABCD", - "metal": "ABCDEFGHIJKLMNOPQRSTUVWXYZ" - }, - "memorySubsystem": { - "supportedMemoryTypes": [ - "Other", - "Other", - "HBM2", - "GDDR6" + "KernelModel": { + "LLVMTarget": "amdgcn", + "LLVMTriple": "amdgcn-amd-amdhsa", + "LLVMFeatures": [ + { + "Name": "gfx950-insts", + "Description": "Additional instructions for GFX950+; umbrella feature that implies the permlane swap, arithmetic-shift-pack, block-scaled FP8/BF8/FP6/BF6/FP4 conversion and minimum3/maximum3 features below" + }, + { + "Name": "mai-insts", + "Description": "Has mAI instructions - the MFMA/SMFMAC matrix arithmetic family executed on the Matrix Cores" + }, + { + "Name": "fp8-cvt-scale-insts", + "Description": "Has fp8 conversion scale instructions, used for block-scaled OCP MXFP8 (E4M3/E5M2) data" + }, + { + "Name": "fp6bf6-cvt-scale-insts", + "Description": "Has fp6 and bf6 conversion scale instructions, new in CDNA 4 for the MXFP6 format" + }, + { + "Name": "fp4-cvt-scale-insts", + "Description": "Has fp4 conversion scale instructions, new in CDNA 4 for the MXFP4 format" + }, + { + "Name": "permlane16-swap", + "Description": "Has v_permlane16_swap_b32 instructions" + }, + { + "Name": "permlane32-swap", + "Description": "Has v_permlane32_swap_b32 instructions" + }, + { + "Name": "ashr-pk-insts", + "Description": "Has Arithmetic Shift Pack instructions" + }, + { + "Name": "fp8-insts", + "Description": "Has fp8 and bf8 instructions" + }, + { + "Name": "ocp-fp8-conversion-insts", + "Description": "Has OCP fp8 (E4M3FN, E5M2) conversion instructions" + }, + { + "Name": "bf16-cvt-insts", + "Description": "Has bf16 conversion instructions" + }, + { + "Name": "bitop3-insts", + "Description": "Has v_bitop3_b32/v_bitop3_b16 instructions" + }, + { + "Name": "prng-inst", + "Description": "Has v_prng_b32 pseudo-random number generation instruction" + }, + { + "Name": "dot12-insts", + "Description": "Has v_dot2_f32_bf16 instructions" + }, + { + "Name": "dot13-insts", + "Description": "Has v_dot2c_f32_bf16 instructions" + }, + { + "Name": "packed-fp32-ops", + "Description": "Has packed FP32 dual-issue instructions (v_pk_fma_f32 and friends), giving 128 FP32 lanes of throughput per CU" + }, + { + "Name": "wavefrontsize64", + "Description": "Fixed 64 work-items per wavefront, implied by the GFX9 feature set. CDNA targets have no wave32 mode and cumode/WGP selection does not apply" + }, + { + "Name": "sramecc", + "Description": "SRAM ECC protection over register files, LDS and caches. Part of the target ID (e.g. gfx950:sramecc+:xnack-) and affects VGPR allocation granularity" + }, + { + "Name": "xnack", + "Description": "Page-fault recovery support enabling unified/shared virtual memory between host and device. Part of the target ID" + }, + { + "Name": "tgsplit", + "Description": "Enable threadgroup split execution, allowing a workgroup's waves to be distributed across compute units" + }, + { + "Name": "kernarg-preload", + "Description": "Hardware supports preloading of kernel arguments into user SGPRs" + } ], - "maxMemoryBandwidth": 0.0, - "maxMemorySize": 0, - "maxBusWidth": 0, - "cacheHierarchy": [ + "SubUnits": [ { - "level": "ABCDEFGHI", - "sizePerUnit": 0, - "totalSize": 0, - "description": "ABCDEFGH" + "Type": "ROCm", + "Version": "7.0" }, { - "level": "ABCDEFGHIJKLMNOPQRS", - "sizePerUnit": 0, - "totalSize": 0, - "description": "ABCDEFGHIJKLMNOPQRS" + "Type": "HIP", + "Version": "7.0" }, { - "level": "ABCDE", - "sizePerUnit": 0, - "totalSize": 0, - "description": "ABCDEFGHIJKLMNOPQRSTUVWXYZAB" + "Type": "OpenCL", + "Version": "2.0" } ] }, - "specializedHardware": { - "rayTracingAccelerators": { - "present": false, - "name": "ABCD", - "countPerComputeUnit": 0, - "performance": { - "raysPerSecond": 0.0 - } + "SpecializedHardware": { + "RayTracingAccelerators": { + "Present": false }, - "aiAccelerators": { - "present": true, - "name": "ABCDEFGHIJKLMNO", - "countPerComputeUnit": 0, - "supportedPrecisions": [ - "INT4" - ], - "performance": { - "fp16TopsPerGPU": 0.0, - "int8TopsPerGPU": 0.0 + "AiAccelerators": { + "Present": true, + "Name": "Matrix Core", + "Count": 1024, + "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "Other"], + "Performance": { + "fp16TopsPerGPU": 2500, + "int8TopsPerGPU": 5000 } - }, - "videoCodecs": { - "encoders": [ - { - "codec": "AV1", - "maxResolution": "ABCDEFGHIJKLMN", - "maxBitrate": 0, - "maxFPS": 0 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "ABCDEFGHIJKL", - "maxBitrate": 0, - "maxFPS": 0 - }, - { - "codec": "VP9", - "maxResolution": "ABCDEFGHIJKLMNOPQRSTUVWX", - "maxBitrate": 0, - "maxFPS": 0 - } - ], - "decoders": [ - { - "codec": "H.265/HEVC", - "maxResolution": "ABCDEFGHIJKLMNOPQRSTUVWXYZ", - "maxBitrate": 0, - "maxFPS": 0 - }, - { - "codec": "H.264", - "maxResolution": "ABCDEFGHIJKLMNOP", - "maxBitrate": 0, - "maxFPS": 0 - }, - { - "codec": "AV1", - "maxResolution": "ABCDEFGHIJKLMNOPQRSTUVWX", - "maxBitrate": 0, - "maxFPS": 0 - }, - { - "codec": "Other", - "maxResolution": "ABCDEFGHI", - "maxBitrate": 0, - "maxFPS": 0 - }, - { - "codec": "AV1", - "maxResolution": "ABCDEFGHIJKLMNOPQRSTUVWXYZAB", - "maxBitrate": 0, - "maxFPS": 0 - }, - { - "codec": "AV1", - "maxResolution": "ABCDEFGHIJKLMN", - "maxBitrate": 0, - "maxFPS": 0 - } - ] } }, - "powerEfficiency": { - "maxTDP": 0, - "powerStates": [ - { - "name": "ABCDEFGHIJKLMNOPQRSTUV", - "description": "ABCDEFGHI" - }, - { - "name": "ABCDEFGHIJKLMNOPQRST", - "description": "ABCDEFGHIJKLMNOPQRSTUVWXYZABC" - }, + "PowerEfficiency": { + "BoostClock": 2400, + "MaxTDP": 1400, + "PowerStates": [ { - "name": "ABCDEFGHIJKL", - "description": "ABCDEFGHIJKLMNOPQRSTUVWXYZAB" + "Name": "Active", + "Description": "All 256 CUs and 8 HBM3E stacks active. MI355X is rated at 1400 W TBP with up to a 2.4 GHz peak engine clock and requires liquid cooling; the air-cooled MI350X is rated at 1000 W TBP with up to a 2.2 GHz peak engine clock. AMD publishes peak engine clock only; no base clock is published for Instinct parts" }, { - "name": "ABCDEFGHIJKLMNOPQR", - "description": "ABCDEFGHIJKLMNOPQRSTUVWXYZABC" + "Name": "Idle", + "Description": "Reduced power state when no workloads are resident" } ], - "clockGating": false, - "dynamicVoltageFrequencyScaling": false - }, - "displayOutputs": { - "maxDisplays": 0, - "maxResolution": "ABCDEFGHIJKLMN", - "maxRefreshRate": 0, - "interfaces": [ - "Other", - "DVI", - "USB-C/DP", - "DVI", - "USB-C/DP" - ], - "hdr": true - }, - "pciExpress": { - "version": "ABCDEFGHIJKLMNOPQRSTUVWXYZ", - "lanes": 0, - "bandwidth": 0.0 + "ClockGating": true, + "DynamicVoltageFrequencyScaling": true }, - "multiGpuSupport": { - "technologies": [ - "Other", - "Other", - "NVLink", - "Infinity Fabric" - ], - "maxGpus": 0, - "interconnectBandwidth": 0.0 + "PciExpress": { + "Version": "5.0", + "Lanes": 16, + "Bandwidth": 128 }, - "softwareFeatures": { - "upscalingTechnologies": [ - "XeSS", - "XeSS", - "FSR" - ], - "meshShading": true, - "variableRateShading": true, - "samplerFeedback": false + "MultiGpuSupport": { + "Technologies": ["Infinity Fabric"], + "MaxGpus": 8, + "InterconnectBandwidth": 1075.2 } } diff --git a/device_lib/apple-gpu-.json b/device_lib/apple-gpu-.json index 97781f2d..cc4e7c32 100644 --- a/device_lib/apple-gpu-.json +++ b/device_lib/apple-gpu-.json @@ -2,44 +2,45 @@ { "Name": "Apple GPU (A7)", "Vendor": "Apple", - "Architecture": "PowerVR G6430", + "Architecture": "Apple1", "ReleaseYear": 2013, "FabricationProcess": { "ProcessNode": 28, - "Manufacturer": "TSMC", + "Manufacturer": "Samsung", "Technology": "28nm HKMG" }, "CoreSubsystem": { - "Name": "PowerVR Rogue", - "subUnits": [ - { - "Type": "USC (Unified Shading Cluster)", + "Name": "USC", + "ChipType": "GPU", + "CoreType": "USC", + "UnitTypes": { + "GPU": { + "Count": 1, + "Size": 4, + "Description": "Integrated GPU of the Apple A7 (S5L8960X) SoC, the first 64-bit smartphone SoC, shipped in iPhone 5s, iPad Air and iPad mini 2/3. Apple announced only a die size of 102 mm2 and 'over 1 billion' transistors and has never named the GPU; die-shot analysis (AnandTech, endorsed by Chipworks) identifies it as the four-cluster Imagination PowerVR Series6 'Rogue' G6430, which is analyst consensus rather than a vendor statement. Apple's retroactive Metal label is MTLGPUFamily.apple1, documented as 'the Apple family 1 GPU features that correspond to the Apple A7 GPUs'; no A7 device ever reported it at run time, since A7 devices stopped at iOS 12 while MTLGPUFamily arrived in iOS 13. A7 has since been dropped from Apple's Metal Feature Set Tables entirely - the oldest entry there is now the A8-series (Apple2). Apple publishes no GPU clock, no memory bandwidth figure, no GPU cache size and no TDP for the A7, so all of those are omitted here rather than estimated", + "Memory": ["System Memory (LPDDR3)"], + "Subunits": ["USC"] + }, + "USC": { "Count": 4, - "Size": 32, - "Description": "Unified shader processors for vertex, pixel, and compute operations", - "Memory": ["L1"], - "SubunitType": "SIMD" + "Description": "Unified Shading Cluster, the PowerVR Rogue shader core handling vertex, pixel and compute work. Imagination's published Series6 table gives the G6430 four USCs and 128 FP32 ALUs (192 FP16), i.e. 32 FP32 ALUs per cluster. AnandTech's analysis describes each USC as a 16-wide scalar SIMD whose pipelines each issue two FP32 MADs per clock, giving 64 FP32 operations per clock per cluster and 256 per clock for the whole GPU. Apple documents that the A7, A8 and A9 GPUs are fully scalar - all floating-point work runs on a scalar processor even for values declared as vectors - and that both mediump and lowp are computed at FP16, a change from the fixed-point lowp of the earlier SGX parts. No absolute FLOPS rating can be given because no GPU clock frequency has ever been published by Apple or Imagination", + "Subunits": ["SIMD Pipeline"] + }, + "SIMD Pipeline": { + "Count": 1, + "Size": 16, + "Description": "16-wide scalar SIMD pipeline, one per USC, with each lane sustaining two FP32 MADs per clock. The 16-lane width is third-party analysis of the PowerVR Rogue USC, not a figure published by Apple or Imagination. The 32-thread SIMD-group width of Apple's own later GPU designs does not apply to this PowerVR-derived architecture" } - ] + } }, "MemorySubsystem": { - "SupportedMemoryTypes": ["LPDDR3"], + "SupportedMemoryTypes": ["Other"], "MemoryTypes": [ { - "Type": "L1", - "Size": 32, - "BankCount": 4, - "MaxMemoryBandwidth": 8.5, - "maxBusWidth": 64, - "description": "L1 cache for shader operations" - }, - { - "Type": "System Memory", + "Type": "System Memory (LPDDR3)", "Size": 1048576, - "BankCount": 2, - "MaxMemoryBandwidth": 8.5, "maxBusWidth": 64, - "description": "Shared system memory with CPU" + "description": "1 GB of LPDDR3 package-on-package memory shared by CPU and GPU over a 64-bit (2 x 32-bit) interface, per Chipworks' teardown; AnandTech notes the A7 stepped back from the 128-bit bus of the A5X/A6X and that all three A7 devices use identical silicon with no increase in memory bandwidth. Apple publishes neither the memory type nor a bandwidth figure, and no primary source establishes the LPDDR3 data rate, so no MaxMemoryBandwidth is recorded (the widely quoted 12.8 GB/s is an unverified derivation). LPDDR3 is reported as 'Other' in SupportedMemoryTypes because the schema enum has no LPDDR3 value. Neither Apple nor Imagination publishes any GPU L1, L2 or texture cache size for the A7; the roughly 4 MB on-die SRAM that die-shot analysis found between the GPU and the DRAM interface was never publicly established to be a GPU cache, so it is deliberately not listed as one" } ] }, @@ -57,29 +58,16 @@ }, "SpecializedHardware": { "RayTracingAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "Performance": { - "RaysPerSecond": 0 - } + "Present": false }, "AiAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "SupportedPrecisions": [], - "Performance": { - "fp16TopsPerGPU": 0, - "int8TopsPerGPU": 0 - } + "Present": false }, "videoCodecs": { "encoders": [ { "codec": "H.264", "maxResolution": "1080p", - "maxBitrate": 25, "maxFPS": 30 } ], @@ -87,44 +75,40 @@ { "codec": "H.264", "maxResolution": "1080p", - "maxBitrate": 25, + "maxFPS": 60 + }, + { + "codec": "Other", + "maxResolution": "720p", + "maxBitrate": 35, "maxFPS": 30 } ] } }, "PowerEfficiency": { - "BaseClock": 450, - "BoostClock": 450, - "MaxTDP": 2, "PowerStates": [ { "Name": "Active", - "Description": "GPU actively processing graphics workloads" + "Description": "GPU actively processing graphics or compute work. Apple has never published a GPU clock frequency or a TDP for any A-series SoC, so BaseClock, BoostClock and MaxTDP are omitted rather than estimated; the frequently repeated 450 MHz figure for the A7 GPU has no primary source, and even contemporary analysts stated they did not have the number" }, { "Name": "Idle", - "Description": "Low power state when GPU is not in use" + "Description": "Low-power state when no GPU work is queued" } ], "ClockGating": true, "DynamicVoltageFrequencyScaling": true }, - "PciExpress": { - "Version": "N/A", - "Lanes": 0, - "Bandwidth": 0 - }, "MultiGpuSupport": { "Technologies": ["None"], - "MaxGpus": 1, - "InterconnectBandwidth": 0 + "MaxGpus": 1 } }, { "Name": "Apple GPU (M1)", "Vendor": "Apple", - "Architecture": "Apple GPU Gen 1", + "Architecture": "Apple7", "ReleaseYear": 2020, "FabricationProcess": { "ProcessNode": 5, @@ -133,490 +117,478 @@ }, "CoreSubsystem": { "Name": "GPU Core", - "subUnits": [ - { - "Type": "Execution Unit", + "ChipType": "GPU", + "CoreType": "GPU Core", + "UnitTypes": { + "GPU": { + "Count": 1, + "Size": 8, + "Description": "Integrated GPU block of the Apple M1 SoC, which Apple states is built on 5-nanometer process technology with 16 billion transistors (Apple names neither the foundry nor the node name; TSMC N5 is the universally reported third-party attribution). Apple ships 8-core and binned 7-core GPU configurations; the 8-core part is described here. Apple's own performance figure is 2.6 teraflops of throughput with nearly 25,000 threads in flight. Metal reports GPU family Apple7 - Apple's Metal Feature Set Tables list the whole M1-series as Apple7, the same family as A14 Bionic, rated for 'Metal 3 & 4' (that rating reflects current OS support; M1 shipped in 2020 against Metal 2). There is no dedicated VRAM: the GPU shares LPDDR4X unified memory with the CPU, the 16-core Neural Engine (Apple-published at 11 trillion operations per second, with the numeric precision unstated) and the media engine, each a separate block on the same die. The M1 GPU has no ray-tracing and no matrix/tensor hardware; Metal nonetheless exposes ray tracing in compute and render pipelines from Apple6 and mesh shading and SIMD-scoped matrix multiply from Apple7, executed on the ordinary shader ALUs. Apple provides no native Vulkan implementation - Vulkan is available only through the MoltenVK translation layer - and OpenGL and OpenCL have been deprecated on macOS since 10.14, so the OpenGL 4.1 and OpenCL 1.2 levels listed under KernelModel are the long-standing macOS ceilings rather than figures Apple has restated for Apple silicon", + "Memory": ["L2 Cache (GPU)", "System Level Cache", "Unified Memory (LPDDR4X)"], + "Subunits": ["GPU Core"] + }, + "GPU Core": { "Count": 8, - "Size": 128, - "Description": "Unified shader processors with 128 ALUs per core", - "Memory": ["L1", "L2"], - "SubunitType": "SIMD" + "Size": 4, + "Description": "Apple GPU core with 128 FP32 ALUs organised as 4 SIMD pipelines of 32 lanes, giving 1024 ALUs across the 8-core GPU. Apple does not publish per-core ALU counts; 128 ALUs per core is the standard third-party figure (AnandTech explicitly labels its Apple GPU table an educated guess) and it is corroborated arithmetically by Apple's own rating, since 8 cores x 128 ALUs x 2 FLOP x 1.278 GHz = 2.62 TFLOPS against Apple's 2.6 TFLOPS. On the Apple7 family FP16 executes at the same peak rate as FP32 rather than at double rate as on A14, per public microbenchmarking", + "Memory": ["Register File", "Threadgroup Memory", "L1 Data Cache", "L1 Instruction Cache"], + "Subunits": ["SIMD Pipeline"] }, - { - "Type": "Tile Memory", - "Count": 8, + "SIMD Pipeline": { + "Count": 4, "Size": 32, - "Description": "On-chip tile memory for efficient rendering", - "Memory": ["Tile"], - "SubunitType": "Memory" + "Description": "32-lane SIMD execution pipeline issuing one instruction per cycle for one 32-thread SIMD-group. The SIMD-group width is a run-time Metal property (threadExecutionWidth) rather than a tabulated feature-set limit; 32 threads per SIMD-group is documented for the M1's G13 architecture by public ISA reverse-engineering, which also records up to 128 32-bit general-purpose registers per thread. The decomposition of the core's 128 ALUs into 4 pipelines of 32 lanes is not published by Apple and comes from the same third-party analysis" } - ] + } }, "MemorySubsystem": { - "SupportedMemoryTypes": ["LPDDR4X"], + "SupportedMemoryTypes": ["Other"], "MemoryTypes": [ { - "Type": "L1", - "Size": 64, - "BankCount": 8, - "MaxMemoryBandwidth": 68.25, - "maxBusWidth": 128, - "description": "L1 cache per GPU core" + "Type": "Register File", + "Size": 208, + "description": "Approximately 208 KB of registers per GPU core across the Apple7 and Apple8 generations. Not published by Apple; measured by a single public microbenchmarking project. Apple's register file is large relative to its very small L1 caches, which is what allows deep occupancy on such a small cache hierarchy" }, { - "Type": "L2", - "Size": 4096, - "BankCount": 1, - "MaxMemoryBandwidth": 68.25, - "maxBusWidth": 128, - "description": "Shared L2 cache" + "Type": "Threadgroup Memory", + "Size": 32, + "description": "Apple's Metal Feature Set Tables cap a single threadgroup at 32 KB of threadgroup memory on the Apple7 family, with a maximum of 1024 threads per threadgroup, 16 B threadgroup memory length alignment, 32 KB maximum explicit imageblock allocation and 128 KB implicit. The physical per-core pool is larger than the per-threadgroup cap - public microbenchmarking measures roughly 60 KB per core - which allows several threadgroups to be resident simultaneously" + }, + { + "Type": "L1 Data Cache", + "Size": 8, + "description": "8 KB L1 data cache per GPU core. Apple publishes no GPU cache sizes at all; 8 KB is measured independently by two public microbenchmarking efforts, one of which notes it is half the size of AMD Vega's L1 at similar latency. Deliberately small compared with contemporary AMD and NVIDIA cores, offset by the large register file and a low-latency L2" + }, + { + "Type": "L1 Instruction Cache", + "Size": 12, + "description": "12 KB L1 instruction cache per GPU core. Not published by Apple; single public microbenchmarking source" }, { - "Type": "Unified Memory", + "Type": "L2 Cache (GPU)", + "Size": 768, + "description": "GPU-wide L2 cache. Apple publishes no figure and the two public measurements disagree: one microbenchmarking project reports 768 KB for the M1 while another measured 'a large 1 MB L2'. The lower of the two is recorded here; treat the real capacity as somewhere in the 0.75-1 MB range. Apple L2 capacity is known to vary non-monotonically across a family, so this value must not be scaled to the Pro and Max dies" + }, + { + "Type": "System Level Cache", + "Size": 8192, + "description": "8 MB SoC-wide system level cache shared by CPU, GPU, Neural Engine and media engine, backing the GPU L2 in front of LPDDR4X. Not published by Apple; measured third-party, with two independent sources agreeing on 8 MB. An earlier third-party estimate of 16 MB made at M1 launch was superseded by this measurement" + }, + { + "Type": "Unified Memory (LPDDR4X)", "Size": 16777216, - "BankCount": 4, "MaxMemoryBandwidth": 68.25, "maxBusWidth": 128, - "description": "Unified memory shared with CPU" + "description": "LPDDR4X unified memory shared with the CPU, Neural Engine and media engine; there is no separate VRAM and no copy across a bus. Apple publishes 8 GB and 16 GB configurations - the 16 GB maximum is recorded here in KB - but publishes no bandwidth figure for the base M1, since the memory-bandwidth line first appears in Apple's specifications with M1 Pro. The 68.25 GB/s over a 128-bit bus (8 x 16-bit channels of LPDDR4X-4266-class memory) is a third-party figure. LPDDR4X is reported as 'Other' in SupportedMemoryTypes because the schema enum has no LPDDR4X value" } ] }, "KernelModel": { + "LLVMTarget": "air64", + "LLVMTriple": "air64-apple-macos", + "LLVMFeatures": [ + { + "Name": "air64", + "Description": "Apple Intermediate Representation (AIR), the LLVM-bitcode-based IR that Apple's metal compiler emits into a .metallib (-target air64-apple-macos, or air64-apple-ios for iOS). AIR is Apple-private, is not an upstream LLVM backend target, and Apple publishes no GPU ISA documentation" + } + ], "SubUnits": [ { "Type": "Metal", - "Version": "2.3" + "Version": "3" }, { - "Type": "OpenGL ES", - "Version": "3.0" + "Type": "Metal", + "Version": "4" + }, + { + "Type": "OpenGL", + "Version": "4.1" }, { - "Type": "Vulkan", + "Type": "OpenCL", "Version": "1.2" } ] }, "SpecializedHardware": { "RayTracingAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "Performance": { - "RaysPerSecond": 0 - } + "Present": false }, "AiAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "SupportedPrecisions": [], - "Performance": { - "fp16TopsPerGPU": 0, - "int8TopsPerGPU": 0 - } + "Present": true, + "Name": "16-core Neural Engine", + "Count": 16 }, "videoCodecs": { "encoders": [ { - "codec": "H.264", - "maxResolution": "4K", - "maxBitrate": 160, - "maxFPS": 60 + "codec": "H.264" }, { - "codec": "H.265/HEVC", - "maxResolution": "4K", - "maxBitrate": 160, - "maxFPS": 60 + "codec": "H.265/HEVC" } ], "decoders": [ { "codec": "H.264", - "maxResolution": "4K", - "maxBitrate": 160, - "maxFPS": 60 + "maxResolution": "4K" }, { "codec": "H.265/HEVC", - "maxResolution": "4K", - "maxBitrate": 160, - "maxFPS": 60 - }, - { - "codec": "VP9", - "maxResolution": "4K", - "maxBitrate": 160, - "maxFPS": 60 + "maxResolution": "4K" } ] } }, "PowerEfficiency": { - "BaseClock": 1000, "BoostClock": 1278, - "MaxTDP": 15, "PowerStates": [ { - "Name": "Performance", - "Description": "Maximum performance for demanding workloads" - }, - { - "Name": "Balanced", - "Description": "Balance between performance and power efficiency" + "Name": "Active", + "Description": "GPU active and clocking up to its maximum. Apple publishes neither GPU clocks nor a TDP for the M1, so no MaxTDP and no BaseClock are recorded here. The 1278 MHz maximum is a third-party figure, and it is the frequency that reproduces Apple's own 2.6 TFLOPS rating from an 8-core, 128-ALU-per-core GPU" }, { - "Name": "Power Saving", - "Description": "Optimized for battery life" + "Name": "Idle", + "Description": "Reduced power state when no GPU work is queued" } ], "ClockGating": true, "DynamicVoltageFrequencyScaling": true }, - "PciExpress": { - "Version": "N/A", - "Lanes": 0, - "Bandwidth": 0 - }, "MultiGpuSupport": { "Technologies": ["None"], - "MaxGpus": 1, - "InterconnectBandwidth": 0 + "MaxGpus": 1 } }, { "Name": "Apple GPU (M1 Pro)", "Vendor": "Apple", - "Architecture": "Apple GPU Gen 1 Pro", + "Architecture": "Apple7", "ReleaseYear": 2021, "FabricationProcess": { "ProcessNode": 5, "Manufacturer": "TSMC", - "Technology": "N5P" + "Technology": "N5" }, "CoreSubsystem": { "Name": "GPU Core", - "subUnits": [ - { - "Type": "Execution Unit", + "ChipType": "GPU", + "CoreType": "GPU Core", + "UnitTypes": { + "GPU": { + "Count": 1, + "Size": 16, + "Description": "Integrated GPU block of the Apple M1 Pro SoC, which Apple states packs 33.7 billion transistors using 5-nanometer process technology (Apple names neither foundry nor node name; the N5/N5P distinction reported by third parties for this die is not established by any source). Apple ships 16-core and binned 14-core GPU configurations; the 16-core part is described here. Metal reports GPU family Apple7, since Apple's Metal Feature Set Tables classify the whole M1-series - base, Pro, Max and Ultra - as Apple7. Apple published no teraflops rating for M1 Pro: its October 2021 announcement contains no FLOPS claim at all, and the widely circulated 5.2 TFLOPS number is not Apple's. Public microbenchmarking measures roughly 5.3 TFLOPS FP32 for the 16-core part, consistent with 2048 ALUs at about 1.296 GHz. The GPU shares LPDDR5 unified memory with the CPU, the 16-core Neural Engine and a media engine that, unlike the base M1's, adds hardware ProRes and ProRes RAW acceleration. There is no ray-tracing and no matrix/tensor hardware; Metal's ray-tracing and mesh-shading APIs run on the ordinary shader ALUs. Apple provides no native Vulkan implementation (MoltenVK translation only), and OpenGL and OpenCL are deprecated on macOS", + "Memory": ["L2 Cache (GPU)", "System Level Cache", "Unified Memory (LPDDR5)"], + "Subunits": ["GPU Core"] + }, + "GPU Core": { "Count": 16, - "Size": 128, - "Description": "Unified shader processors with 128 ALUs per core", - "Memory": ["L1", "L2"], - "SubunitType": "SIMD" + "Size": 4, + "Description": "Apple GPU core with 128 FP32 ALUs organised as 4 SIMD pipelines of 32 lanes, giving 2048 ALUs across the 16-core GPU. Apple does not publish per-core ALU counts; the 128-ALU figure is third-party and is the same core design as the base M1, both being Apple7. On the Apple7 family FP16 executes at the same peak rate as FP32", + "Memory": ["Register File", "Threadgroup Memory", "L1 Data Cache", "L1 Instruction Cache"], + "Subunits": ["SIMD Pipeline"] }, - { - "Type": "Tile Memory", - "Count": 16, + "SIMD Pipeline": { + "Count": 4, "Size": 32, - "Description": "On-chip tile memory for efficient rendering", - "Memory": ["Tile"], - "SubunitType": "Memory" + "Description": "32-lane SIMD execution pipeline issuing one instruction per cycle for one 32-thread SIMD-group. The 32-thread SIMD-group width is a run-time Metal property (threadExecutionWidth) documented for this G13 architecture by public ISA reverse-engineering; the split of 128 ALUs into 4 pipelines of 32 lanes is third-party analysis rather than an Apple disclosure" } - ] + } }, "MemorySubsystem": { "SupportedMemoryTypes": ["LPDDR5"], "MemoryTypes": [ { - "Type": "L1", - "Size": 64, - "BankCount": 16, - "MaxMemoryBandwidth": 200, - "maxBusWidth": 256, - "description": "L1 cache per GPU core" + "Type": "Register File", + "Size": 208, + "description": "Approximately 208 KB of registers per GPU core across the Apple7 and Apple8 generations. Not published by Apple; measured by a single public microbenchmarking project" }, { - "Type": "L2", - "Size": 8192, - "BankCount": 1, - "MaxMemoryBandwidth": 200, - "maxBusWidth": 256, - "description": "Shared L2 cache" + "Type": "Threadgroup Memory", + "Size": 32, + "description": "Apple's Metal Feature Set Tables cap a single threadgroup at 32 KB of threadgroup memory on the Apple7 family, with 1024 maximum threads per threadgroup, 16 B length alignment, 32 KB maximum explicit imageblock allocation and 128 KB implicit. Public microbenchmarking measures a physical per-core pool of roughly 60 KB, larger than the per-threadgroup cap" + }, + { + "Type": "L1 Data Cache", + "Size": 8, + "description": "8 KB L1 data cache per GPU core, unchanged across the Apple7 family. Apple publishes no GPU cache sizes; this figure is measured by two independent public microbenchmarking efforts" + }, + { + "Type": "L1 Instruction Cache", + "Size": 12, + "description": "12 KB L1 instruction cache per GPU core. Not published by Apple; single public microbenchmarking source" }, { - "Type": "Unified Memory", + "Type": "L2 Cache (GPU)", + "Size": 256, + "description": "GPU-wide L2 cache, measured at 256 KB on the M1 Pro by a single public microbenchmarking project. Apple publishes nothing here, and the figure is counter-intuitive: the same source measures 768 KB on the smaller base M1 and 512 KB on the M1 Max, so Apple's L2 capacity varies non-monotonically within the family and must not be inferred by scaling. Treat as indicative, single-source only" + }, + { + "Type": "System Level Cache", + "Size": 24576, + "description": "24 MB SoC-wide system level cache shared by CPU, GPU, Neural Engine and media engine, sitting behind the GPU L2 and in front of LPDDR5. Not published by Apple; measured third-party, with two independent sources agreeing on 24 MB for M1 Pro against 8 MB for M1 and 48 MB for M1 Max" + }, + { + "Type": "Unified Memory (LPDDR5)", "Size": 33554432, - "BankCount": 8, "MaxMemoryBandwidth": 200, "maxBusWidth": 256, - "description": "Unified memory shared with CPU" + "description": "LPDDR5 unified memory shared with the CPU, Neural Engine and media engine, with no separate VRAM and no copy across a bus. Apple publishes 'up to 200GB/s of memory bandwidth with support for up to 32GB' of unified memory; the 32 GB maximum is recorded here in KB. Apple does not state the bus width - the 256-bit LPDDR5-6400 interface is third-party analysis, and its 204.8 GB/s theoretical peak is what Apple rounds down to 200 GB/s" } ] }, "KernelModel": { + "LLVMTarget": "air64", + "LLVMTriple": "air64-apple-macos", + "LLVMFeatures": [ + { + "Name": "air64", + "Description": "Apple Intermediate Representation (AIR), the LLVM-bitcode-based IR that Apple's metal compiler emits into a .metallib (-target air64-apple-macos). AIR is Apple-private and is not an upstream LLVM backend target; Apple publishes no GPU ISA documentation" + } + ], "SubUnits": [ { "Type": "Metal", - "Version": "2.4" + "Version": "3" }, { - "Type": "OpenGL ES", - "Version": "3.0" + "Type": "Metal", + "Version": "4" + }, + { + "Type": "OpenGL", + "Version": "4.1" }, { - "Type": "Vulkan", + "Type": "OpenCL", "Version": "1.2" } ] }, "SpecializedHardware": { "RayTracingAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "Performance": { - "RaysPerSecond": 0 - } + "Present": false }, "AiAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "SupportedPrecisions": [], - "Performance": { - "fp16TopsPerGPU": 0, - "int8TopsPerGPU": 0 - } + "Present": true, + "Name": "16-core Neural Engine", + "Count": 16 }, "videoCodecs": { "encoders": [ { - "codec": "H.264", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 + "codec": "H.264" }, { - "codec": "H.265/HEVC", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 + "codec": "H.265/HEVC" + }, + { + "codec": "Other", + "maxResolution": "8K" } ], "decoders": [ { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 400, - "maxFPS": 30 + "codec": "H.264" }, { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 400, - "maxFPS": 30 + "codec": "H.265/HEVC" }, { - "codec": "VP9", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 + "codec": "Other", + "maxResolution": "8K" } ] } }, "PowerEfficiency": { - "BaseClock": 1000, "BoostClock": 1296, - "MaxTDP": 30, "PowerStates": [ { - "Name": "Performance", - "Description": "Maximum performance for demanding workloads" - }, - { - "Name": "Balanced", - "Description": "Balance between performance and power efficiency" + "Name": "Active", + "Description": "GPU active and clocking up to its maximum. Apple publishes neither GPU clocks nor a TDP for M1 Pro, so no MaxTDP and no BaseClock are recorded. The 1296 MHz maximum is a third-party measurement, corroborated by a second independent source, and applies to both M1 Pro and M1 Max" }, { - "Name": "Power Saving", - "Description": "Optimized for battery life" + "Name": "Idle", + "Description": "Reduced power state when no GPU work is queued" } ], "ClockGating": true, "DynamicVoltageFrequencyScaling": true }, - "PciExpress": { - "Version": "N/A", - "Lanes": 0, - "Bandwidth": 0 - }, "MultiGpuSupport": { "Technologies": ["None"], - "MaxGpus": 1, - "InterconnectBandwidth": 0 + "MaxGpus": 1 } }, { "Name": "Apple GPU (M1 Max)", "Vendor": "Apple", - "Architecture": "Apple GPU Gen 1 Max", + "Architecture": "Apple7", "ReleaseYear": 2021, "FabricationProcess": { "ProcessNode": 5, "Manufacturer": "TSMC", - "Technology": "N5P" + "Technology": "N5" }, "CoreSubsystem": { "Name": "GPU Core", - "subUnits": [ - { - "Type": "Execution Unit", - "Count": 32, - "Size": 128, - "Description": "Unified shader processors with 128 ALUs per core", - "Memory": ["L1", "L2"], - "SubunitType": "SIMD" + "ChipType": "GPU", + "CoreType": "GPU Core", + "UnitTypes": { + "GPU": { + "Count": 1, + "Size": 32, + "Description": "Integrated GPU block of the Apple M1 Max SoC, which Apple states has 57 billion transistors on 5-nanometer process technology (Apple names neither foundry nor node name). Apple ships 32-core and binned 24-core GPU configurations; the 32-core part is described here. Metal reports GPU family Apple7, as Apple's Metal Feature Set Tables classify the entire M1-series as Apple7. Apple published no teraflops rating for M1 Max - the October 2021 announcement contains no FLOPS claim, and the widely circulated 10.4 TFLOPS figure is not Apple's; public microbenchmarking measures roughly 10.6 TFLOPS FP32, consistent with 4096 ALUs at about 1.296 GHz. The GPU shares LPDDR5 unified memory with the CPU, the 16-core Neural Engine and the media engine, which on M1 Max carries two ProRes accelerators and which Apple rates at up to 2x the video encode performance of M1 Pro. There is no ray-tracing and no matrix/tensor hardware; Metal's ray-tracing and mesh-shading APIs execute on the ordinary shader ALUs. Apple provides no native Vulkan implementation (MoltenVK translation only), and OpenGL and OpenCL are deprecated on macOS", + "Memory": ["L2 Cache (GPU)", "System Level Cache", "Unified Memory (LPDDR5)"], + "Subunits": ["GPU Core"] }, - { - "Type": "Tile Memory", + "GPU Core": { "Count": 32, + "Size": 4, + "Description": "Apple GPU core with 128 FP32 ALUs organised as 4 SIMD pipelines of 32 lanes, giving 4096 ALUs across the 32-core GPU. Apple does not publish per-core ALU counts; the 128-ALU figure is third-party, and the core is the same Apple7 design used in the base M1 and M1 Pro. FP16 executes at the same peak rate as FP32 on this family", + "Memory": ["Register File", "Threadgroup Memory", "L1 Data Cache", "L1 Instruction Cache"], + "Subunits": ["SIMD Pipeline"] + }, + "SIMD Pipeline": { + "Count": 4, "Size": 32, - "Description": "On-chip tile memory for efficient rendering", - "Memory": ["Tile"], - "SubunitType": "Memory" + "Description": "32-lane SIMD execution pipeline issuing one instruction per cycle for one 32-thread SIMD-group. The 32-thread SIMD-group width is a run-time Metal property (threadExecutionWidth) documented for this G13 architecture by public ISA reverse-engineering; the 4-pipeline decomposition of the core's 128 ALUs is third-party analysis, not an Apple disclosure" } - ] + } }, "MemorySubsystem": { "SupportedMemoryTypes": ["LPDDR5"], "MemoryTypes": [ { - "Type": "L1", - "Size": 64, - "BankCount": 32, - "MaxMemoryBandwidth": 400, - "maxBusWidth": 512, - "description": "L1 cache per GPU core" + "Type": "Register File", + "Size": 208, + "description": "Approximately 208 KB of registers per GPU core across the Apple7 and Apple8 generations. Not published by Apple; measured by a single public microbenchmarking project" }, { - "Type": "L2", - "Size": 16384, - "BankCount": 1, - "MaxMemoryBandwidth": 400, - "maxBusWidth": 512, - "description": "Shared L2 cache" + "Type": "Threadgroup Memory", + "Size": 32, + "description": "Apple's Metal Feature Set Tables cap a single threadgroup at 32 KB of threadgroup memory on the Apple7 family, with 1024 maximum threads per threadgroup, 16 B length alignment, 32 KB maximum explicit imageblock allocation and 128 KB implicit. Public microbenchmarking measures a physical per-core pool of roughly 60 KB" + }, + { + "Type": "L1 Data Cache", + "Size": 8, + "description": "8 KB L1 data cache per GPU core, unchanged across the Apple7 family. Apple publishes no GPU cache sizes; measured by two independent public microbenchmarking efforts" + }, + { + "Type": "L1 Instruction Cache", + "Size": 12, + "description": "12 KB L1 instruction cache per GPU core. Not published by Apple; single public microbenchmarking source" + }, + { + "Type": "L2 Cache (GPU)", + "Size": 512, + "description": "GPU-wide L2 cache, measured at 512 KB on the M1 Max by a single public microbenchmarking project. Apple publishes nothing here, and the same source measures 768 KB on the base M1 and 256 KB on the M1 Pro, so capacity varies non-monotonically across the family and cannot be derived by scaling core counts. Treat as indicative, single-source only" + }, + { + "Type": "System Level Cache", + "Size": 49152, + "description": "48 MB SoC-wide system level cache shared by CPU, GPU, Neural Engine and media engine, sitting behind the GPU L2 and in front of LPDDR5. Not published by Apple; measured third-party, with two independent sources agreeing on 48 MB for M1 Max against 24 MB for M1 Pro and 8 MB for M1" }, { - "Type": "Unified Memory", + "Type": "Unified Memory (LPDDR5)", "Size": 67108864, - "BankCount": 16, "MaxMemoryBandwidth": 400, "maxBusWidth": 512, - "description": "Unified memory shared with CPU" + "description": "LPDDR5 unified memory shared with the CPU, Neural Engine and media engine, with no separate VRAM and no copy across a bus. Apple publishes 'up to 400GB/s' of memory bandwidth - 2x that of M1 Pro and nearly 6x that of M1 - and support for up to 64 GB of unified memory; the 64 GB maximum is recorded here in KB. Apple does not state the bus width: the 512-bit LPDDR5-6400 interface is third-party analysis, and its 409.6 GB/s theoretical peak is what Apple rounds down to 400 GB/s" } ] }, "KernelModel": { + "LLVMTarget": "air64", + "LLVMTriple": "air64-apple-macos", + "LLVMFeatures": [ + { + "Name": "air64", + "Description": "Apple Intermediate Representation (AIR), the LLVM-bitcode-based IR that Apple's metal compiler emits into a .metallib (-target air64-apple-macos). AIR is Apple-private and is not an upstream LLVM backend target; Apple publishes no GPU ISA documentation" + } + ], "SubUnits": [ { "Type": "Metal", - "Version": "2.4" + "Version": "3" }, { - "Type": "OpenGL ES", - "Version": "3.0" + "Type": "Metal", + "Version": "4" + }, + { + "Type": "OpenGL", + "Version": "4.1" }, { - "Type": "Vulkan", + "Type": "OpenCL", "Version": "1.2" } ] }, "SpecializedHardware": { "RayTracingAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "Performance": { - "RaysPerSecond": 0 - } + "Present": false }, "AiAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "SupportedPrecisions": [], - "Performance": { - "fp16TopsPerGPU": 0, - "int8TopsPerGPU": 0 - } + "Present": true, + "Name": "16-core Neural Engine", + "Count": 16 }, "videoCodecs": { "encoders": [ { - "codec": "H.264", - "maxResolution": "4K", - "maxBitrate": 240, - "maxFPS": 60 + "codec": "H.264" }, { - "codec": "H.265/HEVC", - "maxResolution": "4K", - "maxBitrate": 240, - "maxFPS": 60 + "codec": "H.265/HEVC" + }, + { + "codec": "Other", + "maxResolution": "8K" } ], "decoders": [ { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 480, - "maxFPS": 30 + "codec": "H.264" }, { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 480, - "maxFPS": 30 + "codec": "H.265/HEVC" }, { - "codec": "VP9", - "maxResolution": "4K", - "maxBitrate": 240, - "maxFPS": 60 + "codec": "Other", + "maxResolution": "8K" } ] } }, "PowerEfficiency": { - "BaseClock": 1000, "BoostClock": 1296, - "MaxTDP": 60, "PowerStates": [ { - "Name": "Performance", - "Description": "Maximum performance for demanding workloads" - }, - { - "Name": "Balanced", - "Description": "Balance between performance and power efficiency" + "Name": "Active", + "Description": "GPU active and clocking up to its maximum. Apple publishes neither GPU clocks nor a TDP for M1 Max, so no MaxTDP and no BaseClock are recorded. The 1296 MHz maximum comes from third-party measurement of an M1 Max and is corroborated by a second independent source" }, { - "Name": "Power Saving", - "Description": "Optimized for battery life" + "Name": "Idle", + "Description": "Reduced power state when no GPU work is queued" } ], "ClockGating": true, "DynamicVoltageFrequencyScaling": true }, - "PciExpress": { - "Version": "N/A", - "Lanes": 0, - "Bandwidth": 0 - }, "MultiGpuSupport": { "Technologies": ["None"], - "MaxGpus": 1, - "InterconnectBandwidth": 0 + "MaxGpus": 1 } }, { "Name": "Apple GPU (M2)", "Vendor": "Apple", - "Architecture": "Apple GPU Gen 2", + "Architecture": "Apple8", "ReleaseYear": 2022, "FabricationProcess": { "ProcessNode": 5, @@ -625,162 +597,161 @@ }, "CoreSubsystem": { "Name": "GPU Core", - "subUnits": [ - { - "Type": "Execution Unit", + "ChipType": "GPU", + "CoreType": "GPU Core", + "UnitTypes": { + "GPU": { + "Count": 1, + "Size": 10, + "Description": "Integrated GPU block of the Apple M2 SoC, which Apple states has 20 billion transistors on second-generation 5-nanometer process technology. Apple ships 10-core and binned 8-core GPU configurations; the 10-core part is described here. Metal reports GPU family Apple8 - Apple's Metal Feature Set Tables list the whole M2-series as Apple8, alongside A15 and A16 Bionic. Apple publishes 100 GB/s of unified memory bandwidth for M2 but no teraflops rating; the widely cited 3.6 TFLOPS is third-party arithmetic (10 cores x 128 ALUs x 2 FLOP x 1.398 GHz = 3.58 TFLOPS). The GPU shares LPDDR5 unified memory with the CPU, the 16-core Neural Engine (Apple-published at 15.8 trillion operations per second, precision unstated) and the media engine, which has hardware H.264, HEVC, ProRes and ProRes RAW acceleration. AV1 decode is absent on M2 - Apple states it arrived with the M3 media engine for the first time. There is no ray-tracing and no matrix/tensor hardware; Metal's ray-tracing, mesh-shading and SIMD-scoped matrix multiply APIs execute on the ordinary shader ALUs. Apple provides no native Vulkan implementation (MoltenVK translation only), and OpenGL and OpenCL are deprecated on macOS", + "Memory": ["L2 Cache (GPU)", "System Level Cache", "Unified Memory (LPDDR5)"], + "Subunits": ["GPU Core"] + }, + "GPU Core": { "Count": 10, - "Size": 128, - "Description": "Enhanced unified shader processors with 128 ALUs per core", - "Memory": ["L1", "L2"], - "SubunitType": "SIMD" + "Size": 4, + "Description": "Apple GPU core with 128 FP32 ALUs organised as 4 SIMD pipelines of 32 lanes, giving 1280 ALUs across the 10-core GPU. Apple does not publish per-core ALU counts; the 128-ALU figure is third-party and consistent across the Apple7 and Apple8 generations. Apple8 executes FP16 and FP32 at the same peak rate, and the core contains no matrix or tensor units", + "Memory": ["Register File", "Threadgroup Memory", "L1 Data Cache", "L1 Instruction Cache"], + "Subunits": ["SIMD Pipeline"] }, - { - "Type": "Tile Memory", - "Count": 10, + "SIMD Pipeline": { + "Count": 4, "Size": 32, - "Description": "On-chip tile memory for efficient rendering", - "Memory": ["Tile"], - "SubunitType": "Memory" + "Description": "32-lane SIMD execution pipeline issuing one instruction per cycle for one 32-thread SIMD-group. The SIMD-group width is a run-time Metal property (threadExecutionWidth) rather than a tabulated feature-set limit; the 4 x 32 decomposition of the core's 128 ALUs is not published by Apple and comes from public reverse-engineering of the Apple G13/G14 GPU ISA" } - ] + } }, "MemorySubsystem": { "SupportedMemoryTypes": ["LPDDR5"], "MemoryTypes": [ { - "Type": "L1", - "Size": 64, - "BankCount": 10, - "MaxMemoryBandwidth": 100, - "maxBusWidth": 128, - "description": "L1 cache per GPU core" + "Type": "Register File", + "Size": 208, + "description": "Approximately 208 KB of registers per GPU core across the Apple7 and Apple8 generations. Not published by Apple; measured by public microbenchmarking. The register file is large relative to Apple's very small L1 and L2, and occupancy on Apple8 ranges from 384 to 3072 resident threads per core depending on register pressure" }, { - "Type": "L2", - "Size": 4096, - "BankCount": 1, - "MaxMemoryBandwidth": 100, - "maxBusWidth": 128, - "description": "Shared L2 cache" + "Type": "Threadgroup Memory", + "Size": 32, + "description": "Apple's Metal Feature Set Tables cap a single threadgroup at 32 KB of threadgroup memory on the Apple8 family, with 1024 maximum threads per threadgroup, 16 B length alignment, 32 KB maximum explicit imageblock allocation and 128 KB implicit. The physical per-core pool is larger than the per-threadgroup cap - public microbenchmarking measures roughly 60 KB per core - which allows several threadgroups to be resident at once" + }, + { + "Type": "L1 Data Cache", + "Size": 8, + "description": "8 KB L1 data cache per GPU core with a 128-byte global cache line. Not published by Apple; measured by public microbenchmarking. Deliberately small compared with contemporary AMD and NVIDIA cores, offset by a large register file and a very low-latency L2" }, { - "Type": "Unified Memory", + "Type": "L1 Instruction Cache", + "Size": 12, + "description": "12 KB L1 instruction cache per GPU core. Not published by Apple; measured by public microbenchmarking" + }, + { + "Type": "L2 Cache (GPU)", + "Size": 1536, + "description": "GPU-wide L2 cache, approximately 1.5 MB on the base M2. Apple publishes no GPU cache sizes. This figure rests on a single public microbenchmarking estimate and is explicitly marked approximate by its source; independent measurements of the M1 and of the M2 Pro die report different capacities, so treat it as indicative only. Apple L2 capacity is known to vary non-monotonically across a family" + }, + { + "Type": "System Level Cache", + "Size": 8192, + "description": "8 MB SoC-wide system level cache shared by the CPU, GPU, Neural Engine and media engine, backing the GPU L2 in front of LPDDR5. Not published by Apple; measured by public microbenchmarking" + }, + { + "Type": "Unified Memory (LPDDR5)", "Size": 25165824, - "BankCount": 6, "MaxMemoryBandwidth": 100, "maxBusWidth": 128, - "description": "Unified memory shared with CPU" + "description": "LPDDR5 unified memory shared with the CPU, Neural Engine and media engine; there is no separate VRAM and no copy across a bus. Apple publishes 100 GB/s of unified memory bandwidth and 8 GB, 16 GB and 24 GB configurations; the 24 GB maximum is recorded here in KB. Apple does not publish the bus width or the LPDDR5 speed grade; public die-shot analysis identifies a 128-bit LPDDR5 bus at 6400 MT/s, which yields 102.4 GB/s and is consistent with Apple's rounded 100 GB/s" } ] }, "KernelModel": { + "LLVMTarget": "air64", + "LLVMTriple": "air64-apple-macos", + "LLVMFeatures": [ + { + "Name": "air64", + "Description": "Apple Intermediate Representation (AIR), the LLVM-bitcode-based IR that Apple's metal compiler emits into a .metallib (-target air64-apple-macos, or air64-apple-ios for iPadOS). AIR is Apple-private and is not an upstream LLVM backend target; Apple publishes no GPU ISA documentation" + } + ], "SubUnits": [ { "Type": "Metal", - "Version": "3.0" + "Version": "3" }, { - "Type": "OpenGL ES", - "Version": "3.0" + "Type": "Metal", + "Version": "4" + }, + { + "Type": "OpenGL", + "Version": "4.1" }, { - "Type": "Vulkan", - "Version": "1.3" + "Type": "OpenCL", + "Version": "1.2" } ] }, "SpecializedHardware": { "RayTracingAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "Performance": { - "RaysPerSecond": 0 - } + "Present": false }, "AiAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "SupportedPrecisions": [], - "Performance": { - "fp16TopsPerGPU": 0, - "int8TopsPerGPU": 0 - } + "Present": true, + "Name": "16-core Neural Engine", + "Count": 16 }, "videoCodecs": { "encoders": [ { - "codec": "H.264", - "maxResolution": "4K", - "maxBitrate": 180, - "maxFPS": 60 + "codec": "H.264" }, { - "codec": "H.265/HEVC", - "maxResolution": "4K", - "maxBitrate": 180, - "maxFPS": 60 + "codec": "H.265/HEVC" + }, + { + "codec": "Other" } ], "decoders": [ { "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 360, - "maxFPS": 30 + "maxResolution": "8K" }, { "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 360, - "maxFPS": 30 + "maxResolution": "8K" }, { - "codec": "VP9", - "maxResolution": "4K", - "maxBitrate": 180, - "maxFPS": 60 + "codec": "Other", + "maxResolution": "8K" } ] } }, "PowerEfficiency": { - "BaseClock": 1000, "BoostClock": 1398, - "MaxTDP": 20, "PowerStates": [ { - "Name": "Performance", - "Description": "Maximum performance for demanding workloads" - }, - { - "Name": "Balanced", - "Description": "Balance between performance and power efficiency" + "Name": "Active", + "Description": "GPU active and clocking up to its maximum. Apple publishes neither GPU clocks nor a TDP for the M2, so no MaxTDP is recorded here. The 1398 MHz maximum is a third-party figure with two independent backers; one die-shot analysis instead reports 1406 MHz, and that discrepancy is unresolved. The minimum GPU clock for the base M2 is not publicly established, so no BaseClock is recorded" }, { - "Name": "Power Saving", - "Description": "Optimized for battery life" + "Name": "Idle", + "Description": "Reduced power state when no GPU work is queued" } ], "ClockGating": true, "DynamicVoltageFrequencyScaling": true }, - "PciExpress": { - "Version": "N/A", - "Lanes": 0, - "Bandwidth": 0 - }, "MultiGpuSupport": { "Technologies": ["None"], - "MaxGpus": 1, - "InterconnectBandwidth": 0 + "MaxGpus": 1 } }, { "Name": "Apple GPU (M3)", "Vendor": "Apple", - "Architecture": "Apple GPU Gen 3", + "Architecture": "Apple9", "ReleaseYear": 2023, "FabricationProcess": { "ProcessNode": 3, @@ -789,176 +760,135 @@ }, "CoreSubsystem": { "Name": "GPU Core", - "subUnits": [ - { - "Type": "Execution Unit", - "Count": 10, - "Size": 128, - "Description": "Next-generation unified shader processors with hardware ray tracing support", - "Memory": ["L1", "L2"], - "SubunitType": "SIMD" - }, - { - "Type": "Ray Tracing Unit", + "ChipType": "GPU", + "CoreType": "GPU Core", + "UnitTypes": { + "GPU": { + "Count": 1, + "Size": 10, + "Description": "Integrated GPU block of the Apple M3 SoC, which Apple describes as one of 'the industry's first 3-nanometer chips for a personal computer' with 25 billion transistors (Apple names neither foundry nor node name; TSMC N3B is the third-party attribution, supported by Apple later calling M4's node 'second-generation 3-nanometer'). Apple ships 10-core and binned 8-core GPU configurations; the 10-core part is described here. Metal reports GPU family Apple9, alongside A17 Pro. Apple calls this the next-generation GPU architecture and states it introduces Dynamic Caching, which 'unlike traditional GPUs, allocates the use of local memory in hardware in real time' so that only the exact amount of memory needed is used per task, and brings hardware-accelerated ray tracing and mesh shading to the Mac for the first time. Apple publishes 100 GB/s of unified memory bandwidth and up to 24 GB of unified memory for M3, but no teraflops rating and no GPU clock; the frequently quoted ~1400 MHz figure has no traceable source and is therefore not recorded. The GPU shares unified memory with the CPU, the 16-core Neural Engine (Apple states only that it is up to 60 percent faster than the M1 family; the '18 TOPS' figure in circulation is press arithmetic, not an Apple number) and the media engine, which adds AV1 decode. Apple provides no native Vulkan implementation (MoltenVK translation only), and OpenGL and OpenCL are deprecated on macOS", + "Memory": ["Unified Memory (LPDDR5)"], + "Subunits": ["GPU Core"] + }, + "GPU Core": { "Count": 10, - "Size": 1, - "Description": "Dedicated hardware ray tracing acceleration units", - "Memory": ["L1"], - "SubunitType": "RT" + "Size": 5, + "Description": "Apple9 'next-generation shader core', which Apple describes as pairing the shader ALUs with a texture unit and a new ray tracing unit. 128 FP32 ALUs per core arranged as 4 SIMD pipelines of 32 lanes is carried forward from the Apple7/Apple8 generations and is not confirmed by Apple or by any measurement for Apple9, so it is stated here as an assumption. Apple publishes no per-core register file, L1 or L2 capacity for this generation, and Dynamic Caching makes the register, threadgroup-memory and stack split dynamic in hardware rather than a fixed partition, so no per-core cache sizes are listed", + "Memory": ["Threadgroup Memory"], + "Subunits": ["SIMD Pipeline", "Ray Tracing Unit"] }, - { - "Type": "Tile Memory", - "Count": 10, + "SIMD Pipeline": { + "Count": 4, "Size": 32, - "Description": "On-chip tile memory for efficient rendering", - "Memory": ["Tile"], - "SubunitType": "Memory" + "Description": "32-lane SIMD execution pipeline issuing one instruction per cycle for one 32-thread SIMD-group. The 32-thread SIMD-group width is a run-time Metal property (threadExecutionWidth) and is not tabulated per family in Apple's feature-set tables; the 4-pipeline decomposition of the core's 128 ALUs is inferred from earlier Apple GPU generations rather than documented for Apple9" + }, + "Ray Tracing Unit": { + "Count": 1, + "Description": "Dedicated ray intersection hardware. Apple's Tech Talk 'Explore GPU advancements in M3 and A17 Pro' describes the new shader core as paired with 'a brand new ray tracing unit that accelerates ray intersection requests', which is the basis for one unit per shader core here. Apple publishes no ray-throughput figure, so no rays-per-second rating is recorded. Note that the Metal ray-tracing API itself is available from the Apple6 family; what Apple9 adds is hardware acceleration of it, plus Apple9-exclusive features such as ray tracing with per-component motion interpolation" } - ] + } }, "MemorySubsystem": { "SupportedMemoryTypes": ["LPDDR5"], "MemoryTypes": [ { - "Type": "L1", - "Size": 64, - "BankCount": 10, - "MaxMemoryBandwidth": 100, - "maxBusWidth": 128, - "description": "L1 cache per GPU core" - }, - { - "Type": "L2", - "Size": 4096, - "BankCount": 1, - "MaxMemoryBandwidth": 100, - "maxBusWidth": 128, - "description": "Shared L2 cache" + "Type": "Threadgroup Memory", + "Size": 32, + "description": "Apple's Metal Feature Set Tables cap a single threadgroup at 32 KB of threadgroup memory on the Apple9 family, with 1024 maximum threads per threadgroup, 16 B length alignment, 32 KB maximum explicit imageblock allocation and 256 KB implicit (doubled from the 128 KB of Apple7/Apple8). On Apple9 this storage is drawn from the Dynamic Caching on-chip pool that also backs registers and stack, so the physical pool size per core is not a published constant" }, { - "Type": "Unified Memory", + "Type": "Unified Memory (LPDDR5)", "Size": 25165824, - "BankCount": 6, "MaxMemoryBandwidth": 100, "maxBusWidth": 128, - "description": "Unified memory shared with CPU" + "description": "LPDDR5 unified memory shared by CPU, GPU, Neural Engine and media engine with zero-copy access and no separate VRAM. Apple's Mac technical specifications state 100 GB/s of memory bandwidth, and Apple's announcement states support for up to 24 GB of unified memory; the 24 GB maximum is recorded here in KB. Apple does not state the bus width - the 128-bit figure (8 x 16-bit LPDDR5-6400 controllers) is third-party analysis. Apple publishes no GPU L1/L2 or system-level cache sizes for M3, and unlike the M1 and M2 generations no public microbenchmarking of the M3 GPU cache hierarchy was found, so no cache levels are listed rather than carrying earlier figures forward" } ] }, "KernelModel": { + "LLVMTarget": "air64", + "LLVMTriple": "air64-apple-macos", + "LLVMFeatures": [ + { + "Name": "air64", + "Description": "Apple Intermediate Representation (AIR), the LLVM-bitcode-based IR that Apple's metal compiler emits into a .metallib (-target air64-apple-macos). AIR is Apple-private and is not an upstream LLVM backend target; Apple publishes no GPU ISA documentation" + } + ], "SubUnits": [ { "Type": "Metal", - "Version": "3.1" + "Version": "3" }, { - "Type": "OpenGL ES", - "Version": "3.0" + "Type": "Metal", + "Version": "4" + }, + { + "Type": "OpenGL", + "Version": "4.1" }, { - "Type": "Vulkan", - "Version": "1.3" + "Type": "OpenCL", + "Version": "1.2" } ] }, "SpecializedHardware": { "RayTracingAccelerators": { "Present": true, - "Name": "RT Unit", - "Count": 1, - "Performance": { - "RaysPerSecond": 2.5 - } + "Name": "Ray tracing unit", + "Count": 1 }, "AiAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "SupportedPrecisions": [], - "Performance": { - "fp16TopsPerGPU": 0, - "int8TopsPerGPU": 0 - } + "Present": true, + "Name": "16-core Neural Engine", + "Count": 16 }, "videoCodecs": { "encoders": [ { - "codec": "H.264", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 + "codec": "H.264" }, { - "codec": "H.265/HEVC", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 + "codec": "H.265/HEVC" }, { - "codec": "AV1", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 + "codec": "Other" } ], "decoders": [ { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 400, - "maxFPS": 30 + "codec": "H.264" }, { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 400, - "maxFPS": 30 + "codec": "H.265/HEVC" }, { - "codec": "VP9", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 + "codec": "Other" }, { - "codec": "AV1", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 + "codec": "AV1" } ] } }, "PowerEfficiency": { - "BaseClock": 1000, - "BoostClock": 1424, - "MaxTDP": 22, "PowerStates": [ { - "Name": "Performance", - "Description": "Maximum performance for demanding workloads" - }, - { - "Name": "Balanced", - "Description": "Balance between performance and power efficiency" + "Name": "Active", + "Description": "GPU active. Apple publishes neither GPU clock frequencies nor a chip- or GPU-level TDP for M3, so BaseClock, BoostClock and MaxTDP are all omitted rather than estimated. No third-party measurement of the base M3 GPU clock could be substantiated either, so the 1424 MHz figure that circulates for this part is not recorded" }, { - "Name": "Power Saving", - "Description": "Optimized for battery life" + "Name": "Idle", + "Description": "Low-power state when no GPU work is scheduled; frequency and voltage are managed by the SoC power controller" } ], "ClockGating": true, "DynamicVoltageFrequencyScaling": true }, - "PciExpress": { - "Version": "N/A", - "Lanes": 0, - "Bandwidth": 0 - }, "MultiGpuSupport": { "Technologies": ["None"], - "MaxGpus": 1, - "InterconnectBandwidth": 0 + "MaxGpus": 1 } } ] diff --git a/device_lib/apple-gpu-applegpu_g16s.json b/device_lib/apple-gpu-applegpu_g16s.json index 75da8229..82a1c23a 100644 --- a/device_lib/apple-gpu-applegpu_g16s.json +++ b/device_lib/apple-gpu-applegpu_g16s.json @@ -1,228 +1,143 @@ { - "Name": "Apple GPU G16S", + "Name": "Apple GPU G16S (M4 Pro)", "Vendor": "Apple", - "Type": "gpu", "Architecture": "applegpu_g16s", "ReleaseYear": 2024, "FabricationProcess": { "ProcessNode": 3, "Manufacturer": "TSMC", - "Technology": "N3E" + "Technology": "N3E (Apple: second-generation 3-nanometer)" }, "CoreSubsystem": { - "Name": "Core", - "SubUnits": [ - { - "Type": "Chip", + "Name": "GPU Core", + "ChipType": "Chip", + "CoreType": "GPU Core", + "UnitTypes": { + "Chip": { "Count": 1, - "Size": 10, - "Description": "Chip containing multiple Clusters", - "Memory": ["L2 Cache"], - "SubunitType": "GPU Cluster" + "Size": 20, + "Description": "Integrated GPU block of the Apple M4 Pro SoC (SoC codename H16S / Brava Chop, t6040). Apple GPU generation G16, die-tier suffix S = Pro (matching the G13G = M1, G13S = M1 Pro naming reported by the Mesa/Asahi driver). Ships in 16-core and 20-core GPU bins; Size 20 is the full configuration. Apple GPUs group cores into clusters, but the cluster geometry of the G16 generation is not published, so no cluster level is modeled here", + "Memory": ["L2 / System Level Cache", "Unified Memory (LPDDR5X)"], + "Subunits": ["GPU Core"] }, - { - "Type": "GPU Cluster", - "Count": 10, - "Size": 2, - "Description": "High-level GPU cluster containing multiple cores", - "SubunitType": "Core" - }, - { - "Type": "Core", - "Count": 2, - "Size": 4, - "ISA": "", - "Description": "Individual GPU compute core with 128 ALUs", - "Memory": ["L1 Cache", "Tile Memory", "Vector Register File", "Scalar Register File", "Shared Memory"], - "SubunitType": "SIMD" + "GPU Core": { + "Count": 20, + "Size": 1, + "Description": "Tile-based deferred rendering (TBDR) shader core of Metal GPU family Apple9 (A17 Pro, M3-series and M4-series per Apple's Metal Feature Set Tables). Features Dynamic Caching (on-chip memory allocated dynamically among registers, threadgroup memory and stack), a hardware ray tracing accelerator (Apple: 2x the M3 engine) and hardware mesh shading. There is no dedicated matrix engine in this generation: SIMD-scoped matrix multiply operations (Apple7 and later) execute on the shader ALUs. Apple does not publish per-core ALU counts, cache sizes or clock frequencies", + "Memory": ["L1 Cache (per GPU core)", "Threadgroup Memory", "Imageblock (Tile) Memory"], + "Subunits": ["SIMD Group"] }, - { - "Type": "SIMD", - "Count": 4, + "SIMD Group": { + "Count": 1, "Size": 32, - "Description": "SIMD with Warp Size" - }, - { - "Type": "Neural Engine Core", - "Count": 16, - "Size": 1, - "Description": "Dedicated neural processing unit for AI/ML workloads" + "Description": "Execution is 32-wide: an Apple GPU SIMD group (Metal threadExecutionWidth) is 32 threads, and up to 1024 threads may be launched per threadgroup on Apple9. Count is 1 because Apple does not publish how many 32-wide SIMD pipelines each core contains; the value describes the execution width itself, not an ALU-pipeline count" } - ] + } }, "MemorySubsystem": { "SupportedMemoryTypes": ["LPDDR5"], - "SubUnits": [ + "MemoryTypes": [ { - "Type": "Vector Register File", - "Size": 256, - "BankCount": 4, - "Description": "Vector register file for SIMD operations per GPU core" + "Type": "L1 Cache (per GPU core)", + "description": "Per-core cache and Dynamic Caching pool backing registers, threadgroup memory, tile memory and stack. Apple does not publish its capacity, banking or line size, so no Size is given" }, { - "Type": "Scalar Register File", - "Size": 16, - "BankCount": 2, - "Description": "Scalar register file for control flow and addressing per GPU core" - }, - { - "Type": "Shared Memory", + "Type": "Threadgroup Memory", "Size": 32, - "BankCount": 32, - "Description": "Threadgroup shared memory for inter-thread communication within a GPU core" - }, - { - "Type": "L1 Cache", - "Size": 16, - "BankCount": 4, - "Description": "Per-core L1 cache for fast data access" + "description": "Maximum total threadgroup memory allocation is 32 KB per threadgroup on GPU family Apple9 (Apple Metal Feature Set Tables). Threadgroup memory length alignment is 16 B. On M3 and later this storage is carved out of the core's on-chip memory by Dynamic Caching rather than being statically partitioned" }, { - "Type": "L2 Cache", - "Size": 4096, - "BankCount": 8, - "Description": "Shared L2 cache across GPU cores" + "Type": "Imageblock (Tile) Memory", + "Size": 32, + "description": "On-chip tile storage used by TBDR rendering and imageblocks. Apple9 allows up to 32 KB of explicit imageblock allocation and up to 256 KB of implicit imageblock allocation; explicit imageblock and threadgroup allocations share the same budget" }, { - "Type": "Tile Memory", - "Size": 32, - "BankCount": 1, - "Description": "On-chip tile memory for efficient rendering" + "Type": "L2 / System Level Cache", + "description": "GPU-shared L2 plus the SoC system level cache in front of unified memory. Apple publishes no capacity, bandwidth or banking figures for either level, so no Size is given" }, { - "Type": "System Memory", - "Size": 0, + "Type": "Unified Memory (LPDDR5X)", + "Size": 67108864, "MaxMemoryBandwidth": 273, "maxBusWidth": 256, - "Description": "Unified memory architecture shared with CPU" + "description": "On-package LPDDR5X at up to 8533 MT/s on a 256-bit bus, giving Apple's published 273 GB/s. Unified memory is shared coherently with the CPU, Neural Engine and media engine. Size is the maximum 64 GB configuration (67108864 KB); M4 Pro also ships with 24 GB and 48 GB. LPDDR5X is reported as LPDDR5 because the schema enumeration has no LPDDR5X value" } ] }, "KernelModel": { - "LLVMTriple": "aarch64-apple-macosx", - "LLVMFeatures": [ - { - "Name": "apple-gpu", - "Description": "Apple GPU specific optimizations" - }, + "LLVMTarget": "air64", + "LLVMTriple": "air64-apple-macosx", + "SubUnits": [ { - "Name": "neon", - "Description": "Advanced SIMD (NEON) support" + "Type": "metal", + "Version": "3" }, - { - "Name": "fp16", - "Description": "Half-precision floating-point support" - } - ], - "SubUnits": [ { "Type": "metal", - "Version": "3.2" + "Version": "4" }, { "Type": "opencl", - "Version": "3.0" + "Version": "1.2" } ] }, "SpecializedHardware": { "RayTracingAccelerators": { "Present": true, - "Name": "Ray Tracing Cores", - "Count": 1, - "Performance": { - "RaysPerSecond": 2.5 - } + "Name": "Hardware-accelerated ray tracing (second-generation ray tracing engine)" }, "AiAccelerators": { "Present": true, - "Name": "Neural Engine", + "Name": "Neural Engine (SoC-level block, not part of the GPU)", "Count": 16, - "SupportedPrecisions": ["FP32", "FP16", "INT8"], + "SupportedPrecisions": ["FP16", "INT8"], "Performance": { - "fp16TopsPerGPU": 35.0, - "int8TopsPerGPU": 70.0 + "int8TopsPerGPU": 38 } }, "videoCodecs": { "encoders": [ { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 1000, - "maxFPS": 60 + "codec": "H.264" }, { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 1000, - "maxFPS": 60 + "codec": "H.265/HEVC" }, { - "codec": "AV1", - "maxResolution": "4K", - "maxBitrate": 500, - "maxFPS": 60 + "codec": "Other" } ], "decoders": [ { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 1000, - "maxFPS": 120 + "codec": "H.264" }, { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 1000, - "maxFPS": 120 + "codec": "H.265/HEVC" }, { - "codec": "AV1", - "maxResolution": "8K", - "maxBitrate": 800, - "maxFPS": 60 + "codec": "AV1" }, { - "codec": "VP9", - "maxResolution": "8K", - "maxBitrate": 800, - "maxFPS": 60 + "codec": "Other" } ] } }, "PowerEfficiency": { - "BaseClock": 1300, - "BoostClock": 1580, - "MaxTDP": 40, "PowerStates": [ { "Name": "Active", - "Description": "Full performance mode" - }, - { - "Name": "Automatic", - "Description": "Dynamic performance scaling based on workload" - }, - { - "Name": "Low Power", - "Description": "Reduced performance for battery efficiency" + "Description": "GPU executing work. Apple publishes neither GPU clock frequencies nor a GPU or package TDP for M4 Pro, so BaseClock, BoostClock and MaxTDP are omitted rather than estimated" }, { "Name": "Idle", - "Description": "Minimal power consumption when not in use" + "Description": "Low-power state when no Metal work is queued; power is managed by the SoC power controller" } ], "ClockGating": true, "DynamicVoltageFrequencyScaling": true }, - "PciExpress": { - "Version": "4.0", - "Lanes": 16, - "Bandwidth": 32 - }, "MultiGpuSupport": { "Technologies": ["None"], "MaxGpus": 1, diff --git a/device_lib/apple-gpu-m1m.json b/device_lib/apple-gpu-m1m.json index eafb3bce..7b21c659 100644 --- a/device_lib/apple-gpu-m1m.json +++ b/device_lib/apple-gpu-m1m.json @@ -1,164 +1,147 @@ { - "Name": "Apple GPU (M1 Max)", - "Vendor": "Apple", - "Architecture": "Apple GPU Gen 1 Max", - "ReleaseYear": 2021, - "FabricationProcess": { - "ProcessNode": 5, - "Manufacturer": "TSMC", - "Technology": "N5P" + "Name": "Apple GPU (M1 Max)", + "Vendor": "Apple", + "Architecture": "Apple7", + "ReleaseYear": 2021, + "FabricationProcess": { + "ProcessNode": 5, + "Manufacturer": "TSMC", + "Technology": "N5" + }, + "CoreSubsystem": { + "Name": "GPU Core", + "ChipType": "GPU", + "CoreType": "GPU Core", + "UnitTypes": { + "GPU": { + "Count": 1, + "Size": 32, + "Description": "Integrated GPU block of the M1 Max SoC (57 billion transistors total). Up to 32 GPU cores, 4,096 ALUs, approximately 10.4 TFLOPS FP32 at the peak 1296 MHz clock. Cores share the SoC's system level cache and the unified memory pool with the CPU, Neural Engine, and media engine. The 16-core Neural Engine and the ProRes/video media engines are separate SoC blocks and are not part of the GPU or addressable from Metal. A 24-core GPU binned variant of the same die is also sold; this entry describes the full 32-core configuration. Two M1 Max dies can be joined by Apple's UltraFusion silicon interposer at 2.5 TB/s to form the 64-core-GPU M1 Ultra", + "Memory": ["System Level Cache", "Unified Memory (LPDDR5)"], + "Subunits": ["GPU Core"] + }, + "GPU Core": { + "Count": 32, + "Size": 128, + "Description": "Unified shader core. Apple publishes only the core count; the 128 ALUs per core figure is derived from the widely reported 4,096 total ALUs divided across 32 cores. Each core has its own L1 texture cache and L1 buffer cache (Apple confirms both exist but publishes no capacities) plus 32 KB of on-chip threadgroup/imageblock memory. Threads execute in SIMD groups of 32", + "Memory": ["L1 Texture Cache", "L1 Buffer Cache", "Threadgroup Memory", "Imageblock (Tile) Memory"], + "Subunits": ["ALU"] + }, + "ALU": { + "Count": 128, + "Size": 1, + "Description": "Scalar shader lane executing FP32/FP16/integer operations. Work is issued as 32-wide SIMD groups (Metal threadExecutionWidth is 32 on Apple GPUs), and up to 1024 threads may occupy one threadgroup. Apple does not publish how the 128 lanes of a core are partitioned into issue units, so no finer hierarchy is modeled here. Apple7 adds SIMD-scoped reduction and SIMD-scoped matrix multiply (simdgroup_matrix) operations, which run on these lanes rather than on dedicated matrix hardware" + } + } + }, + "MemorySubsystem": { + "SupportedMemoryTypes": ["LPDDR5"], + "MemoryTypes": [ + { + "Type": "Threadgroup Memory", + "Size": 32, + "description": "32 KB maximum total threadgroup memory allocation per threadgroup on the Apple7 GPU family, per Apple's Metal Feature Set Tables. Shared with explicit imageblock (tile) memory: the sum of the two allocations cannot exceed the imageblock limit" + }, + { + "Type": "Imageblock (Tile) Memory", + "Size": 128, + "description": "On-chip tile storage for the tile-based deferred renderer. Apple7 allows up to 128 KB of implicit imageblock allocation and up to 32 KB of explicit imageblock allocation. Maximum tile size is 32x32 pixels without MSAA (32x16 with 4x MSAA)" + }, + { + "Type": "L1 Texture Cache", + "description": "Per-GPU-core read path cache for texture fetches, described in Apple's 'Metal Compute on MacBook Pro' tech talk. Apple does not publish its capacity and no authoritative measurement is available, so no size is recorded" + }, + { + "Type": "L1 Buffer Cache", + "description": "Per-GPU-core cache for buffer reads, separate from the texture cache, described in Apple's 'Metal Compute on MacBook Pro' tech talk. Capacity is not published by Apple, so no size is recorded" + }, + { + "Type": "System Level Cache", + "Size": 49152, + "MaxMemoryBandwidth": 400, + "description": "48 MB SoC-wide system level cache sitting in front of the memory controllers and shared by the GPU, CPU, and other SoC blocks. Apple does not publish this figure; 48 MB is AnandTech's measured value for M1 Max. There is no publicly established figure for a GPU-private L2, so none is listed" + }, + { + "Type": "Unified Memory (LPDDR5)", + "Size": 67108864, + "MaxMemoryBandwidth": 400, + "maxBusWidth": 512, + "description": "Up to 64 GB of LPDDR5-6400 unified memory on a 512-bit interface delivering Apple's stated 400 GB/s peak bandwidth, shared coherently by CPU, GPU, Neural Engine, and media engine with no separate VRAM copy. Sold in 32 GB and 64 GB configurations; Metal reports a recommended GPU working set of about 21 GB on a 32 GB system and about 48 GB on a 64 GB system" + } + ] + }, + "KernelModel": { + "LLVMTarget": "air64", + "LLVMTriple": "air64-apple-macos13.0", + "LLVMFeatures": [ + { + "Name": "air64", + "Description": "Apple Intermediate Representation. The Metal compiler is Clang/LLVM based and emits LLVM bitcode carrying an air64-apple-macos triple, but air64 is not an upstream LLVM target: it is only accepted by Apple's metal compiler shipped with Xcode" + } + ], + "SubUnits": [ + { + "Type": "Metal", + "Version": "4" + }, + { + "Type": "OpenCL", + "Version": "1.2" + } + ] + }, + "SpecializedHardware": { + "RayTracingAccelerators": { + "Present": false }, - "CoreSubsystem": { - "Name": "GPU Core", - "subUnits": [ - { - "Type": "Execution Unit", - "Count": 32, - "Size": 128, - "Description": "Unified shader processors with 128 ALUs per core", - "Memory": ["L1", "L2"], - "SubunitType": "SIMD" - }, - { - "Type": "Tile Memory", - "Count": 32, - "Size": 32, - "Description": "On-chip tile memory for efficient rendering", - "Memory": ["Tile"], - "SubunitType": "Memory" - } - ] + "AiAccelerators": { + "Present": false }, - "MemorySubsystem": { - "SupportedMemoryTypes": ["LPDDR5"], - "MemoryTypes": [ + "videoCodecs": { + "encoders": [ { - "Type": "L1", - "Size": 64, - "BankCount": 32, - "MaxMemoryBandwidth": 400, - "maxBusWidth": 512, - "description": "L1 cache per GPU core" + "codec": "H.264" }, { - "Type": "L2", - "Size": 16384, - "BankCount": 1, - "MaxMemoryBandwidth": 400, - "maxBusWidth": 512, - "description": "Shared L2 cache" + "codec": "H.265/HEVC" }, { - "Type": "Unified Memory", - "Size": 67108864, - "BankCount": 16, - "MaxMemoryBandwidth": 400, - "maxBusWidth": 512, - "description": "Unified memory shared with CPU" + "codec": "Other" } - ] - }, - "KernelModel": { - "SubUnits": [ + ], + "decoders": [ { - "Type": "Metal", - "Version": "2.4" + "codec": "H.264" }, { - "Type": "OpenGL ES", - "Version": "3.0" + "codec": "H.265/HEVC" }, { - "Type": "Vulkan", - "Version": "1.2" + "codec": "Other", + "maxResolution": "8K" } ] - }, - "SpecializedHardware": { - "RayTracingAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "Performance": { - "RaysPerSecond": 0 - } - }, - "AiAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "SupportedPrecisions": [], - "Performance": { - "fp16TopsPerGPU": 0, - "int8TopsPerGPU": 0 - } + } + }, + "PowerEfficiency": { + "BaseClock": 389, + "BoostClock": 1296, + "PowerStates": [ + { + "Name": "Active", + "Description": "The GPU runs at one of six measured DVFS steps (389, 486, 648, 778, 972, 1296 MHz); Apple publishes neither the frequency table nor a GPU TDP, so no MaxTDP is recorded here. Independent measurements of sustained GPU-limited load on M1 Max disagree (roughly 40-60 W attributed to the GPU within a package total near 90 W), so no single figure is asserted" }, - "videoCodecs": { - "encoders": [ - { - "codec": "H.264", - "maxResolution": "4K", - "maxBitrate": 240, - "maxFPS": 60 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "4K", - "maxBitrate": 240, - "maxFPS": 60 - } - ], - "decoders": [ - { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 480, - "maxFPS": 30 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 480, - "maxFPS": 30 - }, - { - "codec": "VP9", - "maxResolution": "4K", - "maxBitrate": 240, - "maxFPS": 60 - } - ] + { + "Name": "Idle", + "Description": "Low-power state with the GPU clock-gated when no Metal work is in flight" } - }, - "PowerEfficiency": { - "BaseClock": 1000, - "BoostClock": 1296, - "MaxTDP": 60, - "PowerStates": [ - { - "Name": "Performance", - "Description": "Maximum performance for demanding workloads" - }, - { - "Name": "Balanced", - "Description": "Balance between performance and power efficiency" - }, - { - "Name": "Power Saving", - "Description": "Optimized for battery life" - } - ], - "ClockGating": true, - "DynamicVoltageFrequencyScaling": true - }, - "PciExpress": { - "Version": "N/A", - "Lanes": 0, - "Bandwidth": 0 - }, - "MultiGpuSupport": { - "Technologies": ["None"], - "MaxGpus": 1, - "InterconnectBandwidth": 0 - } - } \ No newline at end of file + ], + "ClockGating": true, + "DynamicVoltageFrequencyScaling": true + }, + "MultiGpuSupport": { + "Technologies": ["Other"], + "MaxGpus": 2, + "InterconnectBandwidth": 2500 + } +} diff --git a/device_lib/apple-gpu-m2.json b/device_lib/apple-gpu-m2.json index c4a3f016..f63a4b03 100644 --- a/device_lib/apple-gpu-m2.json +++ b/device_lib/apple-gpu-m2.json @@ -1,164 +1,164 @@ { - "Name": "Apple GPU (M2)", - "Vendor": "Apple", - "Architecture": "Apple GPU Gen 2", - "ReleaseYear": 2022, - "FabricationProcess": { - "ProcessNode": 5, - "Manufacturer": "TSMC", - "Technology": "N5P" + "Name": "Apple GPU (M2)", + "Vendor": "Apple", + "Architecture": "apple8", + "ReleaseYear": 2022, + "FabricationProcess": { + "ProcessNode": 5, + "Manufacturer": "TSMC", + "Technology": "N5P" + }, + "CoreSubsystem": { + "Name": "GPU Core", + "ChipType": "GPU", + "CoreType": "GPU Core", + "UnitTypes": { + "GPU": { + "Count": 1, + "Size": 10, + "Description": "Integrated GPU block of the Apple M2 SoC. Apple publishes up to 10 GPU cores (8-core bins also ship) and 100 GB/s of unified memory bandwidth. There is no dedicated VRAM: the GPU reads and writes the same LPDDR5 unified memory as the CPU, the 16-core Neural Engine (15.8 trillion operations per second) and the media engine, all of which are separate blocks on the same die. Apple reports 20 billion transistors for the whole M2 SoC. The media engine listed under SpecializedHardware is a separate fixed-function block, not part of the GPU cores; its \"Other\" codec entries are ProRes and ProRes RAW, which have a dedicated encode and decode engine. AV1 decode is absent on M2 - Apple states it arrived with the M3 media engine for the first time - and the VP9 decode entry is third-party verified rather than Apple-documented", + "Memory": ["L2 Cache (GPU)", "System Level Cache", "Unified Memory (LPDDR5)"], + "Subunits": ["GPU Core"] + }, + "GPU Core": { + "Count": 10, + "Size": 128, + "Description": "Apple GPU core with 128 FP32 ALUs. Apple does not publish per-core internals; the 128-ALU figure is corroborated by the widely cited 3.6 TFLOPS FP32 rating, since 2 x 128 x 10 cores x 1.398 GHz = 3.58 TFLOPS. Apple8 executes FP16 and FP32 at the same peak rate, and the GPU core contains no matrix or tensor units", + "Memory": ["Register File", "Threadgroup Memory", "L1 Data Cache", "L1 Instruction Cache"], + "Subunits": ["SIMD Pipeline"] + }, + "SIMD Pipeline": { + "Count": 4, + "Size": 32, + "Description": "32-lane SIMD execution pipeline. The Metal SIMD-group width (threadExecutionWidth) on Apple GPUs is 32 threads, and each of the core's four schedulers issues one instruction per cycle for one SIMD-group. The 4 x 32 decomposition of the 128 ALUs is not published by Apple; it comes from public reverse-engineering of the Apple G13/G14 GPU ISA" + } + } + }, + "MemorySubsystem": { + "SupportedMemoryTypes": ["LPDDR5"], + "MemoryTypes": [ + { + "Type": "Register File", + "Size": 208, + "description": "Approximately 208 KB of registers per GPU core across the Apple7/Apple8 generations. Not published by Apple; measured by public microbenchmarking. Apple's register file is large relative to its very small L1 and L2, and occupancy on Apple8 ranges from 384 to 3072 resident threads per core depending on register pressure" + }, + { + "Type": "Threadgroup Memory", + "Size": 32, + "description": "Metal caps a single threadgroup at 32 KB of threadgroup memory on the Apple8 family, with explicit imageblock (tile) allocation also capped at 32 KB per threadgroup and implicit imageblock allocation at 128 KB (Metal Feature Set Tables). Maximum threads per threadgroup is 1024 and threadgroup memory length alignment is 16 B. Physical per-core threadgroup memory is larger than the per-threadgroup cap - public microbenchmarking measures roughly 60 KB per core - which allows several threadgroups to be resident at once" + }, + { + "Type": "L1 Data Cache", + "Size": 8, + "description": "8 KB L1 data cache per GPU core with a 128-byte global cache line. Not published by Apple; measured by public microbenchmarking. Deliberately small compared with contemporary AMD and NVIDIA cores, offset by a large register file and a very low-latency L2" + }, + { + "Type": "L1 Instruction Cache", + "Size": 12, + "description": "12 KB L1 instruction cache per GPU core. Not published by Apple; measured by public microbenchmarking" + }, + { + "Type": "L2 Cache (GPU)", + "Size": 1536, + "description": "GPU-wide L2 cache, approximately 1.5 MB on the base M2. Apple publishes no GPU cache sizes. This figure rests on a single public microbenchmarking estimate and is explicitly marked approximate by its source; independent measurements of the M1 and of the M2 Pro die report different capacities, so treat it as indicative only. Apple L2 capacity is known to vary non-monotonically across a family" + }, + { + "Type": "System Level Cache", + "Size": 8192, + "description": "8 MB SoC-wide system level cache shared by the CPU, GPU, Neural Engine and media engine, backing the GPU L2 in front of LPDDR5. Not published by Apple; measured by public microbenchmarking" + }, + { + "Type": "Unified Memory (LPDDR5)", + "Size": 25165824, + "MaxMemoryBandwidth": 100, + "maxBusWidth": 128, + "description": "LPDDR5 unified memory shared with the CPU, Neural Engine and media engine; there is no separate VRAM and no copy across a bus. Apple publishes 100 GB/s of unified memory bandwidth and 8 GB, 16 GB and 24 GB configurations; the Size given here is the maximum 24 GB configuration in KB. Apple does not publish the bus width or LPDDR5 speed grade; public die-shot analysis identifies a 128-bit LPDDR5 bus at 6400 MT/s, which yields 102.4 GB/s and is consistent with Apple's rounded 100 GB/s. Per-package and per-channel counts for the base M2 are not publicly established and are therefore omitted" + } + ] + }, + "KernelModel": { + "LLVMTarget": "air64", + "LLVMTriple": "air64-apple-macosx", + "LLVMFeatures": [ + { + "Name": "air64", + "Description": "Apple Intermediate Representation (AIR), the LLVM-bitcode-based IR that Apple's Metal compiler emits into a .metallib. The triple air64-apple-macosx is what Apple's toolchain uses (air64-apple-ios / air64-apple-tvos for the other platforms), but air64 is not a registered upstream LLVM target and Apple publishes no GPU ISA documentation" + } + ], + "SubUnits": [ + { + "Type": "Metal", + "Version": "3" + }, + { + "Type": "Metal", + "Version": "4" + }, + { + "Type": "OpenCL", + "Version": "1.2" + }, + { + "Type": "OpenGL", + "Version": "4.1" + } + ] + }, + "SpecializedHardware": { + "RayTracingAccelerators": { + "Present": false }, - "CoreSubsystem": { - "Name": "GPU Core", - "subUnits": [ - { - "Type": "Execution Unit", - "Count": 10, - "Size": 128, - "Description": "Enhanced unified shader processors with 128 ALUs per core", - "Memory": ["L1", "L2"], - "SubunitType": "SIMD" - }, - { - "Type": "Tile Memory", - "Count": 10, - "Size": 32, - "Description": "On-chip tile memory for efficient rendering", - "Memory": ["Tile"], - "SubunitType": "Memory" - } - ] + "AiAccelerators": { + "Present": false }, - "MemorySubsystem": { - "SupportedMemoryTypes": ["LPDDR5"], - "MemoryTypes": [ + "videoCodecs": { + "encoders": [ { - "Type": "L1", - "Size": 64, - "BankCount": 10, - "MaxMemoryBandwidth": 100, - "maxBusWidth": 128, - "description": "L1 cache per GPU core" + "codec": "H.264" }, { - "Type": "L2", - "Size": 4096, - "BankCount": 1, - "MaxMemoryBandwidth": 100, - "maxBusWidth": 128, - "description": "Shared L2 cache" + "codec": "H.265/HEVC" }, { - "Type": "Unified Memory", - "Size": 25165824, - "BankCount": 6, - "MaxMemoryBandwidth": 100, - "maxBusWidth": 128, - "description": "Unified memory shared with CPU" + "codec": "Other" } - ] - }, - "KernelModel": { - "SubUnits": [ + ], + "decoders": [ { - "Type": "Metal", - "Version": "3.0" + "codec": "H.264", + "maxResolution": "8K" }, { - "Type": "OpenGL ES", - "Version": "3.0" + "codec": "H.265/HEVC", + "maxResolution": "8K" }, { - "Type": "Vulkan", - "Version": "1.3" - } - ] - }, - "SpecializedHardware": { - "RayTracingAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "Performance": { - "RaysPerSecond": 0 - } - }, - "AiAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "SupportedPrecisions": [], - "Performance": { - "fp16TopsPerGPU": 0, - "int8TopsPerGPU": 0 - } - }, - "videoCodecs": { - "encoders": [ - { - "codec": "H.264", - "maxResolution": "4K", - "maxBitrate": 180, - "maxFPS": 60 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "4K", - "maxBitrate": 180, - "maxFPS": 60 - } - ], - "decoders": [ - { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 360, - "maxFPS": 30 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 360, - "maxFPS": 30 - }, - { - "codec": "VP9", - "maxResolution": "4K", - "maxBitrate": 180, - "maxFPS": 60 - } - ] - } - }, - "PowerEfficiency": { - "BaseClock": 1000, - "BoostClock": 1398, - "MaxTDP": 20, - "PowerStates": [ - { - "Name": "Performance", - "Description": "Maximum performance for demanding workloads" + "codec": "Other", + "maxResolution": "8K" }, { - "Name": "Balanced", - "Description": "Balance between performance and power efficiency" - }, - { - "Name": "Power Saving", - "Description": "Optimized for battery life" + "codec": "VP9" } - ], - "ClockGating": true, - "DynamicVoltageFrequencyScaling": true - }, - "PciExpress": { - "Version": "N/A", - "Lanes": 0, - "Bandwidth": 0 - }, - "MultiGpuSupport": { - "Technologies": ["None"], - "MaxGpus": 1, - "InterconnectBandwidth": 0 + ] } + }, + "PowerEfficiency": { + "BoostClock": 1398, + "PowerStates": [ + { + "Name": "Active", + "Description": "GPU active and clocking up to its maximum. Apple publishes neither GPU clocks nor a TDP for the M2, so no MaxTDP is recorded here. The 1398 MHz maximum is a third-party figure with two independent backers; one die-shot analysis instead reports 1406 MHz, and that discrepancy is unresolved. The minimum GPU clock for the base M2 is not publicly established (the widely quoted 444 MHz floor is the M2 Pro), so no BaseClock is recorded. Third-party measurement of a MacBook Pro 13-inch M2 gives roughly 13.5 W package power under a GPU-only load and about 35 W peak under a combined CPU+GPU load, settling to a sustained 28-30 W; sustained figures depend on the enclosure, since the MacBook Air M2 is fanless" + }, + { + "Name": "Idle", + "Description": "Reduced power state when no GPU work is queued" + } + ], + "ClockGating": true, + "DynamicVoltageFrequencyScaling": true + }, + "MultiGpuSupport": { + "Technologies": ["None"], + "MaxGpus": 1 } +} diff --git a/device_lib/apple-gpu-m3.json b/device_lib/apple-gpu-m3.json index e17efb4b..265fcbe8 100644 --- a/device_lib/apple-gpu-m3.json +++ b/device_lib/apple-gpu-m3.json @@ -1,184 +1,131 @@ { - "Name": "Apple GPU (M3)", - "Vendor": "Apple", - "Architecture": "Apple GPU Gen 3", - "ReleaseYear": 2023, - "FabricationProcess": { - "ProcessNode": 3, - "Manufacturer": "TSMC", - "Technology": "N3B" + "Name": "Apple GPU (M3)", + "Vendor": "Apple", + "Architecture": "Apple9", + "ReleaseYear": 2023, + "FabricationProcess": { + "ProcessNode": 3, + "Manufacturer": "TSMC", + "Technology": "N3B" + }, + "CoreSubsystem": { + "Name": "GPU Core", + "ChipType": "GPU", + "CoreType": "GPU Core", + "UnitTypes": { + "GPU": { + "Count": 1, + "Size": 10, + "Description": "Integrated GPU block of the Apple M3 system on chip (25 billion transistors, 3 nm). Apple ships M3 with a 10-core GPU; an 8-core binned variant ships in the base MacBook Air and iMac. Metal reports it as MTLGPUFamily.apple9. It is a tile-based deferred renderer whose family 9 shader core introduces Dynamic Caching, hardware-accelerated ray tracing and hardware-accelerated mesh shading, all new to the Mac with M3. The GPU shares LPDDR5 unified memory with the 8-core CPU, the 16-core Neural Engine and the media engine (H.264, HEVC, ProRes and ProRes RAW encode and decode, plus AV1 decode)", + "Memory": ["Unified Memory (LPDDR5)"], + "Subunits": ["GPU Core"] + }, + "GPU Core": { + "Count": 10, + "Size": 128, + "Description": "Apple family 9 \"next-generation shader core\". Each core runs shader, compute, texture and tile work and owns the on-chip memory pool that Dynamic Caching allocates from. Apple does not publish a per-core ALU count for M3: Size 128 is inferred from Apple's own M1 figures (8 GPU cores, 1024 ALUs, 2.6 TFLOPS FP32) and is consistent with third-party FP32 throughput measurements of the 10-core M3 GPU", + "Memory": ["Dynamic Shader Core Memory", "Threadgroup Memory (per threadgroup)", "Tile Memory (per threadgroup)"], + "Subunits": ["SIMD Unit", "Hardware Intersector"] + }, + "SIMD Unit": { + "Count": 4, + "Size": 32, + "Description": "32-wide SIMD execution unit. The width of 32 is Apple's documented SIMD-group size, reported by Metal as threadExecutionWidth on Apple GPUs. Apple has never published the internal issue organization of a GPU core: 4 units of 32 lanes is a modeling choice that reproduces the 128 ALUs per core implied by Apple's M1 figures, which describe the same silicon as 16 execution units of 8 ALUs per core. Neither breakdown has been restated by Apple for M3" + }, + "Hardware Intersector": { + "Count": 1, + "Size": 1, + "Description": "Fixed-function ray tracing unit introduced with Apple family 9 and called the \"hardware intersector\" by Apple. It traverses the acceleration structure independently of the shader cores and regroups intersection-function invocations into coherent SIMD-groups to remove divergence overhead; custom intersection functions still execute as Metal Shading Language code on the shader core. Apple publishes neither the number of intersectors per GPU core nor a ray throughput figure, so this is modeled as one unit per GPU core" + } + } + }, + "MemorySubsystem": { + "SupportedMemoryTypes": ["LPDDR5"], + "MemoryTypes": [ + { + "Type": "Unified Memory (LPDDR5)", + "Size": 25165824, + "MaxMemoryBandwidth": 100, + "maxBusWidth": 128, + "description": "LPDDR5 unified memory shared by the GPU, CPU, Neural Engine and media engine. Apple publishes 100 GB/s of memory bandwidth for M3 and ships 8 GB, 16 GB and 24 GB configurations; Size records the 24 GB maximum. The 128-bit bus width is third-party reported and not published by Apple. There is no separate VRAM: Metal buffers and textures are allocated from this pool and are addressable by CPU and GPU without copies" + }, + { + "Type": "Dynamic Shader Core Memory", + "description": "Unified on-chip memory pool of the Apple family 9 shader core. Dynamic Caching allocates registers, threadgroup memory, tile memory, stack space and buffer data from the same set of larger caches on demand over the lifetime of a shader, instead of reserving a fixed worst-case allocation per shader, which raises occupancy for register-heavy kernels. Apple publishes neither the capacity of this pool nor the sizes of the GPU L1 and L2 caches, so no Size is recorded" + }, + { + "Type": "Threadgroup Memory (per threadgroup)", + "Size": 32, + "description": "Maximum threadgroup (shared) memory one threadgroup may allocate, from Apple's Metal Feature Set Tables for the Apple9 family, with a 16-byte length alignment and a maximum of 1024 threads per threadgroup. This is an API limit rather than the physical capacity of a shader core; on M3 the storage is carved out of the shared on-chip pool by Dynamic Caching" + }, + { + "Type": "Tile Memory (per threadgroup)", + "Size": 32, + "description": "Maximum imageblock/tile memory per threadgroup for the Apple9 family per Apple's Metal Feature Set Tables (with a maximum of 8 KB per thread). The tile-based deferred renderer keeps render-pass attachments and meshlet data here to avoid memory traffic; on M3 it is also dynamically allocated by Dynamic Caching" + } + ] + }, + "KernelModel": { + "LLVMTarget": "air64", + "LLVMTriple": "air64-apple-macos14.0", + "SubUnits": [ + { + "Type": "Metal", + "Version": "3.1" + }, + { + "Type": "Metal Shading Language", + "Version": "3.1" + } + ] + }, + "SpecializedHardware": { + "RayTracingAccelerators": { + "Present": true, + "Name": "Hardware Intersector" }, - "CoreSubsystem": { - "Name": "GPU Core", - "subUnits": [ - { - "Type": "Execution Unit", - "Count": 10, - "Size": 128, - "Description": "Next-generation unified shader processors with hardware ray tracing support", - "Memory": ["L1", "L2"], - "SubunitType": "SIMD" - }, - { - "Type": "Ray Tracing Unit", - "Count": 10, - "Size": 1, - "Description": "Dedicated hardware ray tracing acceleration units", - "Memory": ["L1"], - "SubunitType": "RT" - }, - { - "Type": "Tile Memory", - "Count": 10, - "Size": 32, - "Description": "On-chip tile memory for efficient rendering", - "Memory": ["Tile"], - "SubunitType": "Memory" - } - ] + "AiAccelerators": { + "Present": false }, - "MemorySubsystem": { - "SupportedMemoryTypes": ["LPDDR5"], - "MemoryTypes": [ - { - "Type": "L1", - "Size": 64, - "BankCount": 10, - "MaxMemoryBandwidth": 100, - "maxBusWidth": 128, - "description": "L1 cache per GPU core" - }, + "videoCodecs": { + "encoders": [ { - "Type": "L2", - "Size": 4096, - "BankCount": 1, - "MaxMemoryBandwidth": 100, - "maxBusWidth": 128, - "description": "Shared L2 cache" + "codec": "H.264" }, { - "Type": "Unified Memory", - "Size": 25165824, - "BankCount": 6, - "MaxMemoryBandwidth": 100, - "maxBusWidth": 128, - "description": "Unified memory shared with CPU" + "codec": "H.265/HEVC" } - ] - }, - "KernelModel": { - "SubUnits": [ + ], + "decoders": [ { - "Type": "Metal", - "Version": "3.1" + "codec": "H.264" }, { - "Type": "OpenGL ES", - "Version": "3.0" + "codec": "H.265/HEVC" }, { - "Type": "Vulkan", - "Version": "1.3" + "codec": "AV1" } ] - }, - "SpecializedHardware": { - "RayTracingAccelerators": { - "Present": true, - "Name": "RT Unit", - "Count": 1, - "Performance": { - "RaysPerSecond": 2.5 - } - }, - "AiAccelerators": { - "Present": false, - "Name": null, - "Count": 0, - "SupportedPrecisions": [], - "Performance": { - "fp16TopsPerGPU": 0, - "int8TopsPerGPU": 0 - } + } + }, + "PowerEfficiency": { + "PowerStates": [ + { + "Name": "Active", + "Description": "Full power state with all GPU cores available. Apple publishes no GPU base or boost clock and no GPU or package TDP for M3, so no clock or TDP values are recorded" }, - "videoCodecs": { - "encoders": [ - { - "codec": "H.264", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 - }, - { - "codec": "AV1", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 - } - ], - "decoders": [ - { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 400, - "maxFPS": 30 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 400, - "maxFPS": 30 - }, - { - "codec": "VP9", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 - }, - { - "codec": "AV1", - "maxResolution": "4K", - "maxBitrate": 200, - "maxFPS": 60 - } - ] + { + "Name": "Idle", + "Description": "Reduced power state when no graphics or compute work is scheduled" } - }, - "PowerEfficiency": { - "BaseClock": 1000, - "BoostClock": 1424, - "MaxTDP": 22, - "PowerStates": [ - { - "Name": "Performance", - "Description": "Maximum performance for demanding workloads" - }, - { - "Name": "Balanced", - "Description": "Balance between performance and power efficiency" - }, - { - "Name": "Power Saving", - "Description": "Optimized for battery life" - } - ], - "ClockGating": true, - "DynamicVoltageFrequencyScaling": true - }, - "PciExpress": { - "Version": "N/A", - "Lanes": 0, - "Bandwidth": 0 - }, - "MultiGpuSupport": { - "Technologies": ["None"], - "MaxGpus": 1, - "InterconnectBandwidth": 0 - } - } \ No newline at end of file + ], + "ClockGating": true, + "DynamicVoltageFrequencyScaling": true + }, + "MultiGpuSupport": { + "Technologies": ["None"], + "MaxGpus": 1, + "InterconnectBandwidth": 0 + } +} diff --git a/device_lib/apple-gpu-m4.json b/device_lib/apple-gpu-m4.json index 4f8112ab..a3b5ebf1 100644 --- a/device_lib/apple-gpu-m4.json +++ b/device_lib/apple-gpu-m4.json @@ -1,213 +1,133 @@ { - "$schema": "http://json-schema.org/draft-07/schema#", - "Name": "Apple M4 GPU", - "Vendor": "Apple", - "Architecture": "Apple GPU", - "ReleaseYear": 2024, - "FabricationProcess": { - "ProcessNode": 3, - "Manufacturer": "TSMC", - "Technology": "N3E" - }, - "CoreSubsystem": { - "Name": "GPU Core", - "subUnits": [ - { - "Type": "Cores", - "Count": 10, - "Size": 160, - "Description": "Compute Cores per Chip", - "Memory": ["L2"], - "SubunitType": "Execution Unit" - }, - { - "Type": "Execution Unit", - "Count": 160, - "Size": 128, - "Description": "SIMD execution units for parallel shader processing with 128 ALUs per unit", - "Memory": ["L1", "vector register file", "shared memory"], - "SubunitType": "SIMD" - }, - { - "Type": "Texture Unit", - "Count": 10, - "Size": 4, - "Description": "Texture sampling and filtering units", - "Memory": ["L1"], - "SubunitType": "Texture" - } - ] - }, - "MemorySubsystem": { - "SupportedMemoryTypes": ["LPDDR5"], - "MemoryTypes": [ - { - "Type": "vector register file", - "Size": 8, - "BankCount": 32, - "MaxMemoryBandwidth": 2048, - "maxBusWidth": 512, - "description": "High-speed vector register file for shader operations" - }, - { - "Type": "shared memory", - "Size": 32, - "BankCount": 16, - "MaxMemoryBandwidth": 1024, - "maxBusWidth": 256, - "description": "Shared memory for threadgroup communication and data sharing" - }, - { - "Type": "L1", - "Size": 64, - "BankCount": 8, - "MaxMemoryBandwidth": 512, - "maxBusWidth": 128, - "description": "L1 cache for immediate GPU operations and texture data" - }, - { - "Type": "L2", - "Size": 16384, - "BankCount": 16, - "MaxMemoryBandwidth": 256, - "maxBusWidth": 256, - "description": "L2 cache shared across GPU cores for larger data sets" - }, - { - "Type": "Unified Memory", - "Size": 16777216, - "BankCount": 8, - "MaxMemoryBandwidth": 120, - "maxBusWidth": 256, - "description": "Shared system memory accessible by CPU and GPU with zero-copy access" - } - ] - }, - "KernelModel": { - "SubUnits": [ - { - "Type": "metal", - "Version": "3.1" - }, - { - "Type": "opencl", - "Version": "1.2" - }, - { - "Type": "vulkan", - "Version": "1.3" - }, - { - "Type": "compute", - "Version": "6.0" - } - ] - }, - "SpecializedHardware": { - "RayTracingAccelerators": { - "Present": true, - "Name": "RT Unit", + "Name": "Apple M4 GPU", + "Vendor": "Apple", + "Architecture": "Apple9", + "ReleaseYear": 2024, + "FabricationProcess": { + "ProcessNode": 3, + "Manufacturer": "TSMC", + "Technology": "N3E" + }, + "CoreSubsystem": { + "Name": "GPU Core", + "ChipType": "GPU", + "CoreType": "GPU Core", + "UnitTypes": { + "GPU": { "Count": 1, - "Performance": { - "RaysPerSecond": 2.5 - } + "Size": 10, + "Description": "Integrated GPU block of the Apple M4 SoC (28 billion transistors total). Apple ships a 10-core GPU in Mac mini, iMac, MacBook Pro and most iPad Pro configurations; the entry-level 256 GB/512 GB 11-inch iPad Pro uses a binned 9-core part. The GPU reports Metal GPU family Apple9 and adds Dynamic Caching (on-chip memory allocated to registers, threadgroup memory and stack in hardware at run time), hardware-accelerated ray tracing and hardware-accelerated mesh shading. It shares unified memory with the CPU and Neural Engine at 120 GB/s", + "Memory": ["Unified Memory"], + "Subunits": ["GPU Core"] }, - "AiAccelerators": { - "Present": true, - "Name": "Neural Engine", - "Count": 1, - "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8"], - "Performance": { - "fp16TopsPerGPU": 38, - "int8TopsPerGPU": 76 - } + "GPU Core": { + "Count": 10, + "Size": 4, + "Description": "Apple GPU core: 128 FP32 ALUs organized as 4 SIMD pipelines of 32 lanes, giving 1280 ALUs across the 10-core GPU. The 32-thread SIMD-group width is documented by Apple (Metal Feature Set Tables, Apple9); the 128-ALU-per-core figure follows Apple's published M-series GPU ratios and is corroborated by independent microarchitecture analysis. Apple does not publish per-core register file, L1 or L2 cache capacities for the M3/M4 generation, and Dynamic Caching makes the register/threadgroup split dynamic rather than fixed. Each core also provides the hardware ray-tracing and mesh-shading blocks; Apple does not publish their per-core counts", + "Memory": ["Threadgroup Memory"], + "Subunits": ["SIMD Pipeline"] + }, + "SIMD Pipeline": { + "Count": 4, + "Size": 32, + "Description": "32-wide SIMD execution pipeline. Each pipeline issues one instruction per cycle for one 32-thread SIMD-group, and 4 pipelines per core account for the core's 128 ALUs. The SIMD-group width of 32 and the 1024-thread maximum threadgroup size are Apple-documented Metal limits; the 4-pipeline decomposition is inferred from the ALU count and community microarchitecture analysis, as Apple does not document the dispatch structure" + } + } + }, + "MemorySubsystem": { + "SupportedMemoryTypes": ["LPDDR5"], + "MemoryTypes": [ + { + "Type": "Threadgroup Memory", + "Size": 32, + "description": "Maximum threadgroup (shared) memory allocation per threadgroup is 32 KB for the Apple9 GPU family (Apple Metal Feature Set Tables); maximum threads per threadgroup is 1024 and maximum buffer length is 2 GB. On M3-generation and later GPUs this storage is drawn from the Dynamic Caching on-chip pool that also backs registers and stack, so the physical pool size per core is not a published constant" + }, + { + "Type": "Unified Memory", + "Size": 33554432, + "MaxMemoryBandwidth": 120, + "maxBusWidth": 128, + "description": "LPDDR5X unified memory shared by CPU, GPU, Neural Engine and media engine with zero-copy access. Apple publishes 120 GB/s of memory bandwidth and configurations of 16 GB, 24 GB or 32 GB; the 32 GB maximum is recorded here. The 128-bit bus width is the arithmetic match for 120 GB/s at LPDDR5X-7500 (Apple does not state bus width directly). LPDDR5X is reported as LPDDR5 above because the schema enum has no LPDDR5X value. Apple publishes no GPU L1/L2 or system-level cache sizes for M4, so no cache levels are listed" + } + ] + }, + "KernelModel": { + "LLVMTarget": "air64", + "LLVMTriple": "air64-apple-macos", + "LLVMFeatures": [ + { + "Name": "air64", + "Description": "Apple Intermediate Representation (AIR), the LLVM-bitcode-based IR emitted by Apple's metal compiler (-target air64-apple-macos, or air64-apple-ios for iPadOS). AIR is Apple-private and is not an upstream LLVM backend target; there is no public LLVM target for the Apple GPU ISA" + } + ], + "SubUnits": [ + { + "Type": "metal", + "Version": "4" + }, + { + "Type": "metal-shading-language", + "Version": "4.1" }, - "videoCodecs": { - "encoders": [ - { - "codec": "H.264", - "maxResolution": "4K", - "maxBitrate": 100, - "maxFPS": 60 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "4K", - "maxBitrate": 100, - "maxFPS": 60 - }, - { - "codec": "AV1", - "maxResolution": "4K", - "maxBitrate": 100, - "maxFPS": 60 - } - ], - "decoders": [ - { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 200, - "maxFPS": 60 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 200, - "maxFPS": 60 - }, - { - "codec": "AV1", - "maxResolution": "8K", - "maxBitrate": 200, - "maxFPS": 60 - }, - { - "codec": "VP9", - "maxResolution": "4K", - "maxBitrate": 150, - "maxFPS": 60 - } - ] + { + "Type": "opencl", + "Version": "1.2" + } + ] + }, + "SpecializedHardware": { + "RayTracingAccelerators": { + "Present": true, + "Name": "Ray-tracing engine" + }, + "AiAccelerators": { + "Present": true, + "Name": "16-core Neural Engine", + "Count": 16, + "SupportedPrecisions": ["FP16", "INT8"], + "Performance": { + "int8TopsPerGPU": 38 } }, - "PowerEfficiency": { - "BaseClock": 1398, - "BoostClock": 1598, - "MaxTDP": 22, - "PowerStates": [ + "videoCodecs": { + "encoders": [ { - "Name": "P0", - "Description": "Maximum performance state for demanding GPU workloads" + "codec": "H.264" }, { - "Name": "P1", - "Description": "High performance state with slight power reduction" - }, + "codec": "H.265/HEVC" + } + ], + "decoders": [ { - "Name": "P2", - "Description": "Balanced performance and power state" + "codec": "H.264" }, { - "Name": "P3", - "Description": "Low power state for light GPU tasks" + "codec": "H.265/HEVC" }, { - "Name": "P4", - "Description": "Idle state with minimal power consumption" + "codec": "AV1" } - ], - "ClockGating": true, - "DynamicVoltageFrequencyScaling": true - }, - "PciExpress": { - "Version": "N/A", - "Lanes": 0, - "Bandwidth": 0 - }, - "MultiGpuSupport": { - "Technologies": ["None"], - "MaxGpus": 1, - "InterconnectBandwidth": 0 + ] } - } \ No newline at end of file + }, + "PowerEfficiency": { + "PowerStates": [ + { + "Name": "Active", + "Description": "GPU active. Apple publishes neither GPU clock frequencies nor a chip- or GPU-level TDP for M4, so BaseClock, BoostClock and MaxTDP are omitted rather than estimated; third-party measurements of the M4 GPU peak clock disagree (1.68-1.80 GHz)" + }, + { + "Name": "Idle", + "Description": "Low-power state when no GPU work is scheduled; frequency and voltage are managed by the SoC power controller" + } + ], + "ClockGating": true, + "DynamicVoltageFrequencyScaling": true + }, + "MultiGpuSupport": { + "Technologies": ["None"], + "MaxGpus": 1, + "InterconnectBandwidth": 0 + } +} diff --git a/device_lib/apple-gpu-m4p.json b/device_lib/apple-gpu-m4p.json deleted file mode 100644 index c03fc57a..00000000 --- a/device_lib/apple-gpu-m4p.json +++ /dev/null @@ -1,252 +0,0 @@ -{ - "$schema": "http://json-schema.org/draft-07/schema#", - "Name": "Apple M4 Pro GPU", - "Vendor": "Apple", - "Architecture": "apple-m4", - "ReleaseYear": 2024, - "FabricationProcess": { - "ProcessNode": 3, - "Manufacturer": "TSMC", - "Technology": "N3E" - }, - "CoreSubsystem": { - "Name": "GPU Core Pro", - "subUnits": [ - { - "Type": "Device", - "Size": 20, - "Description": "Compute Cores per Chip", - "Memory": ["L2"], - "SubunitType": "Core" - }, - { - "Type": "Core", - "Size": 80, - "Description": "Compute Core", - "Memory": ["L1", "vector register file", "scalar register file", "shared memory"], - "SubunitType": "Execution Unit" - }, - { - "Type": "Execution Unit", - "Size": 32, - "Description": "SIMD execution units for parallel shader processing with 128 ALUs per unit", - "SubunitType": "SIMD" - }, - { - "Type": "Texture Unit", - "Size": 4, - "Description": "Enhanced texture sampling and filtering units with higher throughput", - "Memory": ["L1"], - "SubunitType": "Texture" - } - ] - }, - "MemorySubsystem": { - "SupportedMemoryTypes": ["LPDDR5"], - "MemoryTypes": [ - { - "Type": "scalar register file", - "Size": 4, - "BankCount": 16, - "MaxMemoryBandwidth": 1024, - "maxBusWidth": 256, - "description": "Scalar register file for control flow and address computation" - }, - { - "Type": "vector register file", - "Size": 12, - "BankCount": 64, - "MaxMemoryBandwidth": 4096, - "maxBusWidth": 512, - "description": "Enhanced vector register file for shader operations with increased capacity" - }, - { - "Type": "shared memory", - "Size": 32, - "BankCount": 32, - "MaxMemoryBandwidth": 2048, - "maxBusWidth": 512, - "description": "Expanded shared memory for threadgroup communication and data sharing" - }, - { - "Type": "L1", - "Size": 128, - "BankCount": 16, - "MaxMemoryBandwidth": 1024, - "maxBusWidth": 256, - "description": "Enhanced L1 cache for immediate GPU operations and texture data" - }, - { - "Type": "L2", - "Size": 32768, - "BankCount": 32, - "MaxMemoryBandwidth": 512, - "maxBusWidth": 512, - "description": "Larger L2 cache shared across GPU cores for enhanced performance" - }, - { - "Type": "Unified Memory", - "Size": 25165824, - "BankCount": 12, - "MaxMemoryBandwidth": 273, - "maxBusWidth": 384, - "description": "Enhanced unified memory with higher bandwidth for CPU-GPU data sharing" - } - ] - }, - "KernelModel": { - "LLVMTarget": "apple-m4", - "LLVMTriple": "air64-apple-macosx14.0.0", - "LLVMFeatures": [ - { - "Name": "gpu-family-apple9", - "Description": "Apple GPU Family 9 feature set for M4 series" - }, - { - "Name": "air-bfloat", - "Description": "BFloat16 support in Apple Intermediate Representation" - }, - { - "Name": "air-pack-struct", - "Description": "Structure packing optimization in AIR" - }, - { - "Name": "air-simd-group-matrix", - "Description": "SIMD group matrix operations support" - }, - { - "Name": "air-mesh-shader", - "Description": "Mesh shader support in Metal" - } - ], - "SubUnits": [ - { - "Type": "metal", - "Version": "3.1" - }, - { - "Type": "opencl", - "Version": "1.2" - }, - { - "Type": "vulkan", - "Version": "1.3" - }, - { - "Type": "compute", - "Version": "6.0" - } - ] - }, - "SpecializedHardware": { - "RayTracingAccelerators": { - "Present": true, - "Name": "RT Unit Pro", - "Count": 2, - "Performance": { - "RaysPerSecond": 5.2 - } - }, - "AiAccelerators": { - "Present": true, - "Name": "Neural Engine", - "Count": 1, - "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "INT4"], - "Performance": { - "fp16TopsPerGPU": 38, - "int8TopsPerGPU": 76 - } - }, - "videoCodecs": { - "encoders": [ - { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 200, - "maxFPS": 60 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 200, - "maxFPS": 60 - }, - { - "codec": "AV1", - "maxResolution": "8K", - "maxBitrate": 200, - "maxFPS": 60 - } - ], - "decoders": [ - { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 400, - "maxFPS": 120 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 400, - "maxFPS": 120 - }, - { - "codec": "AV1", - "maxResolution": "8K", - "maxBitrate": 400, - "maxFPS": 120 - }, - { - "codec": "VP9", - "maxResolution": "8K", - "maxBitrate": 300, - "maxFPS": 60 - } - ] - } - }, - "PowerEfficiency": { - "BaseClock": 1444, - "BoostClock": 1698, - "MaxTDP": 35, - "PowerStates": [ - { - "Name": "P0", - "Description": "Maximum performance state for professional GPU workloads" - }, - { - "Name": "P1", - "Description": "High performance state with optimized power efficiency" - }, - { - "Name": "P2", - "Description": "Balanced performance and power state for sustained workloads" - }, - { - "Name": "P3", - "Description": "Medium power state for moderate GPU tasks" - }, - { - "Name": "P4", - "Description": "Low power state for light GPU operations" - }, - { - "Name": "P5", - "Description": "Idle state with minimal power consumption" - } - ], - "ClockGating": true, - "DynamicVoltageFrequencyScaling": true - }, - "PciExpress": { - "Version": "N/A", - "Lanes": 0, - "Bandwidth": 0 - }, - "MultiGpuSupport": { - "Technologies": ["None"], - "MaxGpus": 1, - "InterconnectBandwidth": 0 - } - } \ No newline at end of file diff --git a/device_lib/aws-npu-trn1.json b/device_lib/aws-npu-trn1.json new file mode 100644 index 00000000..ce0279bf --- /dev/null +++ b/device_lib/aws-npu-trn1.json @@ -0,0 +1,174 @@ +{ + "Name": "Trainium", + "Vendor": "Other", + "Architecture": "trn1", + "ReleaseYear": 2022, + "FabricationProcess": { + "ProcessNode": 7, + "Manufacturer": "TSMC", + "Technology": "N7 (widely reported; AWS does not publish the Trainium process node)" + }, + "CoreSubsystem": { + "Name": "NeuronCore-v2", + "ChipType": "Chip", + "CoreType": "NeuronCore-v2", + "UnitTypes": { + "Chip": { + "Count": 1, + "Size": 2, + "Description": "One Trainium NeuronDevice: 2 NeuronCore-v2, 2 HBM stacks (32 GiB at 820 GB/s), 32 DMA engines, 6 CC-Cores and 4 NeuronLink-v2 links. AWS rates the chip at 190 FP16/BF16/cFP8/TF32 TFLOPS, 47.5 FP32 TFLOPS and 380 INT8 TOPS. A trn1.32xlarge holds 16 chips (32 NeuronCores) wired into a 2D torus over NeuronLink-v2 and delivers up to 3 PFLOPS FP16/BF16; trn1.2xlarge exposes a single chip. Scale-out beyond the instance goes over EFAv2 (800 Gbps on trn1.32xlarge, 1600 Gbps on trn1n.32xlarge)", + "Memory": ["HBM", "SBUF", "PSUM", "TCM"], + "Subunits": ["NeuronCore-v2", "DMA Engine", "CC-Core", "HBM Stack", "NeuronLink-v2 Link", "PCIe Controller"] + }, + "NeuronCore-v2": { + "Count": 2, + "Size": 5, + "Description": "Independent accelerator core with 4 heterogeneous compute engines (Tensor, Vector, Scalar, GpSimd) plus a Sync Engine, all executing concurrently over a shared software-managed state buffer. Cores can be used individually or logically merged (Logical NeuronCore config) so a chip appears as one larger core", + "Memory": ["SBUF", "PSUM", "TCM"], + "Subunits": ["Tensor Engine", "Vector Engine", "Scalar Engine", "GPSIMD Engine", "Sync Engine"] + }, + "Tensor Engine": { + "Count": 1, + "Size": 16384, + "Description": "Systolic array of 128 rows by 128 columns of processing elements per NeuronCore, for matrix multiplication and convolution. Runs at 2.8 GHz, reads operands from SBUF and accumulates into PSUM. Accepts BF16, FP16, TF32 and configurable FP8 (cFP8) inputs at 92 TFLOPS and FP32 inputs at 23 TFLOPS; INT8 inputs produce INT32 results. Accumulation is mixed-precision with FP32 accumulators and FP32 output. On v2 cFP8 runs at the same rate as BF16 - the FP8 rate doubling only arrives with NeuronCore-v3" + }, + "Vector Engine": { + "Count": 1, + "Size": 128, + "Description": "SIMD engine at 1.12 GHz processing 128 elements per cycle, one per SBUF partition, for element-wise and reduction operations (normalization, softmax reductions, accumulation). Operates SBUF-to-SBUF or PSUM-to-SBUF in parallel with the Tensor Engine, at roughly 2.3 FP32 TFLOPS per NeuronCore" + }, + "Scalar Engine": { + "Count": 1, + "Size": 128, + "Description": "128-lane engine at 1.4 GHz for activation functions and transcendentals (GELU, exp, sigmoid, tanh), applied at one element per cycle per lane while the other engines run concurrently, at roughly 2.9 FP32 TFLOPS per NeuronCore" + }, + "GPSIMD Engine": { + "Count": 1, + "Size": 8, + "Description": "8 fully programmable 512-bit wide general-purpose SIMD DSP processors per NeuronCore at 1.4 GHz, together presenting a 128-lane 32-bit data path. Each processor is wired to 16 SBUF partitions and has 64 KB of tightly-coupled memory. Addressable from Neuron Custom C++ Operators, it handles irregular ops (embeddings, gather/scatter, dynamic control flow) that do not map to the fixed-function engines", + "Memory": ["TCM"] + }, + "Sync Engine": { + "Count": 1, + "Size": 1, + "Description": "Engine sequencer inside each NeuronCore that executes the same class of control instructions as the compute-engine sequencers. Its role is to trigger DMA transfers without interfering with compute-engine instruction scheduling and ordering" + }, + "DMA Engine": { + "Count": 32, + "Size": 1, + "Description": "32 descriptor-driven DMA engines per chip for intra- and inter-device data movement. Each processes one transfer at a time at a peak of about 27 GiB/s, for roughly 1 TB/s of aggregate DMA bandwidth with inline memory compression and decompression. They stage tensors between HBM and SBUF/PSUM and overlap data movement with compute" + }, + "CC-Core": { + "Count": 6, + "Size": 1, + "Description": "6 cores per chip dedicated to collective communication (all-reduce, all-gather, reduce-scatter), driving the 4 NeuronLink-v2 links. On NeuronCore-v2 collectives execute on these device-level cores rather than on an engine inside the NeuronCore; NeuronCore-v3 scales the count to 20" + }, + "HBM Stack": { + "Count": 2, + "Size": 1, + "Description": "2 HBM stacks per chip, 16 GiB each, for 32 GiB of device memory at 820 GB/s aggregate. AWS does not state the HBM generation; 820 GB/s across 2 stacks implies about 3.2 Gbps per pin on a 1024-bit stack, which is an HBM2e-class data rate" + }, + "NeuronLink-v2 Link": { + "Count": 4, + "Size": 1, + "Description": "4 NeuronLink-v2 chip-to-chip links per Trainium chip (Inferentia2 has 2), wiring the 16 chips of a trn1.32xlarge into a 2D torus. AWS publishes 'up to 768 GB/s of NeuronLink' for Trn1 instances but does not publish a per-link or per-device-pair figure" + }, + "PCIe Controller": { + "Count": 1, + "Size": 1, + "Description": "Host interface controller attaching the Trainium chip to the EC2 host as a PCIe endpoint. AWS does not publish the PCIe generation or lane count for Trainium" + } + } + }, + "MemorySubsystem": { + "SupportedMemoryTypes": ["HBM2e"], + "MemoryTypes": [ + { + "Type": "SBUF (State Buffer, per NeuronCore)", + "Size": 24576, + "BankCount": 128, + "description": "24 MiB of on-chip software-managed SRAM per NeuronCore, organized as 128 partitions of 192 KiB. Not a cache - no automatic fill, eviction, or coherence. All engines read and write operands here, and data must be explicitly staged from HBM by the DMA engines. BankCount is the partition count. Total 48 MiB per chip across 2 NeuronCores" + }, + { + "Type": "PSUM (Partial Sum Buffer, per NeuronCore)", + "Size": 2048, + "BankCount": 8, + "description": "2 MiB accumulation SRAM per NeuronCore, organized as 128 partitions of 16 KiB, each partition split into 8 banks holding 512 FP32 values (2 KiB) each. Holds FP32 matmul partial sums written by the Tensor Engine, which has exclusive near-memory read-accumulate-write to every 4-byte element, so multi-pass matmuls accumulate in place before being copied back to SBUF. BankCount is banks per partition" + }, + { + "Type": "TCM (GpSimd tightly-coupled memory, per processor)", + "Size": 64, + "description": "64 KB of local data RAM private to each of the 8 GpSimd processors in a NeuronCore, giving 512 KB per NeuronCore. Used as scratch by Neuron Custom C++ Operators" + }, + { + "Type": "HBM (Device DRAM)", + "Size": 33554432, + "BankCount": 2, + "MaxMemoryBandwidth": 820, + "maxBusWidth": 2048, + "description": "32 GiB of device memory per chip across 2 HBM stacks at 820 GB/s. A trn1.32xlarge exposes 512 GB of shared accelerator memory with 9.8 TB/s of total memory bandwidth as published by AWS (16 x 820 GB/s would be 13.1 TB/s; AWS quotes the lower figure). BankCount is the stack count; maxBusWidth assumes the standard 1024-bit interface per HBM stack. AWS does not state the HBM generation, but the implied per-pin data rate is HBM2e class. Moved to and from SBUF explicitly by the DMA engines; there is no hardware cache hierarchy between HBM and SBUF" + } + ] + }, + "KernelModel": { + "SubUnits": [ + { + "Type": "NKI (Neuron Kernel Interface)", + "Version": "0.6.0" + }, + { + "Type": "Neuron Custom C++ Operators (aws-neuronx-gpsimd-customop-lib)", + "Version": "0.24.1.0" + }, + { + "Type": "Neuron Compiler (neuronx-cc)", + "Version": "2.27.5334.0" + }, + { + "Type": "PyTorch NeuronX (torch-neuronx)", + "Version": "2.9.0.2.15.32035" + }, + { + "Type": "AWS Neuron SDK", + "Version": "2.32.0" + } + ] + }, + "SpecializedHardware": { + "RayTracingAccelerators": { + "Present": false + }, + "AiAccelerators": { + "Present": true, + "Name": "NeuronCore-v2 Tensor Engine", + "Count": 2, + "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "Other"], + "Performance": { + "fp16TopsPerGPU": 190, + "int8TopsPerGPU": 380 + } + }, + "videoCodecs": { + "encoders": [], + "decoders": [] + } + }, + "PowerEfficiency": { + "MaxTDP": 350, + "PowerStates": [ + { + "Name": "Active", + "Description": "Full power state with both NeuronCores and HBM stacks active. AWS publishes no per-chip TDP for Trainium and no per-instance power figure; the 350 W here is a third-party estimate, not an AWS number. For scale, the later Trainium2 chip is reported at roughly 500 W" + }, + { + "Name": "Idle", + "Description": "Reduced power state when no workloads are running. AWS does not document the idle power level or the transition policy" + } + ] + }, + "MultiGpuSupport": { + "Technologies": ["Other"], + "MaxGpus": 16, + "InterconnectBandwidth": 768 + } +} diff --git a/device_lib/aws-npu-trn2.json b/device_lib/aws-npu-trn2.json new file mode 100644 index 00000000..78042e27 --- /dev/null +++ b/device_lib/aws-npu-trn2.json @@ -0,0 +1,176 @@ +{ + "Name": "Trainium2", + "Vendor": "Other", + "Architecture": "trn2", + "ReleaseYear": 2024, + "FabricationProcess": { + "ProcessNode": 5, + "Manufacturer": "TSMC", + "Technology": "Not published by AWS; a TSMC 5nm-class node is the commonly reported figure and is recorded here as an inference, not a vendor-stated fact" + }, + "CoreSubsystem": { + "Name": "NeuronCore-v3", + "ChipType": "Chip", + "CoreType": "NeuronCore-v3", + "UnitTypes": { + "Chip": { + "Count": 1, + "Size": 8, + "Description": "One Trainium2 device: eight NeuronCore-v3 cores, 96 GiB of device memory across four 24 GB HBM banks at 2.9 TB/s, 16 CC-Cores and 3.5 TB/s of DMA bandwidth with inline memory compression. Collectively the cores deliver 1299 cFP8 TFLOPS, 667 BF16/FP16/TF32 TFLOPS and 181 FP32 TFLOPS dense, and 2563 sparse TFLOPS. A trn2.48xlarge or trn2u.48xlarge holds 16 chips wired as a 4x4 2D torus over NeuronLink-v3 with 3.2 Tbps of EFAv3 networking; trn2.3xlarge exposes a single chip. Four trn2u.48xlarge instances form a Trn2 UltraServer of 64 chips", + "Memory": ["HBM3", "SBUF", "PSUM"], + "Subunits": ["NeuronCore-v3", "Collective Compute Core", "DMA Engine", "HBM Stack", "NeuronLink Controller", "PCIe Controller"] + }, + "NeuronCore-v3": { + "Count": 8, + "Size": 4, + "Description": "Independent accelerator core with four heterogeneous engines (Tensor, Vector, Scalar, GpSimd) that execute concurrently over a shared software-managed 28 MiB state buffer, with support for control flow, dynamic shapes and a programmable rounding mode (round-nearest-even or stochastic). The Logical NeuronCore setting (LNC) determines how many cores software sees: LNC=2 is the default on Trainium2 and fuses two physical cores sharing one 24 GB HBM bank and address space into a logical core with software id NC_V3d, exposing 4 cores per chip and 64 per trn2.48xlarge; LNC=1 exposes each physical core separately as NC_V3, giving 8 per chip and 128 per instance. LNC is set by the NEURON_LOGICAL_NC_CONFIG runtime variable and the matching neuronx-cc --logical-nc-config flag", + "Memory": ["SBUF", "PSUM"], + "Subunits": ["Tensor Engine", "Vector Engine", "Scalar Engine", "GPSIMD Engine"] + }, + "Tensor Engine": { + "Count": 1, + "Size": 16384, + "Description": "Systolic array of 128 rows by 128 columns of processing elements per NeuronCore for matrix multiplication and convolution, streaming operands from SBUF and accumulating into PSUM. Delivers 79 BF16/FP16/TF32 TFLOPS and 158 cFP8 TFLOPS per core. In double-pumped FP8 mode each PE performs two pairs of FP8 multiplications per cycle, presenting a 256x128 array to the programmer. Supports configurable FP8 (cFP8, in FP8_E4 and FP8_E5 forms with adjustable exponent bias), FP16, BF16, TF32, FP32 and INT8. Structured sparsity in 4:16, 4:12, 4:8, 2:8, 2:4, 1:4 and 1:2 patterns raises throughput to 316 sparse TFLOPS per core", + "Memory": ["SBUF", "PSUM"] + }, + "Vector Engine": { + "Count": 1, + "Size": 128, + "Description": "SIMD unit per NeuronCore operating across the 128 SBUF partitions for element-wise and reduction operations such as normalization, softmax reductions and accumulation, rated at 1 TFLOPS of FP32. Runs SBUF-to-SBUF or PSUM-to-SBUF in parallel with the Tensor Engine", + "Memory": ["SBUF", "PSUM"] + }, + "Scalar Engine": { + "Count": 1, + "Size": 128, + "Description": "Per-element engine for activation functions and transcendentals such as GELU, exp, sigmoid and tanh, applied across the 128 SBUF partitions and rated at 1.2 TFLOPS of FP32, running concurrently with the other engines", + "Memory": ["SBUF", "PSUM"] + }, + "GPSIMD Engine": { + "Count": 1, + "Size": 8, + "Description": "Eight fully programmable 512-bit wide vector processors per NeuronCore with integrated DMA providing 307 GB/s (153 GB/s per direction). Handles irregular operations such as embeddings, gather/scatter and dynamic control flow that do not map to the fixed-function engines. On Trainium2 these are programmed through NKI; the aws-neuronx-gpsimd-customop-lib C++ custom-operator library ships only for Trn1 and Inf2", + "Memory": ["SBUF"] + }, + "Collective Compute Core": { + "Count": 16, + "Size": 1, + "Description": "16 CC-Cores per chip orchestrate collective communication such as all-reduce, all-gather and reduce-scatter among Trainium2 chips within and across instances, overlapping it with compute so the Tensor and Vector engines are not stalled. NeuronCore-v2 had 6", + "Memory": ["SBUF"] + }, + "DMA Engine": { + "Count": 128, + "Size": 1, + "Description": "128 DMA engines per chip, 16 associated with each NeuronCore-v3, moving data between HBM and SBUF and between cores at 3.5 TB/s with inline memory compression and decompression. All staging is explicit; there is no hardware cache between HBM and SBUF" + }, + "HBM Stack": { + "Count": 4, + "Size": 1, + "Description": "Four HBM stacks per chip forming four 24 GB banks, 96 GiB of device memory at 2.9 TB/s aggregate. Each bank is shared by two physical NeuronCore-v3, so under the default LNC=2 configuration one logical NeuronCore maps onto one 24 GB bank in a single shared address space" + }, + "NeuronLink Controller": { + "Count": 1, + "Size": 1, + "Description": "NeuronLink-v3 chip-to-chip interconnect controller providing 1.28 TB/s per chip and enabling memory pooling across up to 64 chips: 1024 GB/s per chip intra-instance across the 4x4 2D torus of a trn2.48xlarge, plus a further 256 GB/s per chip inter-instance for the UltraServer, where chips with the same coordinates in each of the four instances are joined in a ring. NeuronCore-v2 provided 384 GB/s per chip", + "Memory": ["HBM3"] + }, + "PCIe Controller": { + "Count": 1, + "Size": 1, + "Description": "Host interface controller. AWS does not publish the Trainium2 host link width or generation; PCIe 5.0 x16 is recorded here as an inference, not a vendor-stated fact" + } + } + }, + "MemorySubsystem": { + "SupportedMemoryTypes": ["HBM3"], + "MemoryTypes": [ + { + "Type": "SBUF (State Buffer, per NeuronCore)", + "Size": 28672, + "BankCount": 128, + "description": "28 MiB of on-chip software-managed SRAM per NeuronCore, organized as 128 partitions of 224 KiB, up from 24 MiB per core in NeuronCore-v2. Not a cache - there is no automatic fill, eviction or coherence. All engines read and write operands here, and data must be explicitly staged from HBM by the DMA engines. Totals 224 MiB per chip across the eight NeuronCores, against 48 MiB per chip on Trainium" + }, + { + "Type": "PSUM (Partial Sum Buffer, per NeuronCore)", + "Size": 2048, + "BankCount": 8, + "description": "2 MiB accumulation SRAM per NeuronCore, unchanged from NeuronCore-v2: 128 partitions of 16 KiB, each partition holding 8 banks of 512 32-bit values. Holds FP32 matmul partial sums written by the Tensor Engine and supports read-modify-write accumulation, so multi-pass matmuls accumulate in place before being copied back to SBUF" + }, + { + "Type": "HBM3 (Device DRAM)", + "Size": 100663296, + "BankCount": 4, + "MaxMemoryBandwidth": 2900, + "maxBusWidth": 4096, + "description": "96 GiB of device memory per chip across four 24 GB HBM3 stacks at 2.9 TB/s, up from 32 GiB at 0.8 TB/s on Trainium. A trn2.48xlarge exposes 1536 GiB at 46.4 TB/s across its 16 chips, and a Trn2 UltraServer 6144 GiB at 185.6 TB/s across 64. Bus width is derived from four JEDEC-standard 1024-bit HBM stacks and is not published by AWS. Moved to and from SBUF explicitly by the DMA engines; there is no hardware cache hierarchy between HBM and SBUF" + } + ] + }, + "KernelModel": { + "SubUnits": [ + { + "Type": "NKI (Neuron Kernel Interface)", + "Version": "0.6.0" + }, + { + "Type": "NKI Library", + "Version": "2.32.0" + }, + { + "Type": "Neuron Compiler (neuronx-cc)", + "Version": "2.27.5334.0" + }, + { + "Type": "PyTorch on Neuron (torch-xla)", + "Version": "2.12.0" + }, + { + "Type": "Neuron SDK", + "Version": "2.32.0" + } + ] + }, + "SpecializedHardware": { + "RayTracingAccelerators": { + "Present": false + }, + "AiAccelerators": { + "Present": true, + "Name": "NeuronCore-v3 Tensor Engine", + "Count": 8, + "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "Other"], + "Performance": { + "fp16TopsPerGPU": 667, + "int8TopsPerGPU": 1299 + } + }, + "videoCodecs": { + "encoders": [], + "decoders": [] + } + }, + "PowerEfficiency": { + "MaxTDP": 500, + "PowerStates": [ + { + "Name": "Active", + "Description": "Full power state with all eight NeuronCores and HBM stacks active. AWS does not publish a per-chip TDP for Trainium2; the ~500W figure recorded here is a third-party estimate from SemiAnalysis, not a vendor-stated number" + }, + { + "Name": "Idle", + "Description": "Reduced power state when no workloads are running" + } + ], + "ClockGating": true, + "DynamicVoltageFrequencyScaling": true + }, + "PciExpress": { + "Version": "5.0", + "Lanes": 16, + "Bandwidth": 63 + }, + "MultiGpuSupport": { + "Technologies": ["Other"], + "MaxGpus": 64, + "InterconnectBandwidth": 1280 + } +} diff --git a/device_lib/nvidia-gpu-sm_100.json b/device_lib/nvidia-gpu-sm_100.json new file mode 100644 index 00000000..c65b11bf --- /dev/null +++ b/device_lib/nvidia-gpu-sm_100.json @@ -0,0 +1,262 @@ +{ + "Name": "Blackwell", + "Vendor": "NVIDIA", + "Architecture": "sm_100", + "ReleaseYear": 2024, + "FabricationProcess": { + "ProcessNode": 5, + "Manufacturer": "TSMC", + "Technology": "4NP" + }, + "CoreSubsystem": { + "Name": "SM", + "ChipType": "Package", + "CoreType": "SM", + "UnitTypes": { + "Package": { + "Count": 1, + "Size": 2, + "Description": "GB100 dual-die package, the datacenter Blackwell part with compute capability 10.0. The primary configuration described by this file is B200 SXM as shipped in HGX B200 and DGX B200: 148 SMs, 180 GB HBM3E, 7.7 TB/s, configurable up to 1000 W. NVIDIA states 208 billion transistors per package and that each of the two dies is the largest that reticle limits allow; the per-die figure of 104 billion is simply that total halved by reporters, and NVIDIA publishes no die area. The dies are joined by NV-HBI at 10 TB/s into what NVIDIA calls one fully coherent chip, so CUDA enumerates the package as a SINGLE GPU with one L2 and one address space; every count in this file is labelled per-die or per-package. Related SKUs on the same silicon: B100 is a lower-power 700 W part (60 TFLOPS FP32, 192 GB) now effectively delisted by NVIDIA, and the GB200 Grace Blackwell Superchip pairs one Grace CPU with TWO of these packages over NVLink-C2C at 900 GB/s, running a higher power bin of up to 1200 W per GPU for roughly 11 percent more throughput at every precision (80 vs 75 TFLOPS FP32) and exposing 186 GB HBM3E at 8 TB/s. Blackwell adds a second-generation Transformer Engine with micro-tensor scaling, a Decompression Engine, a dedicated RAS engine, and the industry's first TEE-I/O capable GPU for Confidential Computing with inline NVLink protection. Supports up to 7 MIG instances, the smallest being 23 GB, and exposes 16 copy engines on the full GPU, double H100's 8, consistent with the dual-die build. NOTE ON RAY TRACING: this file deliberately omits a RayTracingAccelerators block rather than asserting a value. NVIDIA's Blackwell Datasheet enumerates MIG, the Decompression Engine and the decoders for this GPU but lists no RT cores, and no NVIDIA datacenter Blackwell material mentions ray tracing; the fourth-generation RT core claims made for Blackwell come from the consumer RTX Blackwell whitepaper and describe the GB20x dies (sm_120), not GB100. Since NVIDIA states neither presence nor absence for GB100, neither is recorded. Note on FabricationProcess: NVIDIA markets the node as 4NP, a custom enhancement of the 4N node used for Hopper; ProcessNode is recorded as 5 because TSMC places N4P in its 5 nm family, matching the sm_90 entry for the same node family", + "Memory": ["L2 Cache", "HBM3E (Device DRAM)", "Constant Memory"], + "Subunits": ["Die", "NV-HBI Link", "NVLink 5 Link", "PCIe Controller", "NVDEC Engine", "NVJPEG Engine", "Decompression Engine"] + }, + "Die": { + "Count": 2, + "Size": 37, + "Description": "One GB100 compute die, two per package, each at the reticle limit of TSMC 4NP lithography. NVIDIA does not publish the die area (GH100 was 814 mm2 for comparison), nor the GPC/TPC organization of GB100, so the Graphics Processing Cluster level present in earlier NVIDIA entries is deliberately OMITTED here rather than guessed. The package total of 148 SMs IS published, in NVIDIA's CUDA MPS documentation, which describes partitioning an HGX B200 GPU that normally has 148 SMs. The per-die split is not published by NVIDIA: independent die analysis reports 80 SMs physically present per die with 74 enabled, which is the 74-per-die and 37-TPC figure recorded in this hierarchy at the long-standing 2 SMs per TPC. That split is corroborated arithmetically, since 148 SMs of 128 FP32 lanes at 2 FLOP per clock reproduce NVIDIA's published 75 TFLOPS peak FP32 and 37 TFLOPS FP64 per B200 GPU at a boost clock near 1980 MHz, and imply exactly twice Hopper's per-SM per-clock Tensor Core throughput, which is the documented fifth-generation uplift. Each die carries 4 HBM3E stacks and about 63 MB of the package's 126 MB L2; measurements show the two L2 partitions correspond to the two dies, with near-die latency faster than across NV-HBI", + "Subunits": ["TPC", "Memory Controller", "HBM3E Stack"] + }, + "TPC": { + "Count": 37, + "Size": 2, + "Description": "Texture Processing Cluster containing 2 SMs. 37 enabled per die and 74 per package on B200 (40 per die if the reported 80 physical SMs per die are counted). Derived rather than published: NVIDIA has issued no GB100 block diagram and the terms GPC and TPC appear nowhere in its Blackwell architecture or tuning documentation. NVIDIA has placed 2 SMs per TPC in every architecture since Kepler, and on Blackwell that pair is also the granularity of the tcgen05 CTA-pair MMA, in which two SMs cooperate on one matrix operation", + "Subunits": ["SM"] + }, + "SM": { + "Count": 2, + "Size": 128, + "Description": "Blackwell Streaming Multiprocessor, 74 per die and 148 per B200 package. NVIDIA's CUDA C++ Programming Guide states verbatim for compute capability 10.0 that an SM consists of 128 FP32 cores, 64 FP64 cores, 64 INT32 cores, 4 mixed-precision fifth-generation Tensor Cores with sparsity support, 16 special function units and 4 warp schedulers, organized as 4 processing blocks. NVIDIA does NOT publish load/store or texture unit counts for compute capability 10.0, so unlike the sm_90 entry those units are omitted here rather than carried over from Hopper. Occupancy and storage limits are unchanged from compute capability 9.0: 2048 resident threads (64 warps), 32 resident thread blocks, 255 registers per thread, 65536 32-bit registers per SM, 228 KB of shared memory per SM out of a 256 KB unified data cache, and a 2:1 FP32:FP64 throughput ratio. What is new at 10.0 is the fifth-generation Tensor Core with its dedicated Tensor Memory and tcgen05 instruction family, FP6 and FP4 tensor input types, and a non-portable thread block cluster size of 16 on B200 against the portable maximum of 8, opted into with cudaFuncAttributeNonPortableClusterSizeAllowed", + "Memory": ["Register File (per SM)", "L1 Data Cache / Shared Memory (per SM)", "Tensor Memory (per SM)"], + "Subunits": ["FP32 CUDA Core", "INT32 Core", "FP64 Core", "Tensor Core", "SFU", "Warp Scheduler", "Tensor Memory Accelerator"] + }, + "FP32 CUDA Core": { + "Count": 128, + "Size": 1, + "Description": "Single-precision FP32 lane, 32 per processing block, confirmed by the 128 FP32 add/multiply/FMA operations per clock per SM listed for compute capability 10.0 in NVIDIA's native arithmetic throughput table. 18944 per B200 package across 148 SMs, giving NVIDIA's published 75 TFLOPS of peak non-tensor FP32 per B200 GPU (600 TFLOPS across the 8 GPUs of an HGX B200 board) and 80 TFLOPS per GPU on GB200. Both figures were revised down after launch, when NVIDIA quoted 80 and 90 TFLOPS respectively" + }, + "INT32 Core": { + "Count": 64, + "Size": 1, + "Description": "32-bit integer lane, 16 per processing block, matching the 64 INT32 add operations per clock per SM published for compute capability 10.0 and executing concurrently with the FP32 datapath. Consumer Blackwell (sm_120) doubles this to 128" + }, + "FP64 Core": { + "Count": 64, + "Size": 1, + "Description": "Double-precision lane, 16 per processing block, giving the 2:1 FP32:FP64 ratio published for compute capability 10.0. NVIDIA quotes 296 TFLOPS FP64 per 8-GPU HGX B200 board, i.e. 37 TFLOPS per B200 GPU, and 40 TFLOPS per GPU on GB200. Unlike Hopper the FP64 Tensor Core rate equals the vector FP64 rate rather than doubling it, and sm_100 is the only Blackwell compute capability that retains FP64 tensor cores at all, since 10.3, 11.0 and 12.x all drop them; consumer sm_120 falls to 2 FP64 operations per clock per SM" + }, + "Tensor Core": { + "Count": 4, + "Size": 1, + "Description": "Fifth-generation Tensor Core, one per processing block, 592 per B200 package across 148 SMs. Adds FP6 (E2M3/E3M2) and FP4 (E2M1) input types over Hopper while keeping FP64, TF32, BF16, FP16, FP8 (E4M3/E5M2) and INT8, and is driven by the new tcgen05 instruction family, which reads and writes the per-SM Tensor Memory and can pair two SMs on a single MMA. Supports the OCP microscaling block formats (MXFP8/MXFP6/MXFP4, a UE8M0 power-of-two scale per 32 elements) and NVIDIA's own NVFP4 (a UE4M3 scale per 16 elements). Per B200 GPU NVIDIA publishes 18 PFLOPS FP4, 9 PFLOPS FP8/FP6, 9 POPS INT8, 4.5 PFLOPS FP16/BF16 and 2.2 PFLOPS TF32, ALL WITH 2:4 STRUCTURED SPARSITY; NVIDIA's own footnote reads that dense is one-half of the sparse spec shown, so the DENSE rates are 9 PFLOPS FP4, 4.5 PFLOPS FP8/FP6, 4.5 POPS INT8, 2.25 PFLOPS FP16/BF16 and 1.1 PFLOPS TF32, and NVIDIA's architecture brief prints the same pairs explicitly as dense/sparse. GB200 is about 11 percent higher throughout: 20 PFLOPS FP4 sparse and 10 PFLOPS dense per GPU, 5 PFLOPS FP16/BF16 sparse and 2.5 dense. The second-generation Transformer Engine applies micro-tensor scaling to make FP4 usable, doubling FP4 Tensor Core performance, HBM parameter bandwidth and model size per GPU" + }, + "SFU": { + "Count": 16, + "Size": 1, + "Description": "Special Function Unit, 4 per processing block, computing single-precision transcendentals such as sin, cos, exp2, log2, rcp and rsqrt" + }, + "Warp Scheduler": { + "Count": 4, + "Size": 1, + "Description": "One warp scheduler per processing block. An SM statically distributes its warps among the 4 schedulers, each issuing one instruction per cycle for one ready warp, up to the 64-warp per-SM occupancy limit" + }, + "Tensor Memory Accelerator": { + "Count": 1, + "Size": 1, + "Description": "TMA, the per-SM asynchronous copy engine introduced on Hopper for moving multidimensional tensor tiles between global and shared memory. Blackwell extends it with tile::gather4 and tile::scatter4 modes that gather four non-contiguous rows into one tile or scatter one tile across four rows, the im2col::w and im2col::w::128 convolution modes, and a cta_group qualifier that multicasts into the two-SM CTA pair used by tcgen05. These additions belong to the sm_100 family and are not available on sm_120" + }, + "Memory Controller": { + "Count": 4, + "Size": 1, + "Description": "HBM3E memory controller, one per stack, 4 per die and 8 per package. NVIDIA does not publish the stack count or bus width for GB100; independent reporting describes 4 stacks per die at 1024 bits each for an 8192-bit aggregate interface, which is the organization recorded here" + }, + "HBM3E Stack": { + "Count": 4, + "Size": 1, + "Description": "24 GB HBM3E stack, 4 per die and 8 per package. NVIDIA does not publish the stack count or stack height; the 8-stack organization and the resulting 24 GB per stack are inferred from the 192 GB of physical capacity announced at GTC 2024. Shipping SKUs expose less than the full 192 GB: 180 GB per GPU on HGX B200 and DGX B200 (1440 GB per 8-GPU board) and 186 GB per GPU on GB200 (13.4 TB across the 72 GPUs of an NVL72 rack)", + "Memory": ["HBM3E (Device DRAM)"] + }, + "NV-HBI Link": { + "Count": 1, + "Size": 1, + "Description": "NV-High Bandwidth Interface, one per package: the 10 TB/s die-to-die link that NVIDIA describes as connecting and unifying the two dies into one fully coherent chip. Fast enough that the two dies share a single coherent L2 address space, so software sees one monolithic GPU rather than two devices, with no CUDA-visible NUMA split and no peer-access call between dies. NVIDIA does not state whether the 10 TB/s figure is aggregate or per direction" + }, + "NVLink 5 Link": { + "Count": 18, + "Size": 1, + "Description": "Fifth-generation NVLink port, counted per package since NVLink is exposed at whole-GPU level. NVIDIA states 18 links per GPU giving 1.8 TB/s total, 900 GB/s in each direction; as on Hopper each link uses two high-speed differential pairs per direction, but Blackwell doubles the per-link rate to 50 GB/s each direction, i.e. 100 GB/s bidirectional per link. Through the NVLink 5 Switch these form scale-up domains of 8 GPUs (HGX B200, 14.4 TB/s aggregate) or 72 GPUs (GB200 NVL72, 130 TB/s aggregate across 9 switch trays and 18 compute trays); NVIDIA states the fifth generation scales up to 576 GPUs" + }, + "PCIe Controller": { + "Count": 1, + "Size": 1, + "Description": "PCIe Gen 5 x16 host interface, one per package, 128 GB/s bidirectional. NVIDIA's Blackwell Datasheet and architecture brief both list PCIe Gen5 for HGX B200, GB200 NVL72 and GB200 NVL4, reserving PCIe Gen6 for the later B300; NVIDIA's own MLPerf submissions record the host interconnect as PCIe Gen5 x16. Some launch-day press tables claimed Gen6 for B200, which the shipping documentation does not support. In GB200 the GPU reaches its Grace CPU over NVLink-C2C at 900 GB/s rather than PCIe" + }, + "NVDEC Engine": { + "Count": 7, + "Size": 1, + "Description": "Sixth-generation NVDEC hardware video decoder, 7 per GPU. NVIDIA's Blackwell Datasheet lists 7 NVDEC and NVIDIA's video encode/decode support matrix confirms 7 NVDEC per chip and ZERO NVENC encoders for both HGX B200 and GB200, exactly as on H100. The generation jump from Hopper's fourth-generation NVDEC adds AV1 decode plus H.264 4:2:2 and H.265 4:2:2 and 4:4:4, none of which H100 supported, alongside MPEG-1/2/4, VC-1, VP8, VP9 and H.265 4:2:0 at 8, 10 and 12 bits. NVIDIA quotes 2x H.264, 1.25x HEVC and 1.25x VP9 decode throughput over H100" + }, + "NVJPEG Engine": { + "Count": 7, + "Size": 1, + "Description": "Hardware JPEG decode engine, 7 per GPU per NVIDIA's Blackwell Datasheet, exposed through nvJPEG for data-loading pipelines" + }, + "Decompression Engine": { + "Count": 1, + "Size": 1, + "Description": "Dedicated hardware decompression engine new to Blackwell, decompressing at up to 800 GB/s with support for LZ4, Snappy and Deflate, and feeding compressed data straight into the GPU for database and analytics pipelines. NVIDIA quotes 18x CPU and 6x H100 on query benchmarks. NVIDIA describes the engine but does not publish how many instances a GPU carries, so it is recorded here as one logical engine per package" + } + } + }, + "MemorySubsystem": { + "SupportedMemoryTypes": ["HBM3", "Other"], + "MemoryTypes": [ + { + "Type": "Register File (per SM)", + "Size": 256, + "BankCount": 4, + "description": "65536 32-bit registers (256 KB) per SM, split evenly across the 4 processing blocks, unchanged from compute capability 9.0. A single thread may use at most 255 registers and a single thread block at most 65536. Aggregate register file is about 37 MB per B200 package across 148 SMs" + }, + { + "Type": "L1 Data Cache / Shared Memory (per SM)", + "Size": 256, + "BankCount": 32, + "description": "256 KB of unified L1 data cache, texture cache and shared memory per SM, the same total as Hopper and twice the 128 KB of consumer Blackwell sm_120. The shared-memory carveout is configurable to 0, 8, 16, 32, 64, 100, 132, 164, 196 or 228 KB per SM; a single thread block may address at most 227 KB because CUDA reserves 1 KB, and allocations above 48 KB must use dynamic shared memory with an explicit opt-in. NVIDIA publishes no separate L1 figure: L1 is whatever remains of the 256 KB after the carveout. Shared memory is banked 32 ways and is addressable across a thread block cluster as distributed shared memory" + }, + { + "Type": "Tensor Memory (per SM)", + "Size": 256, + "description": "Tensor Memory (TMEM), dedicated on-chip storage introduced with the fifth-generation Tensor Core and entirely new at compute capability 10.0. NVIDIA documents it only in the PTX ISA and states a geometry rather than a byte total: 512 columns by 128 lanes per CTA with each cell 32 bits, which works out to 256 KB. It is allocated dynamically by a single warp via tcgen05.alloc in granules of 32 columns, always spanning all 128 lanes, and must be explicitly freed before the kernel exits. Access is partitioned, each warp of a warpgroup reaching only its own 32 lanes, so a full warpgroup is needed to cover an allocation. It exists only on the sm_100 family; sm_120 has neither Tensor Memory nor the tcgen05 instructions" + }, + { + "Type": "L2 Cache", + "Size": 129024, + "description": "126 MB of L2 per package, more than double Hopper's 50 MB on H100, and kept coherent across NV-HBI so the two dies present a single L2 address space. NVIDIA publishes only the package total; the roughly 63 MB per die follows from measurements showing the two L2 partitions correspond to the two dies, with accesses to the near partition faster than those crossing NV-HBI. Supports the L2 residency and persistence controls introduced with Ampere" + }, + { + "Type": "HBM3E (Device DRAM)", + "Size": 188743680, + "BankCount": 8, + "MaxMemoryBandwidth": 7700, + "maxBusWidth": 8192, + "description": "180 GB of HBM3E per B200 GPU at 7.7 TB/s across 8 stacks, 4 attached to each die but pooled into one address space. NVIDIA's own figures disagree slightly: the Blackwell Datasheet says 180 GB at 7.7 TB/s per GPU while DGX B200 says 1440 GB and 64 TB/s across 8 GPUs, i.e. 180 GB at 8 TB/s, and an early architecture-brief cell still reads 192 GB while the rest of its own table implies 180. The package physically carries 192 GB, the capacity announced at GTC 2024. GB200 exposes 186 GB per GPU at 8 TB/s, an NVL72 rack totalling 13.4 TB at 576 TB/s. The 8192-bit bus width is not published by NVIDIA and is inferred from 8 stacks of 1024 bits. The Blackwell tuning guide records support for both HBM3 and HBM3E; HBM3E has no dedicated value in this schema's memory-type enum and is reported as Other. ECC is enabled on all SKUs" + }, + { + "Type": "Constant Memory", + "Size": 64, + "description": "64 KB constant memory window per device context at compute capability 10.0, served by a per-SM read-only constant cache with an 8 KB working set" + } + ] + }, + "KernelModel": { + "LLVMTarget": "nvptx64", + "LLVMTriple": "nvptx64-nvidia-cuda", + "LLVMFeatures": [ + { + "Name": "+sm_100", + "Description": "NVPTX subtarget feature selecting the datacenter Blackwell sm_100 target ISA, added in LLVM 20. Portable across compute capability 10.0 and later; excludes architecture-specific instructions" + }, + { + "Name": "+sm_100a", + "Description": "Architecture-specific sm_100a variant exposing the complete Blackwell-only instruction set, including the tcgen05 fifth-generation Tensor Core MMA family, Tensor Memory allocation, and the gather4/scatter4 and im2col::w TMA modes. Per NVIDIA, targets with an a suffix do not follow the onion-layer model, so code built for sm_100a runs on compute capability 10.0 devices only and is neither forward nor backward compatible. Added in LLVM 20" + }, + { + "Name": "+sm_100f", + "Description": "Family-specific sm_100f variant, a target class introduced with compute capability 10.0 that sits between the portable and architecture-specific targets: it permits the subset of architecture-specific features shared across the sm_10x family, so binaries run on compute capability 10.0 and 10.3 but not outside the family. Requires CUDA 12.9 or later and appears in LLVM 21, where the NVPTX backend accepts sm_100f although the Clang driver's offload-arch list does not yet" + }, + { + "Name": "+ptx86", + "Description": "PTX ISA 8.6, the first version in which sm_100 and sm_100a became legal target directives. It shipped in the r565 driver stream rather than in a public toolkit, and LLVM pairs both sm_100 and sm_100a with this feature" + }, + { + "Name": "+ptx87", + "Description": "PTX ISA 8.7, shipped in CUDA 12.8 with the r570 driver, the first publicly released toolkit able to compile for sm_100" + } + ], + "SubUnits": [ + { + "Type": "CUDA", + "Version": "12.8" + }, + { + "Type": "PTX ISA", + "Version": "8.6" + }, + { + "Type": "OpenCL", + "Version": "3.0" + } + ] + }, + "SpecializedHardware": { + "AiAccelerators": { + "Present": true, + "Name": "Fifth-Generation Tensor Core", + "Count": 592, + "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "Other"], + "Performance": { + "fp16TopsPerGPU": 2250, + "int8TopsPerGPU": 4500 + } + }, + "videoCodecs": { + "encoders": [], + "decoders": [ + { + "codec": "H.264" + }, + { + "codec": "H.265/HEVC" + }, + { + "codec": "AV1" + }, + { + "codec": "VP9" + }, + { + "codec": "Other" + } + ] + } + }, + "PowerEfficiency": { + "BoostClock": 1980, + "MaxTDP": 1000, + "PowerStates": [ + { + "Name": "Maximum Performance", + "Description": "Full power state for B200 SXM, which NVIDIA specifies as configurable up to 1000 W per GPU. NVIDIA publishes NO base or boost clock for any GB100 part, and leaves the accelerator frequency field blank in its own MLPerf submissions; the 1980 MHz recorded here is derived from the published 75 TFLOPS peak non-tensor FP32 per B200 GPU over 148 SMs of 128 FP32 lanes at 2 FLOP per clock, and the same arithmetic gives about 2110 MHz for GB200. Independent analysis reports B200 clocks similar to the 1980 MHz H100 SXM5. DGX B200 ships with a 700 W default power cap, adjustable between 200 and 1000 W. B100 is a 700 W part on the same silicon, and GB200 is specified configurable up to 1200 W per GPU under NVL72 liquid cooling" + }, + { + "Name": "Balanced", + "Description": "Reduced clock and voltage operating point selected by the driver when the workload or the power and thermal budget does not require peak clocks" + }, + { + "Name": "Idle", + "Description": "Low power state when no kernels are resident" + } + ], + "ClockGating": true, + "DynamicVoltageFrequencyScaling": true + }, + "PciExpress": { + "Version": "5.0", + "Lanes": 16, + "Bandwidth": 128 + }, + "MultiGpuSupport": { + "Technologies": ["NVLink"], + "MaxGpus": 72, + "InterconnectBandwidth": 1800 + } +} diff --git a/device_lib/nvidia-gpu-sm_120.json b/device_lib/nvidia-gpu-sm_120.json new file mode 100644 index 00000000..bbad63fb --- /dev/null +++ b/device_lib/nvidia-gpu-sm_120.json @@ -0,0 +1,275 @@ +{ + "Name": "Blackwell (GB20x)", + "Vendor": "NVIDIA", + "Architecture": "sm_120", + "ReleaseYear": 2025, + "FabricationProcess": { + "ProcessNode": 4, + "Manufacturer": "TSMC", + "Technology": "4N NVIDIA Custom Process" + }, + "CoreSubsystem": { + "Name": "SM", + "ChipType": "Chip", + "CoreType": "SM", + "UnitTypes": { + "Chip": { + "Count": 1, + "Size": 170, + "Description": "GB202 die: 92.2 billion transistors on 750 mm2, built on the TSMC 4nm 4N NVIDIA Custom Process. This is consumer/workstation Blackwell (compute capability 12.0, sm_120), a different die family from the GB100 datacenter Blackwell of B100/B200/GB200 (compute capability 10.0, sm_100). The primary configuration described by this file is the GeForce RTX 5090: 170 SMs, 21760 CUDA cores, 680 fifth-generation Tensor Cores, 170 fourth-generation RT Cores, 176 ROPs, 96 MB L2, 32 GB GDDR7 on a 512-bit bus, 575 W. The full GB202 die is larger than any GeForce product: 12 GPCs, 96 TPCs, 192 SMs, 24576 CUDA cores, 768 Tensor Cores, 192 RT Cores, 768 texture units and 128 MB of L2. Other sm_120 parts share the same SM and differ only in scale: GeForce RTX 5080 uses a full GB203 (7 GPCs, 42 TPCs, 84 SMs, 10752 CUDA cores, 64 MB L2, 16 GB GDDR7 on 256-bit at 960 GB/s, 2617 MHz boost, 360 W), and GeForce RTX 5070 uses GB205 (full die 5 GPCs, 25 TPCs, 50 SMs, 6400 CUDA cores, 48 MB L2, 192-bit). The RTX PRO Blackwell workstation line is built from the same GB20x dies: RTX PRO 6000 Blackwell Workstation Edition is the largest GB202 configuration shipped, with 24064 CUDA cores (188 SMs, derived from the published core count over 128 cores per SM rather than stated directly), 96 GB of GDDR7 with ECC on the same 512-bit 1792 GB/s interface, 126 TFLOPS FP32, 382 RT TFLOPS, 4 NVENC and 4 NVDEC engines, MIG partitioning and a 600 W board power. GB202 is a graphics part: it has display outputs, RT Cores and NVENC, and unlike the datacenter Blackwell it has no NVLink and only token FP64 hardware", + "Memory": ["L2 Cache", "GDDR7 (Device DRAM)", "Constant Memory"], + "Subunits": ["GPC", "Memory Controller", "GDDR7 Device", "PCIe Controller", "NVENC Engine", "NVDEC Engine", "AI Management Processor"] + }, + "GPC": { + "Count": 11, + "Size": 8, + "Description": "Graphics Processing Cluster, the dominant high-level block in every GB20x GPU. Each GPC holds a dedicated Raster Engine, 8 TPCs (16 SMs) and two ROP partitions of 8 ROP units each, for 16 ROPs per GPC. The full GB202 die has 12 GPCs and 96 TPCs; RTX 5090 enables 11 GPCs but only 85 of the 88 TPCs those GPCs would hold, so the enabled GPCs are not populated uniformly. The Size of 8 is the full-GPC TPC complement", + "Subunits": ["TPC", "Raster Engine", "ROP Partition"] + }, + "TPC": { + "Count": 8, + "Size": 2, + "Description": "Texture Processing Cluster, containing one PolyMorph Engine and 2 SMs. 96 TPCs on the full GB202 die, 85 enabled on RTX 5090 for 170 SMs; 42 on the full GB203 used by RTX 5080 and 25 on the full GB205", + "Subunits": ["SM", "PolyMorph Engine"] + }, + "SM": { + "Count": 2, + "Size": 128, + "Description": "Blackwell Streaming Multiprocessor, organized as 4 processing blocks (partitions) that each hold one warp scheduler, a 64 KB register file slice, one fifth-generation Tensor Core, 32 CUDA-core lanes, 4 SFUs and 1 texture unit. The RTX Blackwell whitepaper describes the SM as 128 CUDA cores, one fourth-generation RT Core, four fifth-generation Tensor Cores, 4 texture units, a 256 KB register file and 128 KB of L1/shared memory. Blackwell's headline SM change over Ada is that the INT32 datapath is fully unified with the FP32 datapath: the 128 lanes execute either FP32 or INT32 in any given clock, which roughly doubles peak integer throughput versus Ada (RTX 5090 reaches 104.8 peak INT32 TOPS, equal to its 104.8 peak non-Tensor FP32 TFLOPS, where Ada's RTX 4090 managed 41.3 INT32 TOPS against 82.6 FP32 TFLOPS). Not every integer instruction attains the 2x rate. Note that the CUDA C++ Programming Guide's compute capability 12.0 section still enumerates the SM as 128 FP32 cores plus 64 separate INT32 cores, which is the Ada layout; the whitepaper's v1.1 addendum explicitly supersedes that, and the whitepaper's own 104.8 INT32 TOPS figure requires 128 integer-capable lanes. NVIDIA also states the SM doubles point-sampled texture performance per cycle versus Ada. Compute capability 12.0 limits: 1536 resident threads (48 warps), 24 resident thread blocks, 65536 32-bit registers per SM, 255 registers per thread, 1024 threads per block, warp size 32. These occupancy limits match Ada sm_89 rather than Hopper sm_90. The CUDA 12.8 and 12.9 guides listed 32 resident blocks per SM for 12.0; the 13.x guides give 24, which is the value recorded here and the one consistent with the 48-warp limit. sm_120 does support thread block clusters, distributed shared memory and the Tensor Memory Accelerator, but NVIDIA publishes no 12.0-specific maximum cluster size, so only the portable maximum of 8 blocks applies and applications should query cudaOccupancyMaxPotentialClusterSize. NVIDIA states no LD/ST unit count for compute capability 12.0 in text (the SM block diagram draws 4 LD/ST blocks per partition, but that is a figure rather than a published figure), so that unit type is omitted from this file rather than guessed", + "Memory": ["Register File (per SM)", "L1 Data Cache / Shared Memory (per SM)"], + "Subunits": ["CUDA Core", "FP64 Core", "Tensor Core", "RT Core", "SFU", "Warp Scheduler", "Texture Unit"] + }, + "CUDA Core": { + "Count": 128, + "Size": 1, + "Description": "Unified FP32/INT32 lane, 32 per processing block, able to issue either an FP32 or an INT32 operation in a given clock. 21760 per RTX 5090 (24576 on the full GB202 die, 10752 on the full GB203, 6400 on the full GB205), giving 104.8 TFLOPS of peak non-Tensor FP32 at the 2407 MHz boost clock, and the same 104.8 TFLOPS for non-Tensor FP16 and BF16" + }, + "FP64 Core": { + "Count": 2, + "Size": 1, + "Description": "Double-precision lane, 2 per SM (384 on the full GB202 die), deliberately minimal. NVIDIA states the FP64 rate is 1/64 of the FP32 rate and that these cores exist only to ensure programs containing FP64 code operate correctly. A very small number of FP64 Tensor Core paths is included for the same correctness reason. sm_120 is not a double-precision compute part" + }, + "Tensor Core": { + "Count": 4, + "Size": 1, + "Description": "Fifth-generation Tensor Core, one per processing block (680 per RTX 5090, 768 on the full GB202 die, 336 on the full GB203). Adds FP4 and FP6 operand support and a second-generation FP8 Transformer Engine over Ada's fourth generation, alongside FP8 (E4M3/E5M2), FP16, BF16, TF32 and INT8, plus 2:4 structured sparsity; a token FP64 matrix path exists for correctness only. NVIDIA does not list INT4 among the Blackwell RTX Tensor Core precisions, so INT4 is not claimed in SupportedPrecisions. RTX 5090 peak rates, quoted as dense/sparse: FP4 with FP32 accumulate 1676/3352 TFLOPS, FP8 with FP16 accumulate 838/1676 TFLOPS, FP8 with FP32 accumulate 419/838 TFLOPS, FP16 with FP16 accumulate 419/838 TFLOPS, FP16 and BF16 with FP32 accumulate 209.5/419 TFLOPS, TF32 104.8/209.5 TFLOPS, INT8 838/1676 TOPS. The 3352 AI TOPS NVIDIA markets for the RTX 5090 is the FP4 figure with sparsity, i.e. 1676 TFLOPS dense. RTX 5080 peaks at 450.2/900.4 dense/sparse INT8 TOPS and 225.1/450.2 FP16-accumulate FP16 TFLOPS. Unlike Hopper sm_90 and datacenter Blackwell sm_100, sm_120 Tensor Cores are driven by warp-level mma rather than warp-group wgmma; sm_120a adds block-scaled mma variants for the .e3m2/.e2m3/.e2m1 FP6 and FP4 types" + }, + "RT Core": { + "Count": 1, + "Size": 1, + "Description": "Fourth-generation RT Core, one per SM (170 per RTX 5090, 192 on the full GB202 die, 84 on the full GB203). Contains dedicated ray-bounding-box and ray-triangle intersection engines and doubles ray-triangle intersection throughput over Ada's third generation. Blackwell adds a Triangle Cluster Intersection Engine and Triangle Cluster Compression for the new Mega Geometry path, hardware Linear Swept Spheres primitives, and an opacity micromap engine for alpha computation. NVIDIA rates RTX 5090 at 317.5 RT TFLOPS and no longer publishes a giga-rays-per-second figure for this generation, so RaysPerSecond is omitted from SpecializedHardware rather than estimated" + }, + "SFU": { + "Count": 16, + "Size": 1, + "Description": "Special Function Unit, 4 per processing block, computing single-precision floating-point transcendentals such as sin, cos, exp2, log2, rcp and rsqrt. The count of 16 per SM is stated by the CUDA C++ Programming Guide's compute capability 12.0 architecture section" + }, + "Warp Scheduler": { + "Count": 4, + "Size": 1, + "Description": "One warp scheduler per processing block. An SM statically distributes its warps among the 4 schedulers, and at each issue slot a scheduler issues one instruction for one of its ready assigned warps. Warp size is 32 threads and up to 48 warps are resident per SM, i.e. 12 per scheduler" + }, + "Texture Unit": { + "Count": 4, + "Size": 1, + "Description": "Texture/sampler unit, one per processing block. 680 per RTX 5090, up from 512 on RTX 4090, giving a bilinear-filtered texel rate of 1636.8 Gigatexels/s. 768 on the full GB202 die" + }, + "PolyMorph Engine": { + "Count": 1, + "Size": 1, + "Description": "Fixed-function geometry engine, one per TPC, handling vertex fetch, tessellation, viewport transform and attribute setup ahead of the GPC's Raster Engine" + }, + "Raster Engine": { + "Count": 1, + "Size": 1, + "Description": "Fixed-function rasterizer, one per GPC, converting the geometry stream into pixel fragments for the SMs" + }, + "ROP Partition": { + "Count": 2, + "Size": 8, + "Description": "Raster Operations partition of 8 ROP units, 2 partitions per GPC for 16 ROPs per GPC. RTX 5090 exposes 176 ROPs across its 11 GPCs, for a pixel fill rate of 423.6 Gigapixels/s; RTX 5080 exposes 112" + }, + "Memory Controller": { + "Count": 16, + "Size": 1, + "Description": "32-bit GDDR7 memory controller. The full GB202 die has 16 for a 512-bit interface and RTX 5090 enables all 16. The full GB203 has 8 (256-bit) and the full GB205 has 6 (192-bit)" + }, + "GDDR7 Device": { + "Count": 16, + "Size": 1, + "Description": "GDDR7 memory device, one per 32-bit controller. RTX 5090 carries 32 GB across the 512-bit interface at 28 Gbps for 1792 GB/s; NVIDIA publishes the capacity and the controller count but not the per-package capacity, so the 2 GB per channel implied by 32 GB over 16 channels is an inference from those two published figures", + "Memory": ["GDDR7 (Device DRAM)"] + }, + "PCIe Controller": { + "Count": 1, + "Size": 1, + "Description": "PCI Express Gen 5 x16 host interface, 128 GB/s bidirectional. This is the only GPU-to-GPU path on GB20x: consumer and workstation Blackwell ship no NVLink and no SLI" + }, + "NVENC Engine": { + "Count": 3, + "Size": 1, + "Description": "Ninth-generation NVENC hardware video encoder, 3 per GB202 as configured on RTX 5090 (4 on RTX PRO 6000 Blackwell; 2 on GB203, 2 on GB205). Encodes H.264 including 4:2:2 and 4:4:4, HEVC up to 8K including 4:2:2, 4:4:4 and 10-bit, and AV1 up to 10-bit. Improves AV1 and HEVC quality by about 5% BD-BR PSNR over the eighth-generation Ada encoder, adds 4:2:2 H.264 and HEVC encoding, and adds an AV1 Ultra Quality mode taking the combined gain to about 15% BD-BR PSNR. No NVIDIA NVENC encodes VP9" + }, + "NVDEC Engine": { + "Count": 2, + "Size": 1, + "Description": "Sixth-generation NVDEC hardware video decoder, 2 per GB202 as configured on RTX 5090 (4 on RTX PRO 6000 Blackwell; 2 on GB203, 1 on GB205). Decodes MPEG-1/2/4, VC-1, VP8, VP9 at 8/10/12-bit, H.264 at 8/10-bit, HEVC at 8/10/12-bit and AV1 at 8/10-bit. Doubles H.264 decode speed so it matches HEVC and AV1 decode, and adds 4:2:2 H.264 and HEVC decode support" + }, + "AI Management Processor": { + "Count": 1, + "Size": 1, + "Description": "AMP, new in RTX Blackwell: a fully programmable hardware scheduler that manages and context-switches concurrent AI and graphics workloads on the GPU without host round trips, so multiple models can be co-resident alongside a rendering workload" + } + } + }, + "MemorySubsystem": { + "SupportedMemoryTypes": ["Other"], + "MemoryTypes": [ + { + "Type": "Register File (per SM)", + "Size": 256, + "BankCount": 4, + "description": "256 KB of register file per SM, i.e. 65536 32-bit registers, split evenly across the 4 processing blocks at 16384 registers (64 KB) each. A single thread may use at most 255 registers and a single thread block at most 65536. Aggregate register file on RTX 5090 is 43520 KB across 170 SMs, matching 170 x 256 KB; the full GB203 has 21504 KB and the full GB205 12800 KB" + }, + { + "Type": "L1 Data Cache / Shared Memory (per SM)", + "Size": 128, + "BankCount": 32, + "description": "128 KB of unified L1 data cache, texture cache and shared memory per SM, the same physical capacity as Ada sm_89. Aggregate on RTX 5090 is 21760 KB across 170 SMs, matching 170 x 128 KB; the full GB203 has 10752 KB and the full GB205 6400 KB, both also 128 KB per SM. Of that array at most 100 KB is addressable as shared memory on compute capability 12.0: the carveout is configurable per kernel to 0, 8, 16, 32, 64 or 100 KB, the remainder serving as L1 and texture cache. A single thread block may address at most 99 KB because 1 KB of the partition is reserved for system use, and any block wanting more than 48 KB must use dynamic shared memory and opt in through cudaFuncSetAttribute with cudaFuncAttributeMaxDynamicSharedMemorySize. The texture-memory working set per SM ranges from 28 KB to 128 KB. Shared memory is banked 32 ways with 32-bit banks. Note that the CUDA 12.9 through 13.3 guides describe the unified array itself as 100 KB, conflating the shared-memory cap with the array size; the 128 KB recorded here is the whitepaper figure and is confirmed by dividing the published per-GPU L1 totals by the SM counts. sm_120 supports distributed shared memory across a thread block cluster" + }, + { + "Type": "L2 Cache", + "Size": 98304, + "maxBusWidth": 512, + "description": "96 MB of GPU-wide L2 on RTX 5090; the full GB202 die carries 128 MB. Backs all SMs and fronts the 16 GDDR7 controllers, and is the large on-die pool that path tracing in particular benefits from. RTX 5080 (full GB203) has 64 MB and the full GB205 48 MB. NVIDIA does not publish an L2 bandwidth figure, so none is recorded here" + }, + { + "Type": "GDDR7 (Device DRAM)", + "Size": 33554432, + "BankCount": 16, + "MaxMemoryBandwidth": 1792, + "maxBusWidth": 512, + "description": "32 GB of GDDR7 across a 512-bit interface at a 28 Gbps data rate, 1792 GB/s on RTX 5090; BankCount here is the count of 32-bit memory channels, not DRAM internal banks. GDDR7 uses PAM3 signalling, three levels carrying 1.5 bits per cycle, with an improved pin-encoding scheme for signal-to-noise. RTX 50-series boards support Enhanced CRC for RAS, and GDDR7 on RTX Blackwell supports single-bit error correction ECC with no performance cost when disabled, plus error detection and replay. RTX 5080 ships 16 GB at 30 Gbps on 256 bits for 960 GB/s. GDDR7 is not one of the enum values this schema offers for SupportedMemoryTypes, so that field records Other" + }, + { + "Type": "Constant Memory", + "Size": 64, + "description": "64 KB constant memory window per device context for compute capability 12.0, served through a read-only per-SM constant cache shared by all functional units, whose working set the CUDA guide gives as 8 KB per SM" + } + ] + }, + "KernelModel": { + "LLVMTarget": "nvptx64", + "LLVMTriple": "nvptx64-nvidia-cuda", + "LLVMFeatures": [ + { + "Name": "+sm_120", + "Description": "NVPTX subtarget feature selecting the consumer/workstation Blackwell sm_120 target ISA (compute capability 12.0)" + }, + { + "Name": "+sm_120a", + "Description": "Architecture-specific sm_120a variant exposing accelerated features that are not forward compatible with later architectures, notably mma with .f16 accumulator at shape .m16n8k16 for the FP8 .e4m3 and .e5m2 types, and mma plus mma.sp::ordered_metadata extended with the .e3m2, .e2m3 and .e2m1 FP6/FP4 types and the .kind, .block_scale and .scale_vec_size qualifiers" + }, + { + "Name": "+sm_120f", + "Description": "Family-specific sm_120f variant, portable across the sm_120 family rather than pinned to a single architecture as sm_120a is. Per the CUDA family-specific compatibility table, compute_120f is compatible with compute capability 12.0 and 12.1, and its feature set is a superset of baseline sm_120 and a subset of sm_120a" + }, + { + "Name": "+ptx87", + "Description": "PTX ISA 8.7, the first release that adds support for the sm_120 and sm_120a target architectures" + }, + { + "Name": "+ptx88", + "Description": "PTX ISA 8.8, required for the family-specific sm_120f target form" + } + ], + "SubUnits": [ + { + "Type": "CUDA", + "Version": "12.8" + }, + { + "Type": "PTX ISA", + "Version": "8.7" + }, + { + "Type": "OpenCL", + "Version": "3.0" + } + ] + }, + "SpecializedHardware": { + "RayTracingAccelerators": { + "Present": true, + "Name": "Fourth-Generation RT Core", + "Count": 170 + }, + "AiAccelerators": { + "Present": true, + "Name": "Fifth-Generation Tensor Core", + "Count": 680, + "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "Other"], + "Performance": { + "fp16TopsPerGPU": 419, + "int8TopsPerGPU": 838 + } + }, + "videoCodecs": { + "encoders": [ + { + "codec": "H.264" + }, + { + "codec": "H.265/HEVC" + }, + { + "codec": "AV1" + } + ], + "decoders": [ + { + "codec": "H.264" + }, + { + "codec": "H.265/HEVC" + }, + { + "codec": "AV1" + }, + { + "codec": "VP9" + } + ] + } + }, + "PowerEfficiency": { + "BaseClock": 2010, + "BoostClock": 2407, + "MaxTDP": 575, + "PowerStates": [ + { + "Name": "Maximum Performance", + "Description": "Full power state for GeForce RTX 5090: 2410 MHz boost per the product page, 2407 MHz in the architecture whitepaper's peak-rate table, over a 575 W total graphics power budget. Base clock is 2010 MHz. RTX 5080 boosts to 2617 MHz at 360 W" + }, + { + "Name": "Balanced", + "Description": "Reduced clock and voltage operating point selected by the driver when the workload or the power and thermal budget does not require peak clocks. RTX Blackwell adds low-latency sleep and accelerated frequency switching, and exploits GDDR7's fast-wakeup clocking for finer-grained memory power management" + }, + { + "Name": "Idle", + "Description": "Low power state when no kernels or rendering work are resident" + } + ], + "ClockGating": true, + "DynamicVoltageFrequencyScaling": true + }, + "PciExpress": { + "Version": "5.0", + "Lanes": 16, + "Bandwidth": 128 + }, + "MultiGpuSupport": { + "Technologies": ["None"], + "MaxGpus": 1 + } +} diff --git a/device_lib/nvidia-gpu-sm_90.json b/device_lib/nvidia-gpu-sm_90.json index f440b9bc..dfd3536b 100644 --- a/device_lib/nvidia-gpu-sm_90.json +++ b/device_lib/nvidia-gpu-sm_90.json @@ -1,274 +1,252 @@ { - "name": "Hopper", - "vendor": "NVIDIA", - "generation": 9, - "releaseYear": 2022, - "fabricationProcess": { - "processNode": 4, - "manufacturer": "TSMC", - "technology": "N4" + "Name": "Hopper", + "Vendor": "NVIDIA", + "Architecture": "sm_90", + "ReleaseYear": 2022, + "FabricationProcess": { + "ProcessNode": 5, + "Manufacturer": "TSMC", + "Technology": "4N" }, - "computeUnits": { - "name": "Streaming Multiprocessor (SM)", - "maxPerGPU": 144, - "subUnits": [ + "CoreSubsystem": { + "Name": "SM", + "ChipType": "Chip", + "CoreType": "SM", + "UnitTypes": { + "Chip": { + "Count": 1, + "Size": 132, + "Description": "GH100 die: 80 billion transistors on 814 mm2, built on the TSMC 4N process customized for NVIDIA. The primary configuration described by this file is H100 SXM5 (132 SMs, 16896 FP32 CUDA cores, 528 Tensor Cores, 80 GB HBM3, 700 W). The full GH100 die is larger than any shipping product: 8 GPCs / 72 TPCs / 144 SMs, 18432 FP32 cores, 576 Tensor Cores, 60 MB L2 and 6 HBM stacks. H100 PCIe ships 114 SMs (57 TPCs, 14592 FP32 cores, 456 Tensor Cores) with 80 GB HBM2e at 2 TB/s and a configurable 300-350 W; H100 NVL ships 132 SMs with 94 GB HBM3 at 3.9 TB/s and 350-400 W. Supports up to 7 MIG instances of 10 GB each (12 GB each on H100 NVL), ECC, and Confidential Computing. GH100 is a compute-only part: no display outputs and no RT cores", + "Memory": ["L2 Cache", "HBM3 (Device DRAM)", "Constant Memory"], + "Subunits": ["GPC", "Memory Controller", "HBM3 Stack", "NVLink 4 Link", "PCIe Controller", "NVDEC Engine", "NVJPEG Engine"] + }, + "GPC": { + "Count": 8, + "Size": 9, + "Description": "Graphics Processing Cluster. The full GH100 die has 8 GPCs of 9 TPCs each (72 TPCs). H100 SXM5 enables 66 of those 72 TPCs and H100 PCIe enables 57, so shipping parts do not populate every GPC uniformly; the Size of 9 is the full-die TPC complement per GPC", + "Subunits": ["TPC"] + }, + "TPC": { + "Count": 9, + "Size": 2, + "Description": "Texture Processing Cluster, containing 2 SMs. 72 TPCs on the full GH100 die, 66 enabled on H100 SXM5 (132 SMs) and 57 on H100 PCIe (114 SMs)", + "Subunits": ["SM"] + }, + "SM": { + "Count": 2, + "Size": 128, + "Description": "Hopper Streaming Multiprocessor, organized as 4 processing blocks (partitions) that each hold one warp scheduler, one dispatch unit, a 64 KB register file slice, one fourth-generation Tensor Core, 32 FP32 lanes, 16 INT32 lanes, 16 FP64 lanes, 4 SFUs, 8 LD/ST units and 1 texture unit. Compute capability 9.0 limits: 2048 resident threads (64 warps), 32 resident thread blocks, 255 registers per thread, 65536 32-bit registers per SM, 228 KB of shared memory per SM. Hopper adds thread block clusters (portable maximum 8 blocks, non-portable 16 on H100) with distributed shared memory across the cluster, and DPX dynamic-programming instructions at 128 operations per cycle per SM. FP32 throughput per SM per clock is 2x that of compute capability 8.0", + "Memory": ["Register File (per SM)", "L1 Data Cache / Shared Memory (per SM)"], + "Subunits": ["FP32 CUDA Core", "INT32 Core", "FP64 Core", "Tensor Core", "SFU", "LD/ST Unit", "Warp Scheduler", "Texture Unit", "Tensor Memory Accelerator"] + }, + "FP32 CUDA Core": { + "Count": 128, + "Size": 1, + "Description": "Single-precision FP32 lane, 32 per processing block. 16896 per H100 SXM5 (14592 on H100 PCIe, 18432 on the full GH100 die), giving 67 TFLOPS of peak non-tensor FP32 on H100 SXM5 and 51 TFLOPS on H100 PCIe" + }, + "INT32 Core": { + "Count": 64, + "Size": 1, + "Description": "32-bit integer lane, 16 per processing block, executing concurrently with the FP32 datapath" + }, + "FP64 Core": { + "Count": 64, + "Size": 1, + "Description": "Double-precision lane, 16 per processing block, giving an FP64:FP32 rate of 1:2. Peak non-tensor FP64 is 34 TFLOPS on H100 SXM5 and 26 TFLOPS on H100 PCIe; the Tensor Cores double this to 67 and 51 TFLOPS respectively" + }, + "Tensor Core": { + "Count": 4, + "Size": 1, + "Description": "Fourth-generation Tensor Core, one per processing block (528 per H100 SXM5, 456 per H100 PCIe, 576 on the full die). Supports FP64, TF32, BF16, FP16, FP8 (E4M3/E5M2) and INT8, plus 2:4 structured sparsity, and is driven by the new asynchronous warp-group MMA (wgmma) instructions. Together with the Transformer Engine, which dynamically chooses between FP8 and FP16 per layer, H100 SXM5 peaks at 989 TFLOPS TF32, 1979 TFLOPS BF16/FP16, 3958 TFLOPS FP8 and 3958 TOPS INT8 with sparsity, i.e. half those figures for dense operands" + }, + "SFU": { + "Count": 16, + "Size": 1, + "Description": "Special Function Unit, 4 per processing block, computing transcendentals such as sin, cos, exp2, log2, rcp and rsqrt" + }, + "LD/ST Unit": { + "Count": 32, + "Size": 1, + "Description": "Load/store unit, 8 per processing block, issuing accesses to shared, local and global memory through the unified L1 data cache" + }, + "Warp Scheduler": { + "Count": 4, + "Size": 1, + "Description": "One warp scheduler plus dispatch unit per processing block, each selecting among up to 16 resident warps of 32 threads for a 64-warp per-SM occupancy limit" + }, + "Texture Unit": { + "Count": 4, + "Size": 1, + "Description": "Texture/sampler unit, one per processing block, sharing the unified L1 data and texture cache" + }, + "Tensor Memory Accelerator": { + "Count": 1, + "Size": 1, + "Description": "TMA, new in Hopper: a per-SM asynchronous copy engine that moves 1D through 5D tensor tiles between global memory and shared memory in both directions from a single thread, generating the addressing itself and freeing warps for compute" + }, + "Memory Controller": { + "Count": 10, + "Size": 1, + "Description": "512-bit HBM memory controller. The full GH100 die has 12 controllers (6144-bit); H100 SXM5 and H100 PCIe enable 10 for a 5120-bit interface" + }, + "HBM3 Stack": { + "Count": 5, + "Size": 1, + "Description": "16 GB HBM3 stack, 5 stacks for 80 GB on H100 SXM5 at 3.35 TB/s. The full GH100 die supports 6 HBM3 or HBM2e stacks. H100 PCIe uses 5 HBM2e stacks (80 GB at 2 TB/s); H100 NVL exposes 94 GB of HBM3 per GPU at 3.9 TB/s", + "Memory": ["HBM3 (Device DRAM)"] + }, + "NVLink 4 Link": { + "Count": 18, + "Size": 1, + "Description": "Fourth-generation NVLink port, 50 GB/s bidirectional each, 18 ports for 900 GB/s total on H100 SXM5. An HGX/DGX H100 baseboard connects 8 GPUs through NVSwitch, and the NVLink Switch System extends the domain to 256 GPUs. H100 PCIe instead uses a 3-way NVLink bridge (600 GB/s) across 2 GPUs; H100 NVL pairs two boards at 600 GB/s" + }, + "PCIe Controller": { + "Count": 1, + "Size": 1, + "Description": "PCIe Gen 5 x16 host interface, 128 GB/s bidirectional, with SR-IOV and Confidential Computing support" + }, + "NVDEC Engine": { + "Count": 7, + "Size": 1, + "Description": "Fourth-generation NVDEC hardware video decoder, 7 per GPU. GH100 ships no NVENC encoder engines" + }, + "NVJPEG Engine": { + "Count": 7, + "Size": 1, + "Description": "Hardware JPEG decode engine, 7 per GPU, exposed through nvJPEG for data-loading pipelines" + } + } + }, + "MemorySubsystem": { + "SupportedMemoryTypes": ["HBM3", "HBM2e"], + "MemoryTypes": [ + { + "Type": "Register File (per SM)", + "Size": 256, + "BankCount": 4, + "description": "65536 32-bit registers (256 KB) per SM, split evenly across the 4 processing blocks. A single thread may use at most 255 registers. Aggregate register file on H100 SXM5 is 33 MB across 132 SMs" + }, { - "name": "CUDA Core", - "countPerComputeUnit": 128, - "description": "FP32 and INT32 processing core" + "Type": "L1 Data Cache / Shared Memory (per SM)", + "Size": 256, + "BankCount": 32, + "description": "256 KB of unified L1 data cache, texture cache and shared memory per SM, up 33% from 192 KB on Ampere. The shared-memory carveout is configurable to 0, 8, 16, 32, 64, 100, 132, 164, 196 or 228 KB per SM; a single thread block may address at most 227 KB because CUDA reserves 1 KB. Shared memory is banked 32 ways, and Hopper thread block clusters can address the shared memory of peer blocks as distributed shared memory" }, { - "name": "Tensor Core (4th gen)", - "countPerComputeUnit": 4, - "description": "Specialized core for matrix operations and AI acceleration" + "Type": "L2 Cache", + "Size": 51200, + "MaxMemoryBandwidth": 3350, + "maxBusWidth": 5120, + "description": "50 MB of GPU-wide L2 on both H100 SXM5 and H100 PCIe; the full GH100 die has 60 MB. Backs all SMs and fronts the HBM controllers, and supports L2 residency controls and compression" }, { - "name": "Special Function Unit (SFU)", - "countPerComputeUnit": 4, - "description": "Units handling transcendental instructions like sine, cosine, etc." + "Type": "HBM3 (Device DRAM)", + "Size": 83886080, + "BankCount": 5, + "MaxMemoryBandwidth": 3350, + "maxBusWidth": 5120, + "description": "80 GB of HBM3 across 5 stacks on a 5120-bit interface, 3.35 TB/s on H100 SXM5. H100 PCIe uses 80 GB of HBM2e on the same 5120-bit interface at 2 TB/s; H100 NVL exposes 94 GB of HBM3 per GPU at 3.9 TB/s. ECC is enabled on all SKUs" + }, + { + "Type": "Constant Memory", + "Size": 64, + "description": "64 KB constant memory window per device context for compute capability 9.0, served through a per-SM constant cache whose size NVIDIA does not publish" } ] }, - "shaderModel": { - "directX": "12 Ultimate", - "vulkan": "1.3", - "openGL": "4.6", - "openCL": "3.0", - "metal": "N/A" - }, - "memorySubsystem": { - "supportedMemoryTypes": ["HBM3"], - "maxMemoryBandwidth": 3350, - "maxMemorySize": 80, - "maxBusWidth": 5120, - "cacheHierarchy": [ + "KernelModel": { + "LLVMTarget": "nvptx64", + "LLVMTriple": "nvptx64-nvidia-cuda", + "LLVMFeatures": [ + { + "Name": "+sm_90", + "Description": "NVPTX subtarget feature selecting the Hopper sm_90 target ISA" + }, + { + "Name": "+sm_90a", + "Description": "Architecture-specific sm_90a variant exposing Hopper-only instructions such as asynchronous warp-group MMA (wgmma) and setmaxnreg. Code compiled for sm_90a is not forward compatible with later architectures" + }, + { + "Name": "+ptx78", + "Description": "PTX ISA 7.8 (CUDA 11.8), the first release able to target sm_90" + }, + { + "Name": "+ptx80", + "Description": "PTX ISA 8.0 (CUDA 12.0), the first release able to target sm_90a" + } + ], + "SubUnits": [ { - "level": "L1 Cache", - "sizePerUnit": 256, - "totalSize": 36864, - "description": "L1 cache per SM with shared memory" + "Type": "CUDA", + "Version": "11.8" }, { - "level": "L2 Cache", - "sizePerUnit": 0, - "totalSize": 51200, - "description": "Shared L2 cache across all SMs" + "Type": "PTX ISA", + "Version": "7.8" + }, + { + "Type": "OpenCL", + "Version": "3.0" } ] }, - "specializedHardware": { - "rayTracingAccelerators": { - "present": true, - "name": "RT Core (3rd gen)", - "countPerComputeUnit": 1, - "performance": { - "raysPerSecond": 200 - } + "SpecializedHardware": { + "RayTracingAccelerators": { + "Present": false }, - "aiAccelerators": { - "present": true, - "name": "Tensor Core (4th gen)", - "countPerComputeUnit": 4, - "supportedPrecisions": ["FP32", "FP16", "BF16", "TF32", "FP8", "INT8", "INT4"], - "performance": { + "AiAccelerators": { + "Present": true, + "Name": "Fourth-Generation Tensor Core", + "Count": 528, + "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "Other"], + "Performance": { "fp16TopsPerGPU": 1979, "int8TopsPerGPU": 3958 } }, "videoCodecs": { - "encoders": [ - { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 600, - "maxFPS": 240 - }, - { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 600, - "maxFPS": 240 - }, - { - "codec": "AV1", - "maxResolution": "8K", - "maxBitrate": 600, - "maxFPS": 240 - } - ], + "encoders": [], "decoders": [ { - "codec": "H.264", - "maxResolution": "8K", - "maxBitrate": 600, - "maxFPS": 240 + "codec": "H.264" }, { - "codec": "H.265/HEVC", - "maxResolution": "8K", - "maxBitrate": 600, - "maxFPS": 240 + "codec": "H.265/HEVC" }, { - "codec": "AV1", - "maxResolution": "8K", - "maxBitrate": 600, - "maxFPS": 240 - }, - { - "codec": "VP9", - "maxResolution": "8K", - "maxBitrate": 600, - "maxFPS": 240 + "codec": "VP9" } ] } }, - "powerEfficiency": { - "maxTDP": 700, - "powerStates": [ + "PowerEfficiency": { + "BoostClock": 1980, + "MaxTDP": 700, + "PowerStates": [ { - "name": "Maximum Performance", - "description": "Full power mode for compute-intensive workloads" + "Name": "Maximum Performance", + "Description": "Full power state for H100 SXM5, up to 700 W configurable TDP. NVIDIA does not publish a base clock for H100; the 1980 MHz boost clock recorded here is derived from the official 67 TFLOPS peak FP32 figure over 16896 FP32 lanes at 2 FLOP/clock. H100 PCIe is rated at 300-350 W configurable and H100 NVL at 350-400 W configurable" }, { - "name": "Balanced", - "description": "Balance between performance and power consumption" + "Name": "Balanced", + "Description": "Reduced clock and voltage operating point selected by the driver when the workload or the power/thermal budget does not require peak clocks" }, { - "name": "Energy Efficient", - "description": "Optimized for energy efficiency" + "Name": "Idle", + "Description": "Low power state when no kernels are resident" } ], - "clockGating": true, - "dynamicVoltageFrequencyScaling": true - }, - "displayOutputs": { - "maxDisplays": 0, - "maxResolution": "N/A", - "maxRefreshRate": 0, - "interfaces": [], - "hdr": false - }, - "pciExpress": { - "version": "5.0", - "lanes": 16, - "bandwidth": 128 - }, - "multiGpuSupport": { - "technologies": ["NVLink 4.0"], - "maxGpus": 8, - "interconnectBandwidth": 900 + "ClockGating": true, + "DynamicVoltageFrequencyScaling": true }, - "softwareFeatures": { - "upscalingTechnologies": [], - "meshShading": true, - "variableRateShading": true, - "samplerFeedback": true + "PciExpress": { + "Version": "5.0", + "Lanes": 16, + "Bandwidth": 128 }, - "additionalFeatures": { - "transformerEngine": { - "present": true, - "description": "Specialized acceleration for transformer models with FP8 precision" - }, - "confidentialComputing": { - "present": true, - "description": "Supports encrypted computation for sensitive data" - }, - "dynamicProgrammingAccelerator": { - "present": true, - "description": "Hardware accelerator for dynamic programming algorithms" - }, - "mig": { - "present": true, - "description": "Multi-Instance GPU technology for workload isolation", - "maxInstances": 7 - }, - "nvdec": { - "present": true, - "description": "Hardware video decoder" - }, - "nvenc": { - "present": true, - "description": "Hardware video encoder" - }, - "cudaCores": 18432, - "tensorCores": 576, - "clockSpeeds": { - "baseClockMHz": 1365, - "boostClockMHz": 1815 - }, - "fp64Performance": { - "peakTeraflops": 34, - "ratioToFP32": 0.5 - }, - "asynchronousCopyEngine": true, - "ecc": true, - "virtualization": { - "present": true, - "technologies": ["vGPU", "MIG"] - } - }, - "variants": [ - { - "name": "H100 SXM5", - "formFactor": "SXM5", - "tdp": 700, - "memorySize": 80, - "nvlinkPorts": 18 - }, - { - "name": "H100 PCIe", - "formFactor": "PCIe", - "tdp": 350, - "memorySize": 80, - "nvlinkPorts": 0 - }, - { - "name": "H100 NVL", - "formFactor": "PCIe", - "tdp": 350, - "memorySize": 94, - "nvlinkPorts": 8 - } - ], - "software": { - "drivers": ["NVIDIA Data Center Driver"], - "sdks": ["CUDA", "cuDNN", "DALI", "TensorRT"], - "frameworks": [ - "TensorFlow", - "PyTorch", - "JAX", - "MXNet", - "RAPIDS", - "Triton Inference Server" - ], - "compilers": ["NVCC", "LLVM/Clang", "OpenACC"] - }, - "inferencePerformance": { - "resnet50ImagesPerSecond": 294000, - "bert": { - "throughput": 49262, - "unit": "queries/second" - }, - "dlrm": { - "throughput": 54000, - "unit": "samples/second" - }, - "recommendationSystems": { - "throughput": 92000, - "unit": "samples/second" - } - }, - "trainingPerformance": { - "resnet50ImagesPerSecond": 43518, - "bert": { - "throughput": 14584, - "unit": "samples/second" - }, - "dlrm": { - "throughput": 14900, - "unit": "samples/second" - } + "MultiGpuSupport": { + "Technologies": ["NVLink"], + "MaxGpus": 256, + "InterconnectBandwidth": 900 } -} \ No newline at end of file +} diff --git a/device_lib/tenstorrent-npu-blackhole.json b/device_lib/tenstorrent-npu-blackhole.json index b18f397d..7196dfbe 100644 --- a/device_lib/tenstorrent-npu-blackhole.json +++ b/device_lib/tenstorrent-npu-blackhole.json @@ -1,7 +1,7 @@ { "Name": "Blackhole", "Vendor": "Other", - "Architecture": "Tensix++", + "Architecture": "blackhole", "ReleaseYear": 2024, "FabricationProcess": { "ProcessNode": 6, @@ -10,72 +10,85 @@ }, "CoreSubsystem": { "Name": "Tensix", - "SubUnits": [ - { - "Type": "Chip", + "ChipType": "Chip", + "CoreType": "Tensix Core", + "UnitTypes": { + "Chip": { "Count": 1, "Size": 120, - "Description": "Full Blackhole die containing 120 active Tensix cores arranged in a 2D mesh with NoC interconnect", - "Memory": ["GDDR6", "SRAM"], - "SubunitType": "Tensix Core" + "Description": "One Blackhole ASIC. Two NoCs join the tiles into a 2D torus, usually drawn as a 2D grid. The die carries 140 Tensix compute tiles in a 14x10 grid, 24 GDDR6 DRAM tiles, 4 L2CPU tiles, 14 Ethernet tiles, 2 PCI Express tiles, 1 ARC tile and 1 Security tile. Tensix tiles are harvested by fusing off whole columns: p100a fuses two columns and exposes 120. p150a/p150b originally exposed all 140, but Tenstorrent reduced them to 120 for cards shipping from January 2026 and for existing cards updated to firmware 19.5.0 or later, which also moved the published SRAM figure from 210 MB to 180 MB and the BLOCKFP8 figure from 774 to 664 TFLOPS. 120 is therefore the count Tenstorrent now publishes for all three cards and the count used throughout this file", + "Memory": ["GDDR6 (Off-chip DRAM)", "SRAM (L1 per Tensix Core)"], + "Subunits": ["Tensix Core", "L2CPU Cluster", "DRAM Controller", "Ethernet Controller", "PCIe Controller", "ARC Management Core", "Security Core"] }, - { - "Type": "Tensix Core", + "Tensix Core": { "Count": 120, "Size": 5, - "Description": "Each Tensix core contains 5 Baby RISC-V CPUs (Data Movement 0, Data Movement 1, Unpack, Math, Pack), a matrix unit (FPU), a vector unit (SFPU), pack/unpack units, 2 NoC router interfaces, and 1.5 MB of local SRAM", - "Memory": ["SRAM (L1)"], - "SubunitType": "Baby RISC-V" + "Description": "Compute tile (worker tile) holding 1536 KiB of L1 SRAM, 5 Baby RISC-V cores, 2 NoC connections, a NoC overlay coprocessor, and the Tensix coprocessor itself: 2 unpackers, 1 matrix unit (FPU), 1 vector unit (SFPU), 1 scalar unit (ThCon) for integer scalar work and 128-bit L1 accesses including atomics, and 4 packers. Compared with Wormhole the Blackhole tile gains clock speed, more L1 bandwidth, an enhanced RISC-V instruction set and extra SFPU instructions; L1 capacity per tile is unchanged. 140 such tiles exist on the die and 120 are exposed on p100a, p150a and p150b", + "Memory": ["SRAM (L1 per Tensix Core)"], + "Subunits": ["Baby RISC-V", "Matrix Unit (FPU)", "Vector Unit (SFPU)"] }, - { - "Type": "Baby RISC-V", + "Baby RISC-V": { "Count": 5, "Size": 1, - "Description": "Small in-order RISC-V cores within each Tensix responsible for instruction dispatch and control. Roles: Data Movement 0 (reader), Data Movement 1 (writer), Unpack (unpacker control), Math (compute engine dispatch), Pack (packer control)" + "Description": "Small in-order 32-bit cores named RISCV B, RISCV T0, RISCV T1, RISCV T2 and RISCV NC. They implement RV32IM plus all of Zicsr, Zaamo, Zba and Zbb, plus parts of Zicntr, F and Zfh; RISCV T2 additionally implements part of V. Instructions are fetched only from L1, through a small per-core L0 instruction cache. The cores are not meant to compute themselves: they oversee the NoC and dispatch instructions into the Tensix coprocessor. TT-Metalium's usual division of labour is two cores driving the NoC (Data Movement 0 and Data Movement 1) and three driving the coprocessor (Unpack, Math and Pack); Tenstorrent's ISA documentation describes this as one natural assignment rather than a fixed property of the hardware", + "Memory": ["Instruction Cache (per Baby RISC-V)", "Private RISC-V Data RAM"] }, - { - "Type": "Matrix Unit (FPU)", + "Matrix Unit (FPU)": { "Count": 1, "Size": 1, - "Description": "Dedicated matrix multiply unit per Tensix core, operates on 32x32 tiles natively. Supports FP8, FP16, BF16, BLOCKFP2/4/8, TF32, and INT8 precisions" + "Description": "Per-Tensix matrix engine performing 8x16 by 16x16 multiply-accumulate, 4096 multiply-adds per cycle. It multiplies low-precision (<= 19-bit) source matrices held in SrcA/SrcB and accumulates into Dst at 16- or 32-bit precision, so FP32 is an accumulator format rather than an input format. Native source formats are TF32, BF16, FP16 and INT8; BFP8/BFP4/BFP2 block-float tiles and FP8 are decompressed by the unpackers on the way in, and higher effective precision is reached by running 2-4 fidelity phases. Tenstorrent's tt-metal GEMM_FLOPS report gives roughly 5.4, 2.7 and 1.35 TFLOP/s per matrix engine at the 1.35 GHz AI clock for LoFi (1 fidelity phase), HiFi2 and HiFi4. Across 120 engines that is 664, 332 and 166 TFLOPS, and the LoFi figure reproduces the only throughput number Tenstorrent publishes for these cards: 664 TFLOPS BLOCKFP8 (774 on the pre-2026 140-core p150). No FP16, BF16 or INT8 figure is published; the values recorded under SpecializedHardware.AiAccelerators.Performance are derived from these per-engine rates and are flagged as such there", + "Memory": ["SRAM (L1 per Tensix Core)"] }, - { - "Type": "Vector Unit (SFPU)", + "Vector Unit (SFPU)": { "Count": 1, "Size": 32, - "Description": "32-wide vector processing unit per Tensix core for element-wise operations such as activation functions, reductions, and transcendentals" + "Description": "Per-Tensix SIMD unit, 32 lanes of 32 bits wide with 8 general-purpose vector registers (LRegs), used for element-wise unary and binary operations: activations, comparisons, reductions and transcendentals, including FP32 multiply-accumulate. Blackhole adds several SFPU instructions over Wormhole. Programmed through Tenstorrent's SFPI C++ intrinsics in the riscv-tt-elf GCC toolchain, which lower to the Xtttensixbh vendor extension", + "Memory": ["SRAM (L1 per Tensix Core)"] }, - { - "Type": "Big RISC-V (SiFive x280)", - "Count": 16, + "L2CPU Cluster": { + "Count": 4, + "Size": 4, + "Description": "L2CPU tile: a cache-coherent cluster of four SiFive X280 CPUs sharing 2 MiB of L3. Each cluster has a direct connection to one local DRAM tile at a fixed place in its address space plus 256 TLB windows onto the NoC, which let it reach any other tile including the other DRAM tiles and the host through the PCI Express tile. All 4 clusters are present and enabled on p100a as well as p150a/p150b, though on p100a GDDR6 harvesting can leave one or two clusters without directly attached DRAM. These are general-purpose 64-bit RISC-V CPUs on the die, not TT-Metalium kernel targets: tt-metal's Blackhole SoC descriptor does not list them and every tt-metal kernel build is 32-bit", + "Memory": ["L3 Cache (per L2CPU Cluster)"], + "Subunits": ["Big RISC-V (SiFive X280)"] + }, + "Big RISC-V (SiFive X280)": { + "Count": 4, "Size": 1, - "Description": "16 SiFive Intelligence x280 64-bit dual-issue in-order CPU cores arranged in 4 clusters of 4. Capable of running Linux and acting as on-chip host processor, eliminating need for external host CPU" + "Description": "SiFive Intelligence X280, a 64-bit dual-issue in-order Linux-capable RISC-V core with RISC-V Vector 1.0. Four per L2CPU cluster and 16 per ASIC, which is the 'Big RISC-V Cores: 16' figure on Tenstorrent's spec sheet for all three cards. They are unrelated to Tenstorrent's own Ascalon cores. Each has private 32 KiB L1 instruction, 32 KiB L1 data and 128 KiB L2 caches. Tenstorrent does not document their intended workload beyond general-purpose on-die compute and host/management duties, and whether shipping firmware boots Linux on them is not stated publicly", + "Memory": ["L1 Instruction Cache (per SiFive X280)", "L1 Data Cache (per SiFive X280)", "L2 Cache (per SiFive X280)"] }, - { - "Type": "DRAM Controller", + "DRAM Controller": { "Count": 8, - "Size": 1, - "Description": "8 GDDR6 DRAM controllers, each managing 4 GB for a total of 32 GB (p150a/p150b) or 28 GB (p100a)" + "Size": 3, + "Description": "The 24 GDDR6 DRAM tiles on the die are arranged in 8 groups of 3, and each group exposes the same 4 GiB of GDDR6 on all three of its tiles, so 32 GiB per ASIC. The three tiles in a group are separate NoC endpoints for one memory, which improves NoC connectivity. p150a and p150b keep all 8 groups for 32 GB at 512 GB/s; p100a fuses one group off, leaving 21 DRAM tiles for 28 GB at 448 GB/s. Tenstorrent does not publish a discrete memory-controller count separate from these DRAM tiles", + "Memory": ["GDDR6 (Off-chip DRAM)"] }, - { - "Type": "Ethernet Controller", + "Ethernet Controller": { "Count": 14, + "Size": 2, + "Description": "14 Ethernet tiles per ASIC, each carrying one bidirectional 400 GbE link with its own MAC/PCS/PHY presented as 3 TX and 3 RX queues, 512 KiB of L1, and 2 Baby RISC-V cores. Compared with Wormhole these gained a second RISC-V core, larger L1, higher clock and 400 GbE in place of 100 GbE. A pair of tiles feeds each QSFP-DD port, giving 800 GbE per port. p150a and p150b enable 12 tiles of which 8 are wired to the card's 4 QSFP-DD 800G ports, for 3.2 Tbps (400 GB/s each way) of chip-to-chip bandwidth; the other 4 are usable for their RISC-V cores and L1 but cannot move Ethernet packets. p100a has no Ethernet tiles available and no external interconnect. Each PCI Express x16 link in use costs four connectable Ethernet tiles, which is why the full ASIC caps at 12 connected", + "Memory": ["SRAM (L1 per Ethernet Core)"] + }, + "PCIe Controller": { + "Count": 2, "Size": 1, - "Description": "14 Ethernet cores on die (8 active on p150), each supporting 400 Gbps for chip-to-chip interconnect via QSFP-DD ports" + "Description": "Two PCI Express 5.0 x16 tiles, each providing a host interface including the address-translation TLBs the host driver uses to reach NoC endpoints. In current products exactly one is in use and the other is permanently idle", + "Memory": ["GDDR6 (Off-chip DRAM)"] }, - { - "Type": "PCIe Controller", + "ARC Management Core": { "Count": 1, "Size": 1, - "Description": "PCIe 5.0 x16 host interface controller" + "Description": "ARC tile running chip and board management firmware: power and clock management via the PLLs, voltage and thermal telemetry, reset control of the Tensix, Ethernet and L2CPU tiles, and firmware update. It runs no customer workload and dispatches none", + "Memory": ["GDDR6 (Off-chip DRAM)"] }, - { - "Type": "ARC Management Core", + "Security Core": { "Count": 1, "Size": 1, - "Description": "System management core for power management, thermal monitoring, and firmware control" + "Description": "Security tile, paired with the ARC tile for chip and board management. Like the ARC tile it executes no customer workload and is not involved in dispatching one", + "Memory": ["GDDR6 (Off-chip DRAM)"] } - ] + } }, "MemorySubsystem": { "SupportedMemoryTypes": ["GDDR6"], @@ -83,8 +96,12 @@ { "Type": "SRAM (L1 per Tensix Core)", "Size": 1536, - "BankCount": 1, - "description": "1.5 MB of local SRAM per Tensix core. Not a cache — no automatic caching, eviction, or coherence. Data must be explicitly managed via NoC transfers. Used for kernel code, circular buffers, tile storage, and inter-core data exchange. Total on-chip SRAM: 180 MB across 120 cores" + "description": "1536 KiB of software-managed scratchpad SRAM per Tensix tile, mapped from MEM_L1_BASE at 0x0000_0000 to 0x0017_FFFF. It is not a cache: there is no automatic fill, eviction or coherence, and data is placed explicitly by NoC transfers. It holds kernel code, circular buffers, tiles and inter-core data; TT-Metalium reports roughly 1464 KB of it as usable after reserved regions. Blackhole matches Wormhole on capacity but has more L1 bandwidth and adds a small write-through L0 data cache in front of it. Tenstorrent publishes 180 MB of on-chip SRAM per card, which is 120 tiles times 1.5 MB; the full 140-tile die holds 210 MB, the figure published for p150 cards before January 2026" + }, + { + "Type": "SRAM (L1 per Ethernet Core)", + "Size": 512, + "description": "512 KiB of L1 SRAM in each of the 14 Ethernet tiles, holding the Ethernet firmware, the TX/RX queues and chip-to-chip staging buffers, and addressable over the NoC exactly like Tensix L1. Twice the 256 KiB found in a Wormhole Ethernet tile. 7 MiB per ASIC across all 14 tiles" }, { "Type": "GDDR6 (Off-chip DRAM)", @@ -92,47 +109,84 @@ "BankCount": 8, "MaxMemoryBandwidth": 512, "maxBusWidth": 256, - "description": "32 GB GDDR6 at 16 GT/s across 8 DRAM controllers (p150a/p150b). 28 GB on p100a (7 controllers active with 448 GB/s). Accessed explicitly by Tensix cores via NoC; no hardware caching layer" + "description": "32 GiB of GDDR6 on p150a and p150b, exposed as 8 independent 4 GiB groups reached through 24 DRAM tiles on the NoC. At 16 GT/s over a 256-bit bus that is 512 GB/s. p100a fuses one group off for 28 GB, and its 448 GB/s implies an effective 224-bit bus. Each L2CPU cluster also has a direct path to one local DRAM tile outside the NoC. DRAM is not cached for Tensix cores - they stage data into L1 explicitly over the NoC" }, { "Type": "Instruction Cache (per Baby RISC-V)", - "Size": 2, - "BankCount": 1, - "description": "Small 0.5–2 KiB per-core instruction cache (approximately 128–512 instructions) to reduce repeated SRAM fetches during kernel execution" + "description": "Baby RISC-V kernel code lives in L1 and can only be executed from L1; each core fetches through a small L0 instruction cache in front of it. On the T cores this cache can fuse up to four adjacent .ttinsn instructions into one 64/96/128-bit instruction issued in a single cycle, otherwise it sustains one 32-bit instruction per cycle. Tenstorrent's TT-Metalium documentation describes it only as roughly 128-512 instructions, i.e. 0.5-2 KiB; no exact per-core Blackhole table is published, so Size is omitted here rather than guessed. Separately, each Baby RISC-V has a tiny non-coherent L0 data cache of 64 bytes (4 lines of 16 bytes), flushed by any fence or atomic" }, { - "Type": "Private RISC-V Memory", - "Size": 4, - "BankCount": 1, - "description": "Small private memory region per Baby RISC-V core for stack and local variables. Isolated per core; shared SRAM is accessible to all cores at identical addresses" + "Type": "Private RISC-V Data RAM", + "Size": 8, + "description": "Private data RAM per Baby RISC-V core with 2-cycle load latency, used for stack and thread-local storage: 8 KiB on RISCV B and RISCV NC, 4 KiB on each of RISCV T0, T1 and T2. Blackhole doubles the Wormhole sizes and adds a second, slower mapping at 0xFFB1_4000-0xFFB1_DFFF that lets the host, a debugger or another core reach it, alongside the fast core-private MEM_LOCAL_BASE mapping. Contrast L1, which all five cores see at the same addresses" + }, + { + "Type": "L1 Instruction Cache (per SiFive X280)", + "Size": 32, + "description": "32 KiB private L1 instruction cache in each of the 16 SiFive X280 cores. Virtually indexed and physically tagged" + }, + { + "Type": "L1 Data Cache (per SiFive X280)", + "Size": 32, + "description": "32 KiB private L1 data cache in each of the 16 SiFive X280 cores. Virtually indexed and physically tagged" + }, + { + "Type": "L2 Cache (per SiFive X280)", + "Size": 128, + "description": "128 KiB of private L2 cache per SiFive X280 core, with hardware prefetchers that firmware or the host configures after the cluster leaves reset" + }, + { + "Type": "L3 Cache (per L2CPU Cluster)", + "Size": 2048, + "description": "2 MiB of L3 cache shared by the four SiFive X280 cores of one L2CPU cluster, 8 MiB per ASIC across the 4 clusters. The same array can be configured as uncached scratchpad instead; Tenstorrent's recommended setup is 2 MiB of cache and no scratchpad" } ] }, "KernelModel": { - "LLVMTarget": "riscv64", - "LLVMTriple": "riscv64-unknown-elf", + "LLVMTarget": "riscv32", + "LLVMTriple": "riscv32-unknown-elf", "LLVMFeatures": [ { - "Name": "rv64gc", - "Description": "RISC-V 64-bit base integer ISA with general extensions (M, A, F, D, C) for Big RISC-V cores" + "Name": "m", + "Description": "RISC-V integer multiply/divide extension. The Tensix and Ethernet Baby RISC-V cores are 32-bit and build against RV32IM with the ilp32 ABI, so there is no C (compressed) or D extension and floating point is soft-float even though the hardware implements part of F and Zfh. The 16 SiFive X280 cores are 64-bit, but tt-metal never compiles kernels for them, so the kernel model described here is the 32-bit Baby RISC-V one" }, { - "Name": "v (Vector Extension)", - "Description": "RISC-V Vector extension support on SiFive x280 Big RISC-V cores" + "Name": "zaamo", + "Description": "RISC-V atomic memory operation instructions, new on Blackhole relative to Wormhole. Present in the Blackhole core definitions tt-metal selects with -mcpu=tt-bh and -mcpu=tt-bh-tensix, whose full ISA strings are rv32im_zaamo_zba_zbb and rv32im_zaamo_zba_zbb_xtttensixbh" + }, + { + "Name": "zba", + "Description": "RISC-V address-generation bit-manipulation extension, new on Blackhole relative to Wormhole" + }, + { + "Name": "zbb", + "Description": "RISC-V basic bit-manipulation extension, new on Blackhole relative to Wormhole" + }, + { + "Name": "xtttensixbh", + "Description": "Tenstorrent's Blackhole Tensix vendor extension, carrying the Tensix coprocessor instructions and the 32-lane SFPU with its 8 LRegs. Kernels are built with the SFPI riscv-tt-elf GCC toolchain, which emits elf32-littleriscv; tt-metal selects -mcpu=tt-bh-tensix for the compute cores and -mcpu=tt-bh for the rest. There is no upstream LLVM equivalent, so LLVM can only target the plain RV32IM scalar core" + }, + { + "Name": "zve32f", + "Description": "RISC-V embedded vector subset with 32-bit float elements. tt-metal builds one Blackhole compute core - the pack TRISC, RISCV T2, the only Baby RISC-V that implements part of V - with -march=rv32im_zmmul_zaamo_zba_zbb_xtttensixbh_zve32f. It is the single place the build uses -march rather than -mcpu, and it is a Zve32f subset, not full RVV" } ], "SubUnits": [ { "Type": "TT-Metalium", - "Version": "1.0" + "Version": "0.76.0" }, { "Type": "TT-NN", - "Version": "1.0" + "Version": "0.76.0" }, { "Type": "TT-Forge", - "Version": "1.0" + "Version": "1.4.0" + }, + { + "Type": "SFPI", + "Version": "7.72.0" } ] }, @@ -144,10 +198,10 @@ "Present": true, "Name": "Tensix Matrix Engine (FPU)", "Count": 120, - "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "INT4", "Other"], + "SupportedPrecisions": ["FP16", "BF16", "INT8", "Other"], "Performance": { "fp16TopsPerGPU": 332, - "int8TopsPerGPU": 664 + "int8TopsPerGPU": 166 } }, "videoCodecs": { @@ -156,17 +210,17 @@ } }, "PowerEfficiency": { - "BaseClock": 1000, + "BaseClock": 1350, "BoostClock": 1350, "MaxTDP": 300, "PowerStates": [ { "Name": "Active", - "Description": "Full power state with all Tensix cores and DRAM controllers active, up to 300W TBP" + "Description": "Full power state with all enabled Tensix, Ethernet, DRAM and L2CPU tiles active. Tenstorrent publishes a single AI clock of 'up to 1.35 GHz' rather than a base/boost pair, so both clock fields carry that figure and the ARC core scales below it under power or thermal limits. Total board power is up to 300 W on p100a, p150a and p150b alike, drawn through one 12+4-pin 12V-2x6 connector" }, { "Name": "Idle", - "Description": "Reduced power state when no workloads are running" + "Description": "Reduced power state when no kernels are dispatched; the ARC management core lowers Tensix clocks and voltage through the PLLs while keeping the PCIe and Ethernet links up" } ], "ClockGating": true, @@ -175,11 +229,10 @@ "PciExpress": { "Version": "5.0", "Lanes": 16, - "Bandwidth": 63 + "Bandwidth": 126 }, "MultiGpuSupport": { "Technologies": ["Other"], - "MaxGpus": 32, "InterconnectBandwidth": 400 } } diff --git a/device_lib/tenstorrent-npu-wormhole.json b/device_lib/tenstorrent-npu-wormhole.json index 73da46cf..90504edf 100644 --- a/device_lib/tenstorrent-npu-wormhole.json +++ b/device_lib/tenstorrent-npu-wormhole.json @@ -1,7 +1,7 @@ { "Name": "Wormhole", "Vendor": "Other", - "Architecture": "Tensix+", + "Architecture": "wormhole_b0", "ReleaseYear": 2023, "FabricationProcess": { "ProcessNode": 12, @@ -10,95 +10,96 @@ }, "CoreSubsystem": { "Name": "Tensix", - "SubUnits": [ - { - "Type": "Chip", + "ChipType": "Chip", + "CoreType": "Tensix Core", + "UnitTypes": { + "Chip": { "Count": 1, "Size": 80, - "Description": "Single Wormhole die (670 mm²) containing 80 Tensix cores arranged in a 10x12 2D grid with NoC interconnect. Of the grid tiles, 80 are Tensix compute cores and the remainder are DRAM, Ethernet, PCIe, and ARC management tiles", - "Memory": ["GDDR6", "SRAM"], - "SubunitType": "Tensix Core" + "Description": "One Wormhole B0 ASIC: a ~670 mm2 GlobalFoundries 12 nm die whose two NoCs form a 2D torus over a 10x12 grid of tiles. The grid holds 80 Tensix compute tiles, 18 GDDR6 DRAM tiles, 16 Ethernet tiles, 1 PCI Express tile and 1 ARC management tile. Tensix tiles are harvested per product: Tenstorrent's spec sheet lists 72 enabled on n150 and 64 per ASIC on n300 (128 across the card's two ASICs); 80 is the full unharvested die and is the count used throughout this file", + "Memory": ["GDDR6 (Off-chip DRAM)", "SRAM (L1 per Tensix Core)"], + "Subunits": ["Tensix Core", "DRAM Controller", "Ethernet Controller", "PCIe Controller", "ARC Management Core"] }, - { - "Type": "Tensix Core", + "Tensix Core": { "Count": 80, "Size": 5, - "Description": "Each Tensix core contains 5 Baby RISC-V CPUs (Data Movement 0, Data Movement 1, Unpack, Math, Pack), a matrix unit (FPU), a vector unit (SFPU), pack/unpack units, 2 NoC router interfaces, and 1.5 MB of local SRAM. n150 cards expose 72 cores; n300 cards expose 64 per die (128 total across 2 dies)", - "Memory": ["SRAM (L1)"], - "SubunitType": "Baby RISC-V" + "Description": "Compute tile (worker tile) holding 5 Baby RISC-V cores, a matrix unit (FPU), a vector unit (SFPU), packers and unpackers, 2 NoC router interfaces, and 1464 KiB of local L1 SRAM (marketed by Tenstorrent as 1.5 MB). Of the 80 tiles on the die, n150 cards expose 72 and n300 cards expose 64 per ASIC (128 per card); the rest are fused off", + "Memory": ["SRAM (L1 per Tensix Core)"], + "Subunits": ["Baby RISC-V", "Matrix Unit (FPU)", "Vector Unit (SFPU)"] }, - { - "Type": "Baby RISC-V", + "Baby RISC-V": { "Count": 5, "Size": 1, - "Description": "Small in-order RISC-V cores within each Tensix responsible for instruction dispatch and control. Roles: Data Movement 0 (reader/RISC0), Data Movement 1 (writer/RISC1), Unpack (RISC2), Math (RISC3), Pack (RISC4). These cores operate concurrently enabling native pipelining and data movement/compute overlap" + "Description": "Small in-order RV32IM cores (RV32I plus hardware multiply and divide; no compressed, atomic or hardware floating-point extensions). They mostly dispatch instructions to the Tensix coprocessor and drive the NoC rather than computing directly. Roles: RISCV B (BRISC, control, exposed by TT-Metalium as Data Movement 0), RISCV NC (NCRISC, network control, Data Movement 1) and RISCV T0/T1/T2 (the TRISCs) driving Unpack, Math and Pack. All five run concurrently, so data movement and compute overlap natively", + "Memory": ["Instruction Cache (per Baby RISC-V)", "Private RISC-V Data RAM"] }, - { - "Type": "Matrix Unit (FPU)", + "Matrix Unit (FPU)": { "Count": 1, "Size": 1, - "Description": "Dedicated matrix multiply unit per Tensix core, operates on 32x32 tiles natively. Supports FP8, FP16, BF16, BLOCKFP2/4/8, TF32, and INT8 precisions" + "Description": "Per-Tensix matrix engine. It multiplies low-precision (<= 19-bit) source matrices held in SrcA/SrcB and accumulates into Dst at 16- or 32-bit precision. Native source formats are TF32, BF16, FP16 and INT8; BFP8/BFP4/BFP2 block-float tiles (8/4/2 bits per datum with a shared 8-bit exponent per 16 datums) are decompressed by the unpackers on the way in, and higher effective precision is reached by running 2-4 fidelity phases. Tenstorrent's ISA documentation gives a theoretical peak of 4.096 TFLOP/s per matrix unit at the 1 GHz AI clock for MVMUL/DOTPV at 1 fidelity phase, halving as phases are added. Tenstorrent's published card figures are 262 TFLOPS FP8, 148 TFLOPS BFP8 and 74 TFLOPS FP16 for n150, and 466 / 262 / 131 TFLOPS for n300; no INT8 TOPS figure is published" }, - { - "Type": "Vector Unit (SFPU)", + "Vector Unit (SFPU)": { "Count": 1, "Size": 32, - "Description": "32-wide vector processing unit per Tensix core for element-wise unary and binary operations such as activation functions, reductions, and transcendentals" + "Description": "Per-Tensix SIMD unit, 32 lanes of 32 bits wide with 8 general-purpose vector registers (LRegs), used for element-wise unary and binary operations: activations, comparisons, reductions and transcendentals. Programmed through Tenstorrent's SFPI C++ intrinsics in the riscv-tt-elf GCC toolchain" }, - { - "Type": "DRAM Controller", + "DRAM Controller": { "Count": 6, "Size": 1, - "Description": "6 GDDR6 DRAM controllers, each managing 2 GB (2 channels of 1 GB each) for a total of 12 GB per die. Multiple DRAM tiles on the NoC can map to the same controller for improved connectivity" + "Description": "The 18 GDDR6 DRAM tiles on the die are arranged in 6 groups of 3 tiles, and each group fronts 2 GDDR6 channels of 1 GiB, so 2 GiB per group and 12 GiB per ASIC. The three tiles in a group are separate NoC endpoints for the same memory, which improves NoC connectivity. Each of the 12 channels is 16 bits wide at 12 GT/s, giving 24 GB/s per channel and 288 GB/s per ASIC" }, - { - "Type": "Ethernet Controller", + "Ethernet Controller": { "Count": 16, "Size": 1, - "Description": "16 Ethernet cores on die, each supporting 100 Gbps for chip-to-chip interconnect, providing 1.6 Tbps total bisection bandwidth for scale-out" + "Description": "16 Ethernet tiles per ASIC, each carrying one 100 GbE link and containing 256 KiB of L1 and one RV32IM Baby RISC-V core. Aggregate chip-to-chip bandwidth is 1.6 Tbps (200 GB/s) per ASIC. On an n300 card two of these links join the two ASICs on-card at 400 GbE; each card also exposes 2x QSFP-DD 400GbE ports and 2x Warp 100 Bridge notches for card-to-card meshes", + "Memory": ["SRAM (L1 per Ethernet Core)"] }, - { - "Type": "PCIe Controller", + "PCIe Controller": { "Count": 1, "Size": 1, - "Description": "PCIe 4.0 x16 host interface controller" + "Description": "Single PCI Express tile providing the PCIe 4.0 x16 host interface, including the address-translation TLBs the host driver uses to reach NoC endpoints. On an n300 only the PCIe-attached ASIC is visible to the host; the second ASIC is reached over the on-card Ethernet link" }, - { - "Type": "ARC Management Core", + "ARC Management Core": { "Count": 1, "Size": 1, - "Description": "System management core for power management, thermal monitoring, and firmware control" + "Description": "ARC tile running the chip and board management firmware: power and clock management, voltage/thermal telemetry, reset control of the Tensix and Ethernet tiles, and firmware update" } - ] + } }, "MemorySubsystem": { "SupportedMemoryTypes": ["GDDR6"], "MemoryTypes": [ { "Type": "SRAM (L1 per Tensix Core)", - "Size": 1536, - "BankCount": 1, - "description": "1.5 MB of local SRAM per Tensix core. Not a cache — no automatic caching, eviction, or coherence. Data must be explicitly managed via NoC transfers. Used for kernel code, circular buffers, tile storage, and inter-core data exchange. Total on-chip SRAM: 120 MB across 80 cores per die" + "Size": 1464, + "BankCount": 16, + "MaxMemoryBandwidth": 256, + "description": "1464 KiB (1,499,136 bytes) of software-managed scratchpad SRAM per Tensix tile, organised as 16 banks of 91.5 KiB; each bank sustains one 128-bit read or one 128-bit write per cycle, which works out to 256 B/cycle, i.e. the 256 GB/s figure above at the 1 GHz AI clock (derived from the bank organisation, not a vendor-published number). It is not a cache: there is no automatic fill, eviction or coherence, and data is placed explicitly by NoC transfers. It holds kernel code, circular buffers, tiles and inter-core data. Tenstorrent markets it as 1.5 MB per Tensix core and 120 MB per ASIC; the exact totals are 114.4 MiB over the full 80 tiles, 103.0 MiB on n150 (72 tiles) and 91.5 MiB per ASIC on n300 (64 tiles)" + }, + { + "Type": "SRAM (L1 per Ethernet Core)", + "Size": 256, + "description": "256 KiB of L1 SRAM in each of the 16 Ethernet tiles, holding the Ethernet firmware, the transmit/receive queues and chip-to-chip staging buffers. 4 MiB per ASIC in total. This memory is addressable over the NoC just like Tensix L1" }, { "Type": "GDDR6 (Off-chip DRAM)", "Size": 12582912, - "BankCount": 6, + "BankCount": 12, "MaxMemoryBandwidth": 288, "maxBusWidth": 192, - "description": "12 GB GDDR6 at 12 GT/s across 6 DRAM controllers per die (2 channels of 1 GB each per controller). n150 cards have 12 GB total, n300 cards have 24 GB total (12 GB per die). Accessed explicitly by Tensix cores via NoC" + "description": "12 GiB of GDDR6 per ASIC, exposed as 12 independent 1 GiB channels reached through 18 DRAM tiles on the NoC. Each channel is 16 bits wide at 12 GT/s for 24 GB/s, so 192 bits and 288 GB/s in aggregate; measured throughput is roughly 92% of theoretical. n150 cards have 12 GB at 288 GB/s, n300 cards 24 GB at 576 GB/s across their two ASICs. DRAM is not cached - Tensix cores stage data into L1 explicitly over the NoC" }, { "Type": "Instruction Cache (per Baby RISC-V)", "Size": 2, "BankCount": 1, - "description": "Small 0.5–2 KiB per-core instruction cache (approximately 128–512 instructions) to reduce repeated SRAM fetches during kernel execution" + "description": "Kernel code lives in L1 and is fetched through a small per-core instruction cache: 2 KiB each on RISCV B, T0, T1 and T2. RISCV NC is the exception, with 16 KiB of local instruction RAM backed by a 0.5 KiB cache" }, { - "Type": "Private RISC-V Memory", + "Type": "Private RISC-V Data RAM", "Size": 4, "BankCount": 1, - "description": "Small private memory region per Baby RISC-V core for stack and local variables. Isolated per core; shared SRAM is accessible to all cores at identical addresses" + "description": "Private data RAM per Baby RISC-V core with 2-cycle load latency, used for stack and thread-local storage: 4 KiB on RISCV B and RISCV NC, 2 KiB on each of RISCV T0, T1 and T2. Private to its own core, unlike L1, which all five cores see at the same addresses" } ] }, @@ -107,22 +108,26 @@ "LLVMTriple": "riscv32-unknown-elf", "LLVMFeatures": [ { - "Name": "rv32im", - "Description": "RISC-V 32-bit base integer ISA with multiply extension for Baby RISC-V cores" + "Name": "m", + "Description": "RISC-V integer multiply/divide extension. The Tensix and Ethernet Baby RISC-V cores implement RV32IM, with hardware multiply and divide both present; there is no C (compressed), A (atomic) or F/D (hardware floating-point) extension, so the ABI is ilp32" + }, + { + "Name": "tt-wh-tensix", + "Description": "Tenstorrent's Wormhole SFPU vector extension. Tensix kernels are built with the SFPI riscv-tt-elf GCC toolchain and select it with -mcpu=tt-wh-tensix (historically -march=rv32iw); it exposes the 32-lane SFPU and its 8 LRegs. There is no upstream LLVM equivalent, so LLVM can only target the plain RV32IM scalar core" } ], "SubUnits": [ { "Type": "TT-Metalium", - "Version": "1.0" + "Version": "0.76.0" }, { "Type": "TT-NN", - "Version": "1.0" + "Version": "0.76.0" }, { "Type": "TT-Forge", - "Version": "1.0" + "Version": "1.4.0" } ] }, @@ -134,10 +139,9 @@ "Present": true, "Name": "Tensix Matrix Engine (FPU)", "Count": 80, - "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "INT4", "Other"], + "SupportedPrecisions": ["FP32", "FP16", "BF16", "INT8", "Other"], "Performance": { - "fp16TopsPerGPU": 74, - "int8TopsPerGPU": 262 + "fp16TopsPerGPU": 74 } }, "videoCodecs": { @@ -152,11 +156,11 @@ "PowerStates": [ { "Name": "Active", - "Description": "Full power state with all Tensix cores and DRAM controllers active, up to 160W TBP (n150) or 300W TBP (n300 dual-die)" + "Description": "Full power state with all enabled Tensix, Ethernet and DRAM tiles active at the 1 GHz AI clock. Total board power is up to 160 W on n150 (single ASIC, the value in MaxTDP) and up to 300 W on n300 (two ASICs); both draw through one 4+4-pin EPS12V connector. Die operating range is 0-75 C, and the card throttles rather than exceeding it" }, { "Name": "Idle", - "Description": "Reduced power state when no workloads are running" + "Description": "Reduced power state when no kernels are dispatched; the ARC management core lowers Tensix clocks and voltage while keeping the PCIe and Ethernet links up" } ], "ClockGating": true, @@ -165,7 +169,7 @@ "PciExpress": { "Version": "4.0", "Lanes": 16, - "Bandwidth": 32 + "Bandwidth": 64 }, "MultiGpuSupport": { "Technologies": ["Other"], diff --git a/docs/JSON_API.md b/docs/JSON_API.md index 69d10761..dd87f03d 100644 --- a/docs/JSON_API.md +++ b/docs/JSON_API.md @@ -40,13 +40,14 @@ Describes the semiconductor manufacturing process. Describes the compute hierarchy of the chip. - **Name** (string, optional): Vendor-specific name for compute units (e.g., "CU", "SM", "Xe-core") -- **SubUnits** (array of objects, optional): Sub-units within each compute unit. - - **Type** (string, required): Type of sub-unit (e.g., "SIMD", "Tensor Core") +- **ChipType** (string, optional): Unit type representing the whole chip (a key in `UnitTypes`) +- **CoreType** (string, optional): Unit type representing a single core (a key in `UnitTypes`, e.g., `"NeuronCore-v3"`, `"Tensix Core"`) +- **UnitTypes** (object, optional): Unit types in the compute hierarchy, keyed by type name (e.g., "SIMD", "Tensor Core"). Each value describes one unit type: - **Count** (integer, required): Number of these units per parent unit - - **Size** (integer, required): Number of sub-units inside this unit (e.g., threads) - - **Description** (string, optional): Description of the sub-unit's function + - **Size** (integer, optional, default `1`): Number of sub-units inside this unit (e.g., threads) + - **Description** (string, optional): Description of the unit's function - **Memory** (array of strings, optional): Memory types embedded in the core (refer to MemorySubsystem) - - **SubunitType** (string, optional): Name of the sub-unit (lookup in this list) + - **Subunits** (array of strings, optional): Names of the unit types contained in this unit (each entry is a key in `UnitTypes`) ### MemorySubsystem @@ -98,15 +99,15 @@ Describes specialized processing units. }, "CoreSubsystem": { "Name": "CU", - "SubUnits": [ - { - "Type": "SIMD", + "CoreType": "SIMD", + "UnitTypes": { + "SIMD": { "Count": 4, "Size": 32, "Description": "Vector ALU", "Memory": ["vector register file"] } - ] + } }, "MemorySubsystem": { "SupportedMemoryTypes": ["GDDR6", "HBM2"], diff --git a/include/knexus-api/_nxs_propertys.h b/include/knexus-api/_nxs_propertys.h index 8fd6182a..5308d349 100644 --- a/include/knexus-api/_nxs_propertys.h +++ b/include/knexus-api/_nxs_propertys.h @@ -107,7 +107,10 @@ KNEXUS_API_PROP(Size, _prop_int, "Number of Sub-Units") KNEXUS_API_PROP(Rank, _prop_int, "Rank") KNEXUS_API_PROP(Shape, _prop_int_vec, "Shape") KNEXUS_API_PROP(SubUnits, _prop_str_vec, "Sub-Unit Vector") -KNEXUS_API_PROP(SubUnitType, _prop_str, "Sub-Unit Type") +KNEXUS_API_PROP(ChipType, _prop_str, "Chip Unit Type Name") +KNEXUS_API_PROP(CoreType, _prop_str, "Core Unit Type Name") +KNEXUS_API_PROP(UnitTypes, _prop_obj_vec, "Unit Type Map (keyed by type name)") +KNEXUS_API_PROP(Subunits, _prop_str_vec, "Contained Sub-Unit Type Names") KNEXUS_API_PROP(Keys, _prop_int_vec, "Node Keys") diff --git a/requirements.txt b/requirements.txt index 454ec52a..027c0d04 100644 --- a/requirements.txt +++ b/requirements.txt @@ -1,4 +1,5 @@ setuptools>=40.8.0 wheel cmake -pybind11>=2.13.6 +# Must match torch-nexus/pyproject.toml's pybind11 pin. +pybind11==3.1.0 diff --git a/schema/device_info_schema.json b/schema/device_info_schema.json index 3f5c355b..2775dd2c 100644 --- a/schema/device_info_schema.json +++ b/schema/device_info_schema.json @@ -50,24 +50,29 @@ "type": "string", "description": "Vendor-specific name for compute units (e.g., CU, SM, Xe-core)" }, - "SubUnits": { - "type": "array", - "description": "Sub-units within each compute unit (e.g. chip, CU, SIMD, Tensor Core", - "items": { + "ChipType": { + "type": "string", + "description": "Unit type representing the whole chip (a key in UnitTypes)" + }, + "CoreType": { + "type": "string", + "description": "Unit type representing a single core (a key in UnitTypes, e.g., NeuronCore-v3, Tensix Core)" + }, + "UnitTypes": { + "type": "object", + "description": "Unit types in the compute hierarchy (e.g. chip, CU, SIMD, Tensor Core), keyed by type name", + "additionalProperties": { "type": "object", - "required": ["Type", "Count", "Size"], + "required": ["Count"], "properties": { - "Type": { - "type": "string", - "description": "Type name of the sub-unit (e.g., SIMD, Tensor Core)" - }, "Count": { "type": "integer", "description": "Number of these units per parent unit" }, "Size": { "type": "integer", - "description": "Number of sub-units inside this unit (e.g. threads)" + "default": 1, + "description": "Number of sub-units inside this unit (e.g. threads). Defaults to 1 when absent" }, "Description": { "type": "string", @@ -80,9 +85,13 @@ "description": "Specify memory type embedded in the Core (Lookup in MemorySubsystem)" } }, - "SubunitType": { - "type": "string", - "description": "Name of the sub-unit (lookup in this list)" + "Subunits": { + "type": "array", + "description": "Names of the unit types contained in this unit (each entry is a key in UnitTypes)", + "items": { + "type": "string", + "description": "Type name of a contained unit (lookup as a key in UnitTypes)" + } } } } diff --git a/src/_info_impl.h b/src/_info_impl.h index 68dd1742..5f0b843c 100644 --- a/src/_info_impl.h +++ b/src/_info_impl.h @@ -57,6 +57,7 @@ class InfoImpl { json getNode(const std::vector &path) const; nxs_property_type getNodeType(json node) const; std::optional getValue(json node, nxs_int propTypeId) const; + bool isUnitTypeEntry(const std::vector &path) const; std::optional getKeys(json node) const; std::optional getProp( const std::vector &path) const; diff --git a/src/info.cpp b/src/info.cpp index 081e78e7..b106113a 100644 --- a/src/info.cpp +++ b/src/info.cpp @@ -111,6 +111,14 @@ std::optional InfoImpl::getValue(json node, return std::nullopt; } +// CoreSubsystem UnitTypes entries default Size to 1 when absent, matching +// "default": 1 in schema/device_info_schema.json. The path of such an entry is +// {..., "UnitTypes", "", "Size"}. +bool InfoImpl::isUnitTypeEntry( + const std::vector &path) const { + return path.size() >= 3 && path[path.size() - 3] == "UnitTypes"; +} + std::optional InfoImpl::getKeys(json node) const { if (node.is_object()) { std::vector keys; @@ -131,6 +139,8 @@ std::optional InfoImpl::getProp( auto node = getNode(path); if (node.is_object()) { if (tail == "Keys") return getKeys(node); + if (tail == "Size" && !node.contains("Size") && isUnitTypeEntry(path)) + return Property((nxs_long)1); } else if (node.is_array()) { if (tail == "Size") return Property((nxs_long)node.size()); // get elem