GPU-Puzzles

GitHub
12.1k 932 中等 1 次阅读 昨天MIT开发框架
AI 解读 由 AI 自动生成，仅供参考

GPU-Puzzles 是一款专为初学者设计的交互式学习工具，旨在帮助用户轻松掌握 CUDA GPU 编程。在机器学习领域，GPU 架构至关重要，但许多从业者往往只停留在调用高层抽象库的层面，缺乏对底层并行计算原理的直观理解。GPU-Puzzles 通过“解谜”的方式解决了这一痛点：它摒弃了枯燥的理论文档，让用户直接在一个类似 Jupyter Notebook 的环境中编写代码，即时构建和测试 GPU 内核。

该工具的核心亮点在于利用 Numba 库，让用户使用熟悉的 Python 语法来编写实际上等同于低级 CUDA 的代码。用户无需配置复杂的本地开发环境，只需在 Google Colab 中开启 GPU 模式即可开始挑战。从简单的向量映射（Map）到更复杂的并行操作，一系列精心设计的谜题引导用户逐步理解线程索引、内存访问等核心概念。通过完成这些练习，用户能在几小时内建立起对支撑现代深度学习算法的并行计算机制的深刻直觉。

GPU-Puzzles 非常适合希望深入理解硬件加速原理的机器学习工程师、数据科学家以及计算机专业学生。对于那些已经熟练使用 PyTorch 或 TensorFlow 等框架，却想揭开“黑盒”面纱、探索底层性能优化之道的开发者来说，这是一个不可多得的入门阶梯。它以寓教于乐的形式，让高门槛的 GPU 编程变得触手可及。

使用场景

某深度学习框架研发团队的初级工程师，在尝试优化自定义算子性能时，因缺乏底层 GPU 架构直觉而陷入瓶颈。

没有 GPU-Puzzles 时

  • 理论脱节：工程师虽熟悉 PyTorch 高层 API，但面对 CUDA 线程块、共享内存等抽象概念时，只能死记硬背文档，无法建立直观的空间映射感。
  • 试错成本高：直接编写原生 CUDA C++ 代码极易出现内存越界或线程同步错误，调试过程繁琐，往往花费数天排查一个索引偏移问题。
  • 环境门槛高：配置本地 GPU 编译环境复杂，新手常在安装驱动、编译器版本匹配上耗费大量精力，迟迟无法进入核心逻辑学习。
  • 反馈周期长：修改代码后需重新编译、部署并运行完整测试脚本，无法即时验证对单个线程行为的理解是否正确。

使用 GPU-Puzzles 后

  • 交互式直觉构建：通过“地图（Map）”、“压缩（Zip）”等拼图游戏，工程师在 Colab 中直接用 Python 语法编写 CUDA 内核，瞬间理解 threadIdx.x 如何对应数据位置。
  • 即时正向反馈：每完成一个谜题，系统立即可视化线程执行流程并评分，将原本晦涩的并行逻辑转化为可见的得分奖励，快速纠正认知偏差。
  • 零配置上手：无需安装本地 CUDA 工具链，依托 Google Colab 云端 GPU 环境，打开链接即可开始编写和调试内核，几分钟内进入心流状态。
  • 平滑过渡实战：从简单的向量加法的谜题入手，逐步掌握归约、卷积等核心算法模式，几小时内即可具备阅读和优化真实深度学习算子源码的能力。

GPU-Puzzles 将枯燥的底层并行计算原理转化为直观的编程闯关体验，让开发者在数小时内跨越从“调用者”到“掌控者”的认知鸿沟。

运行环境要求

操作系统
  • 未说明
GPU

需要支持 CUDA 的 NVIDIA GPU（推荐在 Google Colab 中运行，需开启 GPU 模式）

内存

未说明

依赖
notes该项目是一个交互式教程，旨在通过 Numba 将 Python 代码直接映射为 CUDA 内核来教授 GPU 编程。强烈建议在 Google Colab 环境中运行，需在运行时设置中将硬件加速器更改为 GPU。代码中使用了特定的 chalk 库分支和自定义的 lib.py 文件，需按 README 中的指令预先安装和下载。
python未说明
numba
numpy
chalk (特定分支)
GPU-Puzzles hero image

