この新しい言語はNVIDIAのGPUモノポリオを殺すかもしれない

All Images AI-generated by the author for free with NightCafe Studio - see the footer for the link. アーティストによって生成されたすべての画像は、NightCafe Studioで無料です。

高性能コンピューティングの時代は、単一の名前で定義されています。奇跡

NVIDIAのプラットフォームはGPUのパワーを解除し、事実上の標準となった。

十数年以上にわたり、GPUをプログラミングすることはCUDAでプログラミングすることを意味していました。

しかし、この支配は檻を作り、単一のベンダーへの進歩を閉じ込めた。

しかし、今、2025年半ば、状況は変わっている。

しかし、今、2025年半ば、状況は変わっている。

The computing world is now undergoing a radical transformation towards heterogeneity.

私たちは、専門的なハードウェアの普及を見ています:

• Intel Gaudi Series:

  Intel's Gaudi processors are designed specifically for deep learning training and inference, offering a competitive alternative to Nvidia's GPUs.

  
  

• AMD Instinct MI Series:

  AMD's MI series of GPUs is designed for high-performance computing and AI workloads, providing an alternative to Nvidia's data center GPUs.

  
  

• Groq Tensor Streaming Processor (TSP):

  Groq's TSP architecture is designed for low-latency inference and high throughput, particularly for large language models.

  
  

• Google TPUs (Tensor Processing Units):

  Google's TPUs are custom-designed chips optimized for machine learning workloads, particularly in Google's cloud infrastructure.

  
  

• AWS Trainium:

  AWS Trainium is a chip designed for machine learning training, offering high performance and cost-effectiveness.

そして、カスタマイズされたシリコンチップを構築するスタートアップ企業が日々増えています。

この新しい多様な風景は、新しいプログラミング哲学を必要とします。

MLIR(Multi-Level Intermediate Representation)とMojoプログラミング言語

This is not just another competitor; they represent a fundamental paradigm shift.

This is a revolution in how we design, optimize, and deploy software for any hardware.

This article will deeply explore the architectural chasm between CUDA and MLIR.

1. 私たちは、具体的で実践的な比較を提供するために、完全な作業コードの例を使用します。
2. MLIRが、その敬虔な前任者であるLLVMに比べて、なぜ突破口であるかを解き明かす。
3. 私たちは、Mojoが優れた長期的なソリューションであると主張します。
4. 私たちは、なぜこの新しいスタックがコストとスピードのゲーム変換器であるのかを分析します。

この影響は、重要な新興分野、例えば、Generative AI, Quantum Computing, and evenBlockchain.

未来を見つめ、カバーする。mining ASICsで、Neuromorphic Computingそして、specialized hardwareGPU が悪く処理する少量のデータストリーム

これは、ある時代の終わりと新しい時代の始まりの物語です。

これは、ある時代の終わりと新しい時代の始まりの物語です。

この変化の大きさを理解するためには、まず、understand the four key players.

1. CUDA: The Powerful, Proprietary Incumbent

CUDA stands for Compute Unified Device Architecture.

これはNVIDIAのパラレルコンピューティングプラットフォームとプログラミングモデルです。

これは、開発者がNVIDIA GPUで実行されるC++のようなコードを書くことを可能にします。

CUDA's Strengths:

Its ecosystem of libraries is mature and unmatched:

• Mathematical Libraries:
  • cuBLAS: For basic linear algebra subprograms (BLAS).
  • cuRAND: For random number generation.
  • cuFFT: For Fast Fourier Transforms.
  • cuSPARSE: For sparse matrix operations.
  • cuTENSOR: For tensor operations.
  • cuSOLVER: For dense and sparse direct solvers.
• Parallel Algorithm Libraries:
  • nvGRAPH: For graph algorithms.
  • Thrust: For parallel algorithms and data structures.
• Communication Libraries:
  • NVSHMEM: For partitioned global address space (PGAS) programming.
  • NCCL: For multi-GPU and multi-node collective communication.
• Deep Learning Libraries:
  • cuDNN: For deep neural network computations.
  • TensorRT: For optimized deep learning inference.
  • Riva: For conversational AI.
  • DALI: For data loading and augmentation for deep learning.

ハードウェアに対する直接的で低レベルのコントロールを提供し、専門家にとってピークパフォーマンスを可能にします。

Its long history has built a massive community with vast documentation and support.

CUDA's Fatal Flaw: The Cage

Vendor Lock-In: CUDA code runs only on NVIDIA GPUs.

これは、開発者や全体の産業を単一の高価なハードウェアサプライヤーに結びつける。

それは競争を抑制し、仕事に最適なハードウェアを選択する自由を制限します。

「The Two-Language Problem: A Major Bottleneck in AI and Scientific Computing」

研究者はPythonのような高レベルの言語でプロトタイプを開発し、そのシンプルさとイテレーションのスピードのために開発した。

But for production, performance-critical code must be completely rewritten in low-level C++/CUDA.

これは痛ましく、費用がかかる切断を引き起こし、研究から展開への道のりを遅らせる。