快速开始

GPU 拼图

GPU 架构对机器学习至关重要，而且似乎每天都在变得更加重要。然而，你完全可以在不接触 GPU 代码的情况下成为机器学习专家。在抽象层面上工作很难培养直观的理解。

这个笔记本旨在以完全交互式的方式教授初学者 GPU 编程。它没有提供包含概念的文字说明，而是直接让你动手编写和构建 GPU 内核。这些练习使用 NUMBA，它可以将 Python 代码直接映射到 CUDA 内核。虽然看起来像 Python，但实际上与编写低级 CUDA 代码几乎相同。 我认为，在几个小时内，你就可以从基础知识入手，理解当今驱动 99% 深度学习的真正算法。如果你想阅读官方文档，可以在这里找到：

NUMBA CUDA 指南

我建议在 Colab 中完成这些练习，因为它很容易上手。请务必先复制一份自己的副本，在设置中启用 GPU 模式（Runtime / Change runtime type，然后将 Hardware accelerator 设置为 GPU），然后再开始编码。

在 Colab 中打开

（如果你喜欢这种拼图风格，也可以看看我的 Tensor Puzzles，适用于 PyTorch。）

操作指南

!pip install -qqq git+https://github.com/danoneata/chalk@srush-patch-1
!wget -q https://github.com/srush/GPU-Puzzles/raw/main/robot.png https://github.com/srush/GPU-Puzzles/raw/main/lib.py

import numba
import numpy as np
import warnings
from lib import CudaProblem, Coord

warnings.filterwarnings(
    action="ignore", category=numba.NumbaPerformanceWarning, module="numba"
)

拼图 1：映射

实现一个“内核”（GPU 函数），将向量 a 的每个位置加 10，并将结果存储在向量 out 中。每个位置对应一个线程。

警告 这段代码看起来像 Python，但实际上是 CUDA！你不能使用标准的 Python 工具，比如列表推导式，也不能询问 NumPy 数组的形状或大小（如果需要大小，会作为参数传入）。这些拼图只需要进行简单的操作，基本上是 +、*、简单的数组索引、for 循环和 if 语句。你可以使用局部变量。如果你遇到错误，那很可能是因为你尝试了过于复杂的操作 :).

提示：可以把函数 call 看作是为每个线程运行一次。唯一的区别是每次调用时 cuda.threadIdx.x 都会变化。

def map_spec(a):
    return a + 10


def map_test(cuda):
    def call(out, a) -> None:
        local_i = cuda.threadIdx.x
        # 填充此处（大约 1 行）

    return call


SIZE = 4
out = np.zeros((SIZE,))
a = np.arange(SIZE)
problem = CudaProblem(
    "Map", map_test, [a], out, threadsperblock=Coord(SIZE, 1), spec=map_spec
)
problem.show()

# Map
 
   分数（每线程最大值）：
   | 全局读取 | 全局写入 | 共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试失败。
你的结果：[0. 0. 0. 0.]
标准答案：[10 11 12 13]

拼图 2 - Zip

实现一个内核，将 a 和 b 的每个位置相加，并将结果存储在 out 中。每个位置对应一个线程。

def zip_spec(a, b):
    return a + b


def zip_test(cuda):
    def call(out, a, b) -> None:
        local_i = cuda.threadIdx.x
        # 填充此处（大约 1 行）

    return call


SIZE = 4
out = np.zeros((SIZE,))
a = np.arange(SIZE)
b = np.arange(SIZE)
problem = CudaProblem(
    "Zip", zip_test, [a, b], out, threadsperblock=Coord(SIZE, 1), spec=zip_spec
)
problem.show()

# Zip
 
   分数（每线程最大值）：
   | 全局读取 | 全局写入 | 共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg



problem.check()

测试失败。
你的结果：[0. 0. 0. 0.]
标准答案：[0 2 4 6]

拼图 3 - 保护

实现一个内核，将向量 a 的每个位置加 10，并将结果存储在向量 out 中。线程数量多于位置数量。

def map_guard_test(cuda):
    def call(out, a, size) -> None:
        local_i = cuda.threadIdx.x
        # 填充此处（大约 2 行）

    return call


SIZE = 4
out = np.zeros((SIZE,))
a = np.arange(SIZE)
problem = CudaProblem(
    "Guard",
    map_guard_test,
    [a],
    out,
    [SIZE],
    threadsperblock=Coord(8, 1),
    spec=map_spec,
)
problem.show()

# Guard
 
   分数（每线程最大值）：
   | 全局读取 | 全局写入 | 共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试失败。
你的结果：[0. 0. 0. 0.]
标准答案：[10 11 12 13]

拼图 4 - 二维映射

实现一个内核，将矩阵 a 的每个位置加 10，并将结果存储在矩阵 out 中。输入矩阵 a 是正方形的二维矩阵。线程数量多于位置数量。

def map_2D_test(cuda):
    def call(out, a, size) -> None:
        local_i = cuda.threadIdx.x
        local_j = cuda.threadIdx.y
        # 填充此处（大约 2 行）

    return call


SIZE = 2
out = np.zeros((SIZE, SIZE))
a = np.arange(SIZE * SIZE).reshape((SIZE, SIZE))
problem = CudaProblem(
    "Map 2D", map_2D_test, [a], out, [SIZE], threadsperblock=Coord(3, 3), spec=map_spec
)
problem.show()

# Map 2D
 
   分数（每线程最大值）：
   | 全局读取 | 全局写入 | 共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试失败。
你的结果：[[0. 0.]
 [0. 0.]]
标准答案：[[10 11]
 [12 13]]

谜题 5 - 广播

实现一个内核，将 a 和 b 相加，并将结果存储在 out 中。 输入 a 和 b 是向量。线程的数量多于元素的位置数。

def broadcast_test(cuda):
    def call(out, a, b, size) -> None:
        local_i = cuda.threadIdx.x
        local_j = cuda.threadIdx.y
        # 填充我（大约 2 行）

    return call


SIZE = 2
out = np.zeros((SIZE, SIZE))
a = np.arange(SIZE).reshape(SIZE, 1)
b = np.arange(SIZE).reshape(1, SIZE)
problem = CudaProblem(
    "Broadcast",
    broadcast_test,
    [a, b],
    out,
    [SIZE],
    threadsperblock=Coord(3, 3),
    spec=zip_spec,
)
problem.show()

# 广播
 
   每线程得分（最大值）：
   | 全局读取 | 全局写入 | 共享内存读取 | 共享内存写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试未通过。
您的结果：[[0. 0.]
 [0. 0.]]
规范结果：[[0 1]
 [1 2]]

谜题 6 - 块

实现一个内核，将 a 的每个位置加上 10，并将结果存储在 out 中。 每个块中的线程数量少于 a 的大小。

提示：块是一组线程。每个块的线程数量是有限的，但我们可以有多个不同的块。变量 cuda.blockIdx 告诉我们当前所在的块编号。

def map_block_test(cuda):
    def call(out, a, size) -> None:
        i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x
        # 填充我（大约 2 行）

    return call


SIZE = 9
out = np.zeros((SIZE,))
a = np.arange(SIZE)
problem = CudaProblem(
    "Blocks",
    map_block_test,
    [a],
    out,
    [SIZE],
    threadsperblock=Coord(4, 1),
    blockspergrid=Coord(3, 1),
    spec=map_spec,
)
problem.show()

# 块
 
   每线程得分（最大值）：
   | 全局读取 | 全局写入 | 共享内存读取 | 共享内存写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试未通过。
您的结果：[0. 0. 0. 0. 0. 0. 0. 0. 0.]
规范结果：[10 11 12 13 14 15 16 17 18]

谜题 7 - 二维块

以 2D 方式实现相同的内核。每个块中的线程数量在两个方向上都少于 a 的大小。

def map_block2D_test(cuda):
    def call(out, a, size) -> None:
        i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x
        # 填充我（大约 4 行）

    return call


SIZE = 5
out = np.zeros((SIZE, SIZE))
a = np.ones((SIZE, SIZE))

problem = CudaProblem(
    "Blocks 2D",
    map_block2D_test,
    [a],
    out,
    [SIZE],
    threadsperblock=Coord(3, 3),
    blockspergrid=Coord(2, 2),
    spec=map_spec,
)
problem.show()