プログラミングの複雑さ:

CUDAは強力ですが、悪名高い複雑さと言葉です。

開発者は強制的に メモリマネージャーは、CPU(ホスト)とGPU(デバイス)の間でデータを転送します。

開発者はまた、ハードウェアスケジュールヤーであり、トレードブロック、グリッド、および同期を管理する必要があります。

この複雑さは、急激な学習曲線であり、微妙なバグの頻繁な源である。

2. LLVM: The Foundation and Its "Semantic Gap”

LLVM プロジェクトは、モジュラーおよび再利用可能なコンパイラ技術のコレクションです。

その核心はLLVM Intermediate Representation(IR)であり、低レベルのアセンブリのような言語です。

LLVMは、現代のコンパイラバックエンド、特にCPUの標準となった。

コンパイラフロントエンド(Clang for C++ など)は、ソースコードを LLVM IR に翻訳します。

LLVM バックエンドは、この IR を最適化し、特定の CPU のマシンコードに変換します。

このモジュラリズムはその時代に革命的でした。

しかし、LLVMはCPU中心の世界のために設計されました。

そのIRは異質なハードウェアの新しい世界にとってあまりにも低いレベルです。

It loses crucial high-level information from the source code, a problem known as the "semantic gap."

For example, when compiling a TensorFlow model, the knowledge that an operation is a Convolution is lost.

LLVM IR only sees a generic collection of loops and arithmetic instructions.

This prevents the compiler from performing powerful, domain-specific optimizations.

It no longer understands the programmer's high-level intent.

This is the essence of the “semantic gap problem.”

And this problem is what MLIR has Solved.

3. MLIR: The Universal Translator for Hardware

MLIRは、CPU、GPU、およびそのTPUのためのTensorFlowをコンパイルする必要性からGoogleで生まれた。

彼らはLLVMの単一、低レベルのIRが十分ではなかったことに気づいた。

MLIRの突破口は、複数のIRを定義し、構成するための統一されたインフラストラクチャです。

これらの構成可能な IR は、方言と呼ばれます。

MLIR is like a universal translator, fluent in everything from high-level concepts to low-level machine details.

高レベルの方言は、ドメイン特有の概念を直接表すことができます。

For example, a "TensorFlow dialect" has an operation for tf.conv2d.

A "Linear Algebra dialect" has an operation for linalg.matmul.

これにより、LLVMが排除する重要なセマンティック情報が保持されます。

これにより、強力なコンパイラ戦略と呼ばれる漸進的な減少♪♪♪

1. コンパイラは、高レベルの方言表示から始まります。
2. この方言で高レベルのドメイン特有の最適化を実行します。
3. 次に、それは一連の中間方言を通じてコードを徐々に「下げる」のです。
4. 各中間方言は、独自の特定の最適化を実行します。
5. 最後に、最終的なマシンコード生成のためのLLVM IR 方言のような低レベルの方言に達する。

This process preserves high-level context for as long as possible.

This enables vastly superior optimizations for any hardware target.

MLIRは、高レベルの言語と多様なシリコンの間の欠けているリンクです。

4. Mojo: The User-Friendly Face of MLIR's Power

MLIRが強力で複雑なエンジンであれば、Mojoはスリムで直感的なユーザーインターフェイスです。

Mojo was created by Chris Lattner, the original architect of LLVM and the Swift language.

Mojoは、LLVMとSwift言語のオリジナルアーキテクターであるChris Lattnerによって作成されました。

最初の原則から、MLIR時代の完璧な言語になるように設計されています。

この点で、今では最も技術的に進歩した言語です。

Even Rust is based on LLVM and has all of LLVM’s shortcomings.

Mojo is the only major programming language today based on MLIR.

Mojo's Key Features:

「Superset of Python」

• Mojo は、既存の Python エコシステムとの完全な互換性を目的としています。
• これは殺し屋の特徴です!
• 開発者はNumPy、Pandas、MatplotlibなどのPythonライブラリをインポートして使用することができます。
• It completely bypasses the "cold start" problem that new languages face by tapping into Python's vast ecosystem.

True Systems プログラミングの特徴:

• Pythonとは異なり、Mojoは強力な静的入力を持つコンパイル言語です。
• This eliminates entire classes of runtime errors and enables C++-level performance optimizations.
• モダンなメモリ管理コンセプトを導入し、ゴミコレクターのオーバーヘッドなしでメモリのセキュリティを確保するために、所有権と貸し出し(Rustから)を導入します。

First-Class MLIR Integration:

• Mojo exposes the full power of MLIR directly to the developer.
• プログラマーは、ほとんどのアプリケーションのための高レベルのパイソンコードを書くことができます。
• 最大のパフォーマンスが必要な場合、特定の MLIR 方言を使用し、低レベルのカーネルを書き込むことができます。
• 重要なことは、これらはすべて同じファイル内で、同じ言語で行うことができます。

Mojo elegantly solves the "two-language problem."

Full Code Examples and Analysis

理論は一つ、実践はもう一つ。

以下の完全な作業コードの例 -