# 二维块
 
   每线程得分（最大值）：
   | 全局读取 | 全局写入 | 共享内存读取 | 共享内存写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试未通过。
您的结果：[[0. 0. 0. 0. 0.]
 [0. 0. 0. 0. 0.]
 [0. 0. 0. 0. 0.]
 [0. 0. 0. 0. 0.]
 [0. 0. 0. 0. 0.]]
规范结果：[[11. 11. 11. 11. 11.]
 [11. 11. 11. 11. 11.]
 [11. 11. 11. 11. 11.]
 [11. 11. 11. 11. 11.]
 [11. 11. 11. 11. 11.]]

谜题 8 - 共享内存

实现一个内核，将 a 的每个位置加上 10，并将结果存储在 out 中。 每个块中的线程数量少于 a 的大小。

警告：每个块只能拥有固定大小的共享内存，该块内的线程可以读写。这必须是一个字面意义上的 Python 常量，而不是变量。在写入共享内存后，需要调用 cuda.syncthreads 来确保线程之间不会发生交叉。

（这个示例实际上并不需要共享内存或 syncthreads，但这是一个演示。）

TPB = 4
def shared_test(cuda):
    def call(out, a, size) -> None:
        shared = cuda.shared.array(TPB, numba.float32)
        i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x
        local_i = cuda.threadIdx.x

        if i < size:
            shared[local_i] = a[i]
            cuda.syncthreads()

        # 填充我（大约 2 行）

    return call


SIZE = 8
out = np.zeros(SIZE)
a = np.ones(SIZE)
problem = CudaProblem(
    "Shared",
    shared_test,
    [a],
    out,
    [SIZE],
    threadsperblock=Coord(TPB, 1),
    blockspergrid=Coord(2, 1),
    spec=map_spec,
)
problem.show()

# 共享内存
 
   每线程得分（最大值）：
   | 全局读取 | 全局写入 | 共享内存读取 | 共享内存写入 |
   |             1 |             0 |             0 |             1 |

svg

problem.check()

测试未通过。
您的结果：[0. 0. 0. 0. 0. 0. 0. 0.]
规范结果：[11. 11. 11. 11. 11. 11. 11. 11.]

谜题 9 - 池化

实现一个内核，将 a 的最后 3 个位置相加，并将结果存储在 out 中。 每个位置对应 1 个线程。每个线程只需要 1 次全局读取和 1 次全局写入。

提示：请记住要小心同步。

def pool_spec(a):
    out = np.zeros(*a.shape)
    for i in range(a.shape[0]):
        out[i] = a[max(i - 2, 0) : i + 1].sum()
    return out


TPB = 8
def pool_test(cuda):
    def call(out, a, size) -> None:
        shared = cuda.shared.array(TPB, numba.float32)
        i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x
        local_i = cuda.threadIdx.x
        # 填充我（大约 8 行）

    return call


SIZE = 8
out = np.zeros(SIZE)
a = np.arange(SIZE)
problem = CudaProblem(
    "Pooling",
    pool_test,
    [a],
    out,
    [SIZE],
    threadsperblock=Coord(TPB, 1),
    blockspergrid=Coord(1, 1),
    spec=pool_spec,
)
problem.show()

# 池化
 
   每线程得分（最大值）：
   | 全局读取 | 全局写入 | 共享内存读取 | 共享内存写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试未通过。
您的结果：[0. 0. 0. 0. 0. 0. 0. 0.]
规范结果：[ 0.  1.  3.  6.  9. 12. 15. 18.]

谜题10 - 点积

实现一个内核，计算 a 和 b 的点积，并将结果存储在 out 中。 每个线程对应一个位置。每个线程只需要2次全局读取和1次全局写入。

注意：对于这个问题，你不需要担心共享内存的读取次数。我们稍后再处理这个挑战。

def dot_spec(a, b):
    return a @ b

TPB = 8
def dot_test(cuda):
    def call(out, a, b, size) -> None:
        shared = cuda.shared.array(TPB, numba.float32)

        i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x
        local_i = cuda.threadIdx.x
        # 填充我 (大约9行)
    return call


SIZE = 8
out = np.zeros(1)
a = np.arange(SIZE)
b = np.arange(SIZE)
problem = CudaProblem(
    "Dot",
    dot_test,
    [a, b],
    out,
    [SIZE],
    threadsperblock=Coord(SIZE, 1),
    blockspergrid=Coord(1, 1),
    spec=dot_spec,
)
problem.show()

# Dot
 
   分数 (每线程最大值)：
   |  全局读取 | 全局写入 |  共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试失败。
你的结果：[0.]
标准答案：140

谜题11 - 一维卷积

实现一个内核，计算 a 和 b 之间的一维卷积，并将结果存储在 out 中。 你需要处理一般情况。每个线程只需要2次全局读取和1次全局写入。

def conv_spec(a, b):
    out = np.zeros(*a.shape)
    len = b.shape[0]
    for i in range(a.shape[0]):
        out[i] = sum([a[i + j] * b[j] for j in range(len) if i + j < a.shape[0]])
    return out


MAX_CONV = 4
TPB = 8
TPB_MAX_CONV = TPB + MAX_CONV
def conv_test(cuda):
    def call(out, a, b, a_size, b_size) -> None:
        i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x
        local_i = cuda.threadIdx.x

        # 填充我 (大约17行)

    return call


# 测试1

SIZE = 6
CONV = 3
out = np.zeros(SIZE)
a = np.arange(SIZE)
b = np.arange(CONV)
problem = CudaProblem(
    "1D Conv (Simple)",
    conv_test,
    [a, b],
    out,
    [SIZE, CONV],
    Coord(1, 1),
    Coord(TPB, 1),
    spec=conv_spec,
)
problem.show()

# 1D Conv (Simple)
 
   分数 (每线程最大值)：
   |  全局读取 | 全局写入 |  共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试失败。
你的结果：[0. 0. 0. 0. 0. 0.]
标准答案：[ 5.  8. 11. 14.  5.  0.]

测试2

out = np.zeros(15)
a = np.arange(15)
b = np.arange(4)
problem = CudaProblem(
    "1D Conv (Full)",
    conv_test,
    [a, b],
    out,
    [15, 4],
    Coord(2, 1),
    Coord(TPB, 1),
    spec=conv_spec,
)
problem.show()

# 1D Conv (Full)
 
   分数 (每线程最大值)：
   |  全局读取 | 全局写入 |  共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试失败。
你的结果：[0. 0. 0. 0. 0. 0. 0. 0. 0. 0. 0. 0. 0. 0. 0.]
标准答案：[14. 20. 26. 32. 38. 44. 50. 56. 62. 68. 74. 80. 41. 14.  0.]

谜题12 - 前缀和

实现一个内核，对 a 进行求和，并将结果存储在 out 中。 如果 a 的大小大于块大小，则只存储每个块的和。

我们将使用共享内存中的并行前缀和算法来实现这一点。 也就是说，算法的每一步都应该将剩余数字的一半相加。请按照下图所示进行操作：