二つのパラダイム間の深い違いを示すでしょう。

Example 1: Matrix Multiplication

これは、ハイパフォーマンスコンピューティングの「Hello, World!」であり、各プラットフォームの核心哲学を明確に示しています。

The Full CUDA Implementation

これは、マトリックス倍増のための完全でコンパイル可能なCUDAプログラムです。

(キューダ C++)

// Filename: matmul.cu
// To compile: nvcc matmul.cu -o matmul_cuda

#include <iostream>
#include <vector>
#include <cuda_runtime.h>

// Helper to check for CUDA errors
#define CUDA_CHECK(err) { \
    cudaError_t err_code = err; \
    if (err_code != cudaSuccess) { \
        std::cerr << "CUDA Error: " << cudaGetErrorString(err_code) << " at line " << __LINE__ << std::endl; \
        exit(EXIT_FAILURE); \
    } \
}

// CUDA Kernel for Matrix Multiplication (Device Code)
__global__ void matrixMulKernel(float* C, const float* A, const float* B, int N) {
    // Calculate the global row and column index of the element
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    // Boundary check to avoid accessing out-of-bounds memory
    if (row < N && col < N) {
        float p_value = 0.0f;
        // Each thread computes one element of the result matrix C
        for (int k = 0; k < N; ++k) {
            p_value += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = p_value;
    }
}

// Main function (Host Code)
int main() {
    const int N = 256;
    const int size = N * N * sizeof(float);

    // Step 1. Allocate host memory
    std::vector<float> h_A(N * N);
    std::vector<float> h_B(N * N);
    std::vector<float> h_C(N * N);

    // Initialize host matrices
    for (int i = 0; i < N * N; ++i) {
        h_A[i] = static_cast<float>(rand()) / RAND_MAX;
        h_B[i] = static_cast<float>(rand()) / RAND_MAX;
    }

    // Step 2. Allocate device memory
    float *d_A, *d_B, *d_C;
    CUDA_CHECK(cudaMalloc((void**)&d_A, size));
    CUDA_CHECK(cudaMalloc((void**)&d_B, size));
    CUDA_CHECK(cudaMalloc((void**)&d_C, size));

    // Step 3. Copy matrices from host to device
    std::cout << "Copying data from host to device..." << std::endl;
    CUDA_CHECK(cudaMemcpy(d_A, h_A.data(), size, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_B, h_B.data(), size, cudaMemcpyHostToDevice));

    // Step 4. Define kernel launch configuration
    // Use 16x16 threads per block, a common choice
    dim3 threadsPerBlock(16, 16);
    // Calculate the number of blocks needed in each dimension
    dim3 numBlocks((N + threadsPerBlock.x - 1) / threadsPerBlock.x, (N + threadsPerBlock.y - 1) / threadsPerBlock.y);

    // Step 5. Launch the kernel on the device
    std::cout << "Launching kernel..." << std::endl;
    matrixMulKernel<<<numBlocks, threadsPerBlock>>>(d_C, d_A, d_B, N);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize()); // Wait for the kernel to finish

    // Step 6. Copy the result matrix back from device to host
    std::cout << "Copying result from device to host..." << std::endl;
    CUDA_CHECK(cudaMemcpy(h_C.data(), d_C, size, cudaMemcpyDeviceToHost));

    // Step 7. Free device memory
    CUDA_CHECK(cudaFree(d_A));
    CUDA_CHECK(cudaFree(d_B));
    CUDA_CHECK(cudaFree(d_C));

    std::cout << "CUDA Matrix Multiplication finished successfully." << std::endl;
    // (Optional: Add verification step here)

    return 0;
}

Analysis of the CUDA Code:

コードは、ボイラープレートと低レベルの管理によって支配されています。

ステップ1、2、3、6、7は、CPU/GPUの境界線を越えてメモリを管理するためのものです。

これは退屈で、エラーの傾向があり、コアアルゴリズムを暗くします。

The globalkeyword, blockIdx, threadIdx, and the <<<...>>> syntax are CUDA-specific hardware abstractions. Keyword, blockIdx, threadIdx, and the <<<...>>> syntax are CUDA-specific hardware abstractions.

このコードは、基本的にNVIDIAのハードウェアアーキテクチャに永久的に結びついている。

実際のアルゴリズム、つまり3つのロープは、全体のコードのほんのわずかな部分である。

The programmer's mental overhead is spent on hardware management, not on the problem itself.

The Full Mojo Implementation

この Mojo バージョンは、驚くべきシンプルさとパワーで同じ結果を達成します。

(ムービー)

# Filename: matmul.mojo
# To run: mojo matmul.mojo

from memory import DType, Tensor
from random import rand
from time import now

fn matmul_naive(C: Tensor[DType.float32], A: Tensor[DType.float32], B: Tensor[DType.float32]):
    """A naive, high-level implementation of matrix multiplication."""
    let N = A.dim(0)
    let M = A.dim(1)
    let P = B.dim(1)

    for i in range(N):
        for j in range(P):
            var sum: Float32 = 0.0
            for k in range(M):
                sum += A.load(i, k) * B.load(k, j)
            C.store(i, j, sum)

fn main():
    let N = 256
    
    # 1. Allocate and initialize tensors.
    # Mojo's Tensor handles memory allocation automatically.
    # The compiler will place it in the most appropriate memory space.
    var A = Tensor[DType.float32](N, N)
    var B = Tensor[DType.float32](N, N)
    var C = Tensor[DType.float32](N, N)

    for i in range(N):
        for j in range(N):
            A.store(i, j, rand[DType.float32]())
            B.store(i, j, rand[DType.float32]())

    print("Starting Mojo Matrix Multiplication...")
    
    let start_time = now()
    
    # 2. Call the function.
    # The MLIR-based compiler optimizes this high-level code.
    # It can automatically tile, vectorize, and parallelize this code
    # for the target hardware (CPU, GPU, etc.).
    matmul_naive(C, A, B)

    let end_time = now()
    let duration_ms = (end_time - start_time) / 1_000_000.0

    print("Mojo Matrix Multiplication finished successfully.")
    print("Execution time:", duration_ms, "ms")
    # (Optional: Print a corner of the result matrix to verify)
    print("Result C[0,0]:", C.load(0,0))
}

And that is all!

The Mojo Approach is Far Superior

Programmability and Focus:

• Mojo コードはクリーンで、アルゴリズムを直接表現します。
• プログラマーは何(数学)に焦点を当て、どのように(メモリ転送)するかではなく、
• 手動のcudaMalloc、cudaMemcpy、またはcudaFreeはありません。
• That entire class of errors is gone.

Abstraction with Performance:

• シンプルな巣立ったループは、実行されるものではありません。
• The MLIR-based compiler performs sophisticated transformations.
• そのため、この単純なコードを高度に最適化されたカーネルに変換します。
• It can apply tiling, vectorization, and parallelization automatically.
• プログラマーは、コンパイラをガイドするためにヒント(例えば @vectorize または @parallelize)を追加し、複雑さのないコントロールを実現できます。

Portability (The Ultimate Advantage):

• This is the crucial point.
• 同じ matmul.mojo ファイルは、NVIDIA GPU、AMD GPU、AVX512 搭載の Intel CPU、または Google TPU で実行するように再コンパイルできます。
• 論理は同じで、コンパイラバックエンドは変わります。
• CUDAコードは、新しいハードウェアターゲットごとに完全かつ費用がかかる書き換えが必要になります。
• Mojoは「パフォーマンス・ポータビリティ」を提供し、ベンダー・ロック・インを破り、コードを未来に保証します。

MLIR-based Mojo is undeniably set to replace LLVM-based CUDA, and developers will enjoy the change!

Mojoの詳細については、下の記事を参照してください。

Example 2: Gen AI and the Transformer Attention Mechanism

The "attention" mechanism is the heart of models like GPT-4 and is a major computational bottleneck.

Optimizing it is critical.

The CUDA Implementation (Conceptual FlashAttention)

FlashAttention は、GPU の遅いメインメモリ(HBM)とスピードインチップメモリ(SRAM)の間のデータの移動を手動で専門的にオーケストラ化するマニュアルです。

The real code is thousands of lines long and incredibly complex.

アルゴリズムの完全な実装のコンポーネントへのリンクは以下の通りです。

https://github.com/Dao-AILab/flash-attention/blob/main/csrc/flash_attn/src/flash_fwd_kernel.h

https://github.com/Dao-AILab/flash-attention/blob/main/csrc/flash_attn/flash_api.cpp

全部で約3000本の長さです。

The repository contains thousands of files.

学習曲線とオンボード曲線は両方とも急激です。

A simplified version (AI-generated) is given below:

(キューダ C++)

// This is a simplified conceptual view of a FlashAttention-style CUDA kernel.
// The actual implementation is far more complex.

template<typename Kernel_traits>
__global__ void flash_attention_fwd_kernel(Flash_fwd_params params) {

    // 1. Incredibly complex setup code
    // Calculates dozens of pointers and indices for HBM and shared memory (SRAM)
    const int block_row_idx = blockIdx.x;
    const int head_idx = blockIdx.y;
    // ... many more calculations ...

    // 2. Explicitly allocate shared memory tiles for Q, K, V
    // The developer must manage this limited resource manually.
    extern __shared__ char smem[];
    float* sQ = (float*)smem;
    float* sK = sQ + kTileM * kTileK;
    float* sV = sK + kTileN * kTileK;

    // 3. Main loop over the sequence, manually loading blocks
    for (int k_block_idx = 0; k_block_idx < params.k_num_blocks; ++k_block_idx) {

        // Manually orchestrate asynchronous loads from HBM into SRAM
        // to hide memory latency. This is extremely difficult to get right.
        load_qkv_block_from_hbm(params, ...);
        __syncthreads(); // Hard synchronization barrier

        // Manually perform matrix multiplication in fast SRAM
        compute_sram_matmul(sQ, sK, ...);

        // Recompute softmax "online" to avoid writing the huge intermediate
        // attention score matrix back to slow HBM. This is the core trick.
        compute_online_softmax(...);
        __syncthrows();

        // Update the output block
        update_output_block(sV, ...);
    }

    // 4. Manually write the final output block back to HBM
    store_output_to_hbm(params, ...);
}