TPB = 8
def sum_spec(a):
    out = np.zeros((a.shape[0] + TPB - 1) // TPB)
    for j, i in enumerate(range(0, a.shape[-1], TPB)):
        out[j] = a[i : i + TPB].sum()
    return out


def sum_test(cuda):
    def call(out, a, size: int) -> None:
        cache = cuda.shared.array(TPB, numba.float32)
        i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x
        local_i = cuda.threadIdx.x
        # 填充我 (大约12行)

    return call


# 测试1

SIZE = 8
out = np.zeros(1)
inp = np.arange(SIZE)
problem = CudaProblem(
    "Sum (Simple)",
    sum_test,
    [inp],
    out,
    [SIZE],
    Coord(1, 1),
    Coord(TPB, 1),
    spec=sum_spec,
)
problem.show()

# Sum (Simple)
 
   分数 (每线程最大值)：
   |  全局读取 | 全局写入 |  共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试失败。
你的结果：[0.]
标准答案：[28.]

测试2

SIZE = 15
out = np.zeros(2)
inp = np.arange(SIZE)
problem = CudaProblem(
    "Sum (Full)",
    sum_test,
    [inp],
    out,
    [SIZE],
    Coord(2, 1),
    Coord(TPB, 1),
    spec=sum_spec,
)
problem.show()

# Sum (Full)
 
   分数 (每线程最大值)：
   |  全局读取 | 全局写入 |  共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试失败。
你的结果：[0. 0.]
标准答案：[28. 77.]

谜题13 - 轴向求和

实现一个内核，对 a 的每一列进行求和，并将结果存储在 out 中。

TPB = 8
def sum_spec(a):
    out = np.zeros((a.shape[0], (a.shape[1] + TPB - 1) // TPB))
    for j, i in enumerate(range(0, a.shape[-1], TPB)):
        out[..., j] = a[..., i : i + TPB].sum(-1)
    return out


def axis_sum_test(cuda):
    def call(out, a, size: int) -> None:
        cache = cuda.shared.array(TPB, numba.float32)
        i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x
        local_i = cuda.threadIdx.x
        batch = cuda.blockIdx.y
        # 填充我 (大约12行)

    return call


BATCH = 4
SIZE = 6
out = np.zeros((BATCH, 1))
inp = np.arange(BATCH * SIZE).reshape((BATCH, SIZE))
problem = CudaProblem(
    "Axis Sum",
    axis_sum_test,
    [inp],
    out,
    [SIZE],
    Coord(1, BATCH),
    Coord(TPB, 1),
    spec=sum_spec,
)
problem.show()

# Axis Sum
 
   分数 (每线程最大值)：
   |  全局读取 | 全局写入 |  共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试失败。
你的结果：[[0.]
 [0.]
 [0.]
 [0.]]
标准答案：[[ 15.]
 [ 51.]
 [ 87.]
 [123.]]

谜题 14 - 矩阵乘法！

实现一个核函数，用于将方阵 a 和 b 相乘，并将结果存储在 out 中。

提示：这里最高效的算法是先将一个块复制到共享内存中，然后再计算每个单独的行-列点积。如果矩阵可以完全放入共享内存，则很容易做到这一点。请先实现这种情况，然后更新你的代码以计算部分点积，并迭代地移动你复制到共享内存中的那一部分。 你应该能够在只进行 6 次全局读取的情况下完成更复杂的情况。

def matmul_spec(a, b):
    return a @ b


TPB = 3
def mm_oneblock_test(cuda):
    def call(out, a, b, size: int) -> None:
        a_shared = cuda.shared.array((TPB, TPB), numba.float32)
        b_shared = cuda.shared.array((TPB, TPB), numba.float32)

        i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x
        j = cuda.blockIdx.y * cuda.blockDim.y + cuda.threadIdx.y
        local_i = cuda.threadIdx.x
        local_j = cuda.threadIdx.y
        # 填充此处（大约 14 行）

    return call

# 测试 1

SIZE = 2
out = np.zeros((SIZE, SIZE))
inp1 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE))
inp2 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE)).T

problem = CudaProblem(
    "矩阵乘法（简单）",
    mm_oneblock_test,
    [inp1, inp2],
    out,
    [SIZE],
    Coord(1, 1),
    Coord(TPB, TPB),
    spec=matmul_spec,
)
problem.show(sparse=True)

# 矩阵乘法（简单）
 
   分数（每线程最大值）：
   | 全局读取 | 全局写入 | 共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试失败。
你的结果：[[0. 0.]
 [0. 0.]]
标准答案：[[ 1  3]
 [ 3 13]]

测试 2

SIZE = 8
out = np.zeros((SIZE, SIZE))
inp1 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE))
inp2 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE)).T

problem = CudaProblem(
    "矩阵乘法（完整）",
    mm_oneblock_test,
    [inp1, inp2],
    out,
    [SIZE],
    Coord(3, 3),
    Coord(TPB, TPB),
    spec=matmul_spec,
)
problem.show(sparse=True)

# 矩阵乘法（完整）
 
   分数（每线程最大值）：
   | 全局读取 | 全局写入 | 共享读取 | 共享写入 |
   |             0 |             0 |             0 |             0 |

svg

problem.check()

测试失败。
你的结果：[[0. 0. 0. 0. 0. 0. 0. 0.]
 [0. 0. 0. 0. 0. 0. 0. 0.]
 [0. 0. 0. 0. 0. 0. 0. 0.]
 [0. 0. 0. 0. 0. 0. 0. 0.]
 [0. 0. 0. 0. 0. 0. 0. 0.]
 [0. 0. 0. 0. 0. 0. 0. 0.]
 [0. 0. 0. 0. 0. 0. 0. 0.]
 [0. 0. 0. 0. 0. 0. 0. 0.]]