Analysis of the CUDA/FlashAttention Approach:

• これは、マニュアル、ハードウェア特有のエンジニアリングの傑作です。
• GPU を手動でプログラムされたマシンのように扱うことで信じられないほどのパフォーマンスを達成します。
• これにより、コードはほぼ読み取り不可能で、維持不可能で、移植不可能になります。
• Only a handful of world-class experts can write or modify such code.
• それは閉鎖的な生態系内のパフォーマンスの頂点を表していますが、複雑さと硬さの頂点でもあります。

The Conceptual Mojo Implementation

Mojoのバージョンは同じことを表現します。アルゴリズムアイデア (tiling, online softmax) at a high level, delegating the hardware orchestration to the MLIR compiler.

(笑)

from memory import DType, Tensor
from algorithm import parallelize

struct AttentionParams:
    var is_causal: Bool
    # ... other model parameters

# This function is a high-level, portable description of the FlashAttention algorithm.
fn flash_attention[T: DType](Q: Tensor[T], K: Tensor[T], V: Tensor[T], params: AttentionParams) -> Tensor[T]:
    # Define problem dimensions from input tensors
    let num_batches = Q.dim(0)
    let num_heads = Q.dim(2)
    let seqlen_q = Q.dim(1)
    let seqlen_k = K.dim(1)
    
    # Define tunable tiling parameters. The compiler can use these as hints.
    alias BLOCK_M: Int = 128
    alias BLOCK_N: Int = 64

    # The output tensor
    var O = Tensor[T](Q.dims)

    # The @parallelize decorator tells the compiler to map this function
    # over the available hardware parallelism (e.g., CUDA thread blocks or CPU cores).
    @parallelize(num_batches * num_heads)
    fn compute_head(batch_idx: Int, head_idx: Int):
        
        # Define per-worker accumulators. The compiler will map these
        # to the fastest available memory (e.g., registers or SRAM).
        var o_i = Tensor[T](seqlen_q, V.dim(3))
        var l_i = Tensor[T](seqlen_q) # Stores the denominator of the softmax
        var m_i = Tensor[T](seqlen_q) # Stores the max of each row for stable softmax
        o_i.zero()
        l_i.fill(0.0)
        m_i.fill(-50000.0) # Negative infinity

        # Loop over blocks of the Key/Value sequence
        for j in range(0, seqlen_k, BLOCK_N):
            # 1. Load tiles of K and V.
            # The compiler is responsible for generating the optimal code
            # to move this data from main memory to fast memory.
            let k_j = K.load_tile[BLOCK_N](batch_idx, j, head_idx)
            let v_j = V.load_tile[BLOCK_N](batch_idx, j, head_idx)
            
            # Loop over blocks of the Query sequence
            for i in range(0, seqlen_q, BLOCK_M):
                # 2. Load tile of Q.
                let q_i = Q.load_tile[BLOCK_M](batch_idx, i, head_idx)
                
                # 3. Compute attention scores for the tile. This is a simple matmul.
                let s_ij = q_i @ k_j.transpose()
                
                # Causal masking for decoder models like GPT
                if params.is_causal:
                    # Algorithmic logic, no hardware specifics
                    apply_causal_mask(s_ij, i, j)

                # 4. Perform the "online softmax" update.
                # This is pure mathematical logic, not memory management.
                let m_ij = row_max(s_ij)
                let p_ij = exp(s_ij - m_ij)
                let l_ij = row_sum(p_ij)
                
                let m_new = max(m_i, m_ij)
                let l_new = exp(m_i - m_new) * l_i + exp(m_ij - m_new) * l_ij

                # Update output tile
                o_i = (l_i / l_new * exp(m_i - m_new)) * o_i + (exp(m_ij - m_new) / l_new) * (p_ij @ v_j)

                # Update softmax stats
                l_i = l_new
                m_i = m_new

        # 5. Store the final output. The compiler manages the write-back.
        O.store_tile(batch_idx, head_idx, o_i)
    
    compute_head()
    return O

1 ファイル

100個未満の場所

No brain-racking dependencies.

Of course, this is just the algorithm, but in the repository, the same algorithm took nearly 3000 LOC with CUDA!

So now you understand the difference:

Mojo is Game-Changing for AI:

Separation of Concerns:

• Mojo コードはアルゴリズムを説明します。
• CUDAコードは、ハードウェアのマニュアル実装を記述します。
• これは深い違いです。
• Mojoプログラマーは、アルゴリズムの改善に焦点を当てることができます:
• MLIRコンパイラはそれをシリコンにマッピングすることに焦点を当てています。

Research Velocity and Maintainability:

• An AI researcher can easily understand and modify this Mojo code to test a new idea.
• CUDAコードを変更することは、希少なスキルを必要とする大規模で時間がかかるエンジニアリングプロジェクトです。
• これは、研究開発サイクルを劇的に加速させます。

Hardware Freedom:(最も重要なもの)

• This Mojo code is not tied to NVIDIA.
• It can be compiled to run on:
  • AMD GPUs
  • Google TPUs
  • Intel Gaudi
  • Custom AI chips.
  • Any architecture there is!
• MLIR の言語は、新しいハードウェアをサポートするために拡張できます。
• Mojo コードを真に未来証明にします。

This breaks the NVIDIA monopoly on high-performance AI and will drive down costs.

Specialized Hardware and Future Domains

The limitations of the CUDA model become even more apparent when we look beyond traditional dense workloads to the future of computing.

MLIR/Mojo is designed for this future.

Blockchain, Mining, and ASICs

Bitcoin のような Proof-of-Work ブロックチェーンには膨大なハッシュパワーが必要です。

目標は、他のデータとハッシュされた場合、特定のターゲットの下の結果を生成する「ノンセ」を見つけることです。

This is a brute-force search, perfect for parallel hardware.

Initially, miners used CPUs, then GPUs for their superior parallelism.

The CUDA code for a SHA-256 miner is low-level, focused on bitwise and integer operations.

However, for a stable, unchanging algorithm like SHA-256, the ultimate hardware is an ASIC.

An ASIC (Application-Specific Integrated Circuit) is a chip designed for one single purpose - to implement an algorithm in hardware.

ASIC(Application-Specific Integrated Circuit)は、ハードウェアにアルゴリズムを実装するための1つの目的のために設計されたチップです。

An SHA-256 ASIC has the hashing logic literally baked into the silicon.

それは、その1つのタスクのためのGPUよりも何千倍も効率的です。

ここでCUDAの物語は終わりますが、MLIR/Mojoの物語はさらに興味深くなります。

チップを設計するプロセスは、High Level Synthesis(HLS)と呼ばれています。

HLS ツールは、アルゴリズムの高レベルの説明を、チップを製造するために使用される低レベルのハードウェア説明言語 (Verilog または VHDL など) に変換します。

MLIR, through projects like CIRCT (Circuit IR for Compilers and Tools), is designed to be the backbone of next-generation HLS.

1. 開発者はMojoでハッシュアルゴリズムを書くことができます。
2. GPU マイニングでは、GPU バックエンドを使用してコンパイルします。
3. ASIC を作成するには、HLS バックエンドを使用して同じ Mojo コードをコンパイルできます。
4. The MLIR infrastructure would lower the high-level Mojo logic into Verilog.

これは、高レベルのソフトウェアからカスタマイズされたシリコンデザインまで、全体を統一します。

それは、GPUや新しいASICであれ、最も効率的なハードウェアに新しいアルゴリズムの迅速なプロトタイプ化と展開を可能にします。

CUDAには答えがありません。

CUDA has no answer to this.

It is a software-only solution for a single vendor's programmable hardware.

Neuromorphic Computing and Sparse Data

NVIDIA GPUs are masters of SIMT: Single Instruction, Multiple Thread.

これは、何千ものトレードが異なるデータ(例えば、マトリックスの倍数)で同じ命令を実行している場合に非常に効率的であることを意味します。

However, they are very inefficient at workloads with heavy branching or irregular data access.

理由は「トレード・ディヴェーゲンス」です。

If threads in a group (a "warp") take different branches of an if/else statement, the hardware must execute both paths serially, with threads in the inactive path simply turned off.

これは多くの重要な問題のためにパフォーマンスを殺します。

Neuromorphic Computing:

これは脳にインスピレーションを与えたコンピュータパラダイムです。

インテルのLoihiのような神経形チップは、時計や密集マトリックスの数学に基づいていません。

They are event-driven.

「ニューロン」は、その入力潜在力が限界を超えるときにのみ「スピーク」を発射する。

これらのピークは他の「シナプス」に移動し、その後他のニューロンが燃える可能性があります。

これは極めて稀少で、分支が重く、非同期的なプロセスです。

Trying to simulate this on a GPU is horrifically inefficient due to constant thread divergence.

MLIRはそのための完璧なソリューションです。

MLIRはそのための完璧なソリューションです。

1. MLIR内で「神経形方言」を作成することができる。
2. この方言には、Spike、Synapse、NeuronUpdateのための一流の操作があります。
3. 開発者は、これらの高レベルの概念を使用してMojoで神経形アルゴリズムを書くことができます。
4. MLIRコンパイラは、Loihiのような特定の神経形チップのバックエンドを備え、これらのコンセプトをチップのネイティブ、イベント駆動指示に翻訳します。

これは、完全に非伝統的な形式のコンピューティングのためのポータブルで高レベルのプログラミングモデルを可能にします。

CUDAモデルはこの分野では無関係です。

The CUDA model is not relevant in this domain.

Sparse and Graph Data:

多くの現実世界の問題には、わずかなデータが含まれています:ソーシャルネットワーク、推奨エンジン、科学的シミュレーション。

これらを密集したマトリックスとして表すことは無駄である。