标准答案：[[  140   364   588   812  1036  1260  1484  1708]
 [  364  1100  1836  2572  3308  4044  4780  5516]
 [  588  1836  3084  4332  5580  6828  8076  9324]
 [  812  2572  4332  6092  7852  9612 11372 13132]
 [ 1036  3308  5580  7852 10124 12396 14668 16940]
 [ 1260  4044  6828  9612 12396 15180 17964 20748]
 [ 1484  4780  8076 11372 14668 17964 21260 24556]
 [ 1708  5516  9324 13132 16940 20748 24556 28364]]

常见问题

相似工具推荐

openclaw

OpenClaw 是一款专为个人打造的本地化 AI 助手，旨在让你在自己的设备上拥有完全可控的智能伙伴。它打破了传统 AI 助手局限于特定网页或应用的束缚，能够直接接入你日常使用的各类通讯渠道，包括微信、WhatsApp、Telegram、Discord、iMessage 等数十种平台。无论你在哪个聊天软件中发送消息，OpenClaw 都能即时响应，甚至支持在 macOS、iOS 和 Android 设备上进行语音交互，并提供实时的画布渲染功能供你操控。 这款工具主要解决了用户对数据隐私、响应速度以及“始终在线”体验的需求。通过将 AI 部署在本地，用户无需依赖云端服务即可享受快速、私密的智能辅助，真正实现了“你的数据，你做主”。其独特的技术亮点在于强大的网关架构，将控制平面与核心助手分离，确保跨平台通信的流畅性与扩展性。 OpenClaw 非常适合希望构建个性化工作流的技术爱好者、开发者，以及注重隐私保护且不愿被单一生态绑定的普通用户。只要具备基础的终端操作能力（支持 macOS、Linux 及 Windows WSL2），即可通过简单的命令行引导完成部署。如果你渴望拥有一个懂你

★ 349.3k|★★★☆☆|1周前
Agent开发框架图像

stable-diffusion-webui

stable-diffusion-webui 是一个基于 Gradio 构建的网页版操作界面，旨在让用户能够轻松地在本地运行和使用强大的 Stable Diffusion 图像生成模型。它解决了原始模型依赖命令行、操作门槛高且功能分散的痛点，将复杂的 AI 绘图流程整合进一个直观易用的图形化平台。 无论是希望快速上手的普通创作者、需要精细控制画面细节的设计师，还是想要深入探索模型潜力的开发者与研究人员，都能从中获益。其核心亮点在于极高的功能丰富度：不仅支持文生图、图生图、局部重绘（Inpainting）和外绘（Outpainting）等基础模式，还独创了注意力机制调整、提示词矩阵、负向提示词以及“高清修复”等高级功能。此外，它内置了 GFPGAN 和 CodeFormer 等人脸修复工具，支持多种神经网络放大算法，并允许用户通过插件系统无限扩展能力。即使是显存有限的设备，stable-diffusion-webui 也提供了相应的优化选项，让高质量的 AI 艺术创作变得触手可及。

★ 162.1k|★★★☆☆|2周前
开发框架图像Agent

everything-claude-code

everything-claude-code 是一套专为 AI 编程助手（如 Claude Code、Codex、Cursor 等）打造的高性能优化系统。它不仅仅是一组配置文件，而是一个经过长期实战打磨的完整框架，旨在解决 AI 代理在实际开发中面临的效率低下、记忆丢失、安全隐患及缺乏持续学习能力等核心痛点。 通过引入技能模块化、直觉增强、记忆持久化机制以及内置的安全扫描功能，everything-claude-code 能显著提升 AI 在复杂任务中的表现，帮助开发者构建更稳定、更智能的生产级 AI 代理。其独特的“研究优先”开发理念和针对 Token 消耗的优化策略，使得模型响应更快、成本更低，同时有效防御潜在的攻击向量。 这套工具特别适合软件开发者、AI 研究人员以及希望深度定制 AI 工作流的技术团队使用。无论您是在构建大型代码库，还是需要 AI 协助进行安全审计与自动化测试，everything-claude-code 都能提供强大的底层支持。作为一个曾荣获 Anthropic 黑客大奖的开源项目，它融合了多语言支持与丰富的实战钩子（hooks），让 AI 真正成长为懂上