これらを密集したマトリックスとして表すことは無駄である。

これらをGPUで処理すると、不規則なメモリアクセスパターンが生じ、GPUのメモリコールシングリングの最適化とパフォーマンスが低下する。

Again, MLIR provides the answer.

1. 「graph dialect」または「sparse tensor dialect」は、これらのデータ構造をネイティブに表すことができます。
2. コンパイラは、スパルシティに対処するための専門的な最適化を適用することができます。
3. たとえば、ノードを再編成してメモリの位置を改善したり、圧縮されたストレージ形式を使用したりすることができます。

4. 
   

This allows a high-level algorithm written in Mojo to be efficiently compiled for sparse data on any hardware.

これは今日、極めて困難なことである。

そして、Cudaと一緒に不可能に近づく。

そして、Cudaと一緒に不可能に近づく。

Quantum Computing Simulation

Simulating a quantum computer on a classical computer is essential for developing and testing quantum algorithms.

最も一般的な方法は、ステータスベクターシミュレーションです。

N-qubit 量子システムの状態は 2^N 複合数のベクターによって表される。

For just 50 qubits, this vector has 2^50 (over a quadrillion) elements, requiring petabytes of memory.

量子アルゴリズムは「ゲート」の連続である。

各ゲートは、巨大なステータスベクターを非常に大きな、非常に希少なマトリックスに倍増することに等しい。

これは、コンピュータ密集性とメモリ帯域幅の両方に関連するワークロードです。

NVIDIA は、高性能の CUDA ベースのソリューションである cuQuantum ライブラリで、ここに大きく投資しています。

cuQuantum は NVIDIA GPU で非常に高速ですが、古典的な CUDA 制限があります。

1. ベンダーロックイン:あなたの量子シミュレーションは、NVIDIAハードウェアに結びついている。
2. 低レベルの最適化:コンパイラは、マトリックスベクターの倍数のみを表示します。
3. No Domain Advantage: 量子力学の最適化はなく、LLVM(セマンティック・ギャップ)に基づいています。

The MLIR/Mojo Advantage for Quantum Simulation:

MLIRアプローチは、コンパイラ内のより高いレベルのインテリジェンスを可能にします。

1. MLIR では「量子方言」を定義することができる。
2. この方言は、ゲートをマトリックスとして表すのではなく、それらをその量子物体として表すだろう:Hadamard、CNOT、Toffoli。
3. 開発者は、これらの高レベルのオブジェクトを使用してMojoで量子回路を書きます。
4. MLIRコンパイラは、いかなるマトリックスも生成される前に量子特有の最適化を実行することができます。

たとえば、コンパイラは、ハダマードゲート(H)を連続して2回適用することは、アイデンティティ操作であり、完全に排除できることを知るでしょう。

それは、ゲートの特定の順序が、より効率的な単一のゲートに「合併」することができることを知るだろう。

たとえば、コンパイラは、ハダマードゲート(H)を連続して2回適用することは、アイデンティティ操作であり、完全に排除できることを知るでしょう。

それは、ゲートの特定の順序が、より効率的な単一のゲートに「合併」することができることを知るだろう。

This is an entire class of optimization that is invisible to the CUDA compiler, which only sees generic matrices, thanks to LLVM.

これらの高レベルのアルジェビア的単純化を実行した後、MLIRコンパイラはその後、単純化された回路をターゲットハードウェアの狭いマトリックス操作の最適化された順序に下げる。

Because this is all built on MLIR, the same high-level quantum circuit written in Mojo could be compiled to run on an NVIDIA GPU, an AMD GPU, or a CPU cluster.

これは、より高いパフォーマンス(よりスマートな最適化による)と完全なハードウェアの自由の両方を提供します。

Nvidiaは、量子シミュレーションハードウェアとソフトウェアステックに大きく投資しています。

しかし、そのCUDA-Qプラットフォームは依然としてLLVMベースです。

MLIR-based Mojo can not just offer advanced optimization - it also offers simpler programming.

Final Verdict: Today vs. The Inevitable Future

The Verdict Today (2025):

1. クダは山の王であり、山は広い。
2. 成熟した生態系、広範な図書館、そして巨大なコミュニティは強力な資産です。
3. NVIDIAのハードウェアにすでに投資しており、すぐに製品を配信する必要があるチームにとって、CUDAは実践的な選択肢です。
4. 統治の10年の惰性は強力な力である。
5. ムービーはまだ若い。
6. その生態系は信じられないほどのスピードで成長しているが、まだCUDAの戦闘テストされた図書館の広さに匹敵することはできない。

The Verdict for the Long Run:

1. 未来は異質である。
2. これは推測ではなく、現実です。
3. カスタムAIシリコンの増加とAMDとIntelの新たな競争は、ベンダーのロックインを許容できないビジネスおよび技術的リスクにした。
4. 将来の問題、すなわち、希少なデータ、神経形人工知能、ブロックチェーンマイニング、量子コンピューティングは、今日のGPUの固いSIMTモデルに適していません。
5. MLIRは、この問題を解決するために設計された唯一の既存の業界サポートアーキテクチャです。
6. Google、Apple、Intel、AMD、ARMによる採用は、コンパイラの将来の中心的な役割を明確に示している。
7. Mojo は、この力を活用するために作られた唯一の言語です。

ムービー:

• 2 言語問題の解決
• 実用性とパフォーマンスを組み合わせる
• MLIRエコシステム全体へのゲートウェイを提供します。

CUDAからMLIRに基づく世界への移行は徐々に進むが、それは避けられない。

これは、閉鎖的でハードウェア中心のモデルから、ソフトウェアによって定義されたオープンな未来への根本的な転換です。

Mojoの欠点

1. Mojoはまだ開発中です。
2. まだ授業もありません。
3. 第三者図書館は少ないが、信じられないほどのペースで成長している。
4. Python が使われているすべての場所にアプリケーションがありますが、Python で進化する必要があります。
5. 言語全体はまだオープンソースではありませんが、専門家によると、それは間もなく変わるでしょう。
6. Windows(まだ)をサポートしていません。
7. そして、Android、iOS、Edge IOTシステムへのポートが必要です。

しかし、長期的には勝者になるのでしょうか。

I believe it will, and developers will be happier with Mojo than CUDA.

結論

CUDA built the impressive palace of today's high-performance computing.

But it is a cage.

MLIR and Mojo are handing every developer the key to unlock it and build the future on any foundation they choose.

そして、その財団はMLIRとMojoになる予定です。

そして、その財団はMLIRとMojoになる予定です。

The simplest reason - the budget.

だからこそ、Nvidiaがピヨットしない限り、そしてすぐに:

This will be the end of the dominance of Nvidia - unless they embrace MLIR as well!

参照

Official Project Pages

• MLIR (Multi-Level Intermediate Representation)
  • Text description: The official homepage for the MLIR project, hosted by LLVM. This is the canonical source for documentation, talks, and the project's overall mission statement.
  • https://mlir.llvm.org/
• Mojo Programming Language
  • The official documentation for the Mojo programming language from Modular, the company that created it. This is the primary resource for learning the language.[2]
  • https://docs.modular.com/mojo/
• NVIDIA CUDA Toolkit
  • The official portal from NVIDIA for downloading the CUDA Toolkit, which includes the compilers, libraries, and tools necessary for CUDA development.
  • https://developer.nvidia.com/cuda-toolkit
• LLVM Compiler Infrastructure Project
  • The main homepage for the LLVM project, which provides an overview of the entire ecosystem, including Clang, LLDB, and other sub-projects. MLIR is a part of this larger project.
  • https://llvm.org/
• Chris Lattner's Homepage
  • The personal homepage of Chris Lattner, the creator of LLVM, Clang, Swift, MLIR, and Mojo. It provides his work history and links to his talks and papers, offering direct insight into the creation of these technologies.
  • https://nondot.org/sabre/

AI and Attention Mechanism (FlashAttention)

• FlashAttention Original Paper (arXiv)
  • The original scientific paper, "FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness," which introduced the algorithm. This is the primary source for understanding the technical details and performance benefits.
  • https://arxiv.org/abs/2205.14135
• FlashAttention-2 Paper (arXiv)
  • The follow-up paper describing FlashAttention-2, which details further optimizations for parallelism and work partitioning to achieve even greater speedups on modern GPUs.
  • https://arxiv.org/abs/2307.08691
• FlashAttention GitHub Repository
  • The official GitHub repository containing the source code for the FlashAttention and FlashAttention-2 CUDA kernels.
  • https://github.com/Dao-AILab/flash-attention

Quantum Computing Simulation

• NVIDIA cuQuantum Official Page
  • NVIDIA's official product page for the cuQuantum SDK, outlining its features for accelerating quantum computing simulations on GPUs.
  • https://developer.nvidia.com/cuquantum
• NVIDIA cuQuantum Documentation
  • The detailed technical documentation for the cuQuantum SDK, providing a high-level overview and API references for the libraries.
  • https://docs.nvidia.com/cuda/cuquantum/index.html

Specialized Hardware (Neuromorphic & ASICs)

• Intel Neuromorphic Computing Overview
  • Intel's official overview of their neuromorphic computing research, which discusses the goals of the program and the Loihi research chips.
  • https://www.intel.com/content/www/us/en/research/neuromorphic-computing.html
• CIRCT (Circuit IR Compilers and Tools) Project
  • The official homepage for the CIRCT project, an LLVM/MLIR incubator looking to apply compiler technology to hardware design, including High-Level Synthesis (HLS) for FPGAs and ASICs.
  • https://circt.llvm.org/
• CIRCT GitHub Repository
  • The official GitHub repository for the CIRCT project, containing the source code, dialects, and tools for hardware compiler design.
  • https://github.com/llvm/circt

Google AI Studio は、この記事の概要と研究に使用されました。

https://aistudio.google.com/

All pictures were generated by the author with NightCafe Studio for free, available at the link below:

https://creator.nightcafe.studio/