★ 160.8k|★★☆☆☆|今天
开发框架Agent语言模型

ComfyUI

ComfyUI 是一款功能强大且高度模块化的视觉 AI 引擎，专为设计和执行复杂的 Stable Diffusion 图像生成流程而打造。它摒弃了传统的代码编写模式，采用直观的节点式流程图界面，让用户通过连接不同的功能模块即可构建个性化的生成管线。 这一设计巧妙解决了高级 AI 绘图工作流配置复杂、灵活性不足的痛点。用户无需具备编程背景，也能自由组合模型、调整参数并实时预览效果，轻松实现从基础文生图到多步骤高清修复等各类复杂任务。ComfyUI 拥有极佳的兼容性，不仅支持 Windows、macOS 和 Linux 全平台，还广泛适配 NVIDIA、AMD、Intel 及苹果 Silicon 等多种硬件架构，并率先支持 SDXL、Flux、SD3 等前沿模型。 无论是希望深入探索算法潜力的研究人员和开发者，还是追求极致创作自由度的设计师与资深 AI 绘画爱好者，ComfyUI 都能提供强大的支持。其独特的模块化架构允许社区不断扩展新功能，使其成为当前最灵活、生态最丰富的开源扩散模型工具之一，帮助用户将创意高效转化为现实。

★ 109.2k|★★☆☆☆|昨天
开发框架图像Agent

gemini-cli

gemini-cli 是一款由谷歌推出的开源 AI 命令行工具，它将强大的 Gemini 大模型能力直接集成到用户的终端环境中。对于习惯在命令行工作的开发者而言，它提供了一条从输入提示词到获取模型响应的最短路径，无需切换窗口即可享受智能辅助。 这款工具主要解决了开发过程中频繁上下文切换的痛点，让用户能在熟悉的终端界面内直接完成代码理解、生成、调试以及自动化运维任务。无论是查询大型代码库、根据草图生成应用，还是执行复杂的 Git 操作，gemini-cli 都能通过自然语言指令高效处理。 它特别适合广大软件工程师、DevOps 人员及技术研究人员使用。其核心亮点包括支持高达 100 万 token 的超长上下文窗口，具备出色的逻辑推理能力；内置 Google 搜索、文件操作及 Shell 命令执行等实用工具；更独特的是，它支持 MCP（模型上下文协议），允许用户灵活扩展自定义集成，连接如图像生成等外部能力。此外，个人谷歌账号即可享受免费的额度支持，且项目基于 Apache 2.0 协议完全开源，是提升终端工作效率的理想助手。

★ 100.8k|★★☆☆☆|1周前
插件Agent图像

markitdown

MarkItDown 是一款由微软 AutoGen 团队打造的轻量级 Python 工具，专为将各类文件高效转换为 Markdown 格式而设计。它支持 PDF、Word、Excel、PPT、图片（含 OCR）、音频（含语音转录）、HTML 乃至 YouTube 链接等多种格式的解析，能够精准提取文档中的标题、列表、表格和链接等关键结构信息。 在人工智能应用日益普及的今天，大语言模型（LLM）虽擅长处理文本，却难以直接读取复杂的二进制办公文档。MarkItDown 恰好解决了这一痛点，它将非结构化或半结构化的文件转化为模型“原生理解”且 Token 效率极高的 Markdown 格式，成为连接本地文件与 AI 分析 pipeline 的理想桥梁。此外，它还提供了 MCP（模型上下文协议）服务器，可无缝集成到 Claude Desktop 等 LLM 应用中。 这款工具特别适合开发者、数据科学家及 AI 研究人员使用，尤其是那些需要构建文档检索增强生成（RAG）系统、进行批量文本分析或希望让 AI 助手直接“阅读”本地文件的用户。虽然生成的内容也具备一定可读性，但其核心优势在于为机器

★ 93.4k|★★☆☆☆|1周前
插件开发框架