Giao diện

và

và
Đây là một chủ đề khá nâng cao, nằm ở giao điểm của AI Compiler, GPU Programming, PyTorch, Triton và Machine Learning Systems. Tuy nhiên, nếu học đúng lộ trình thì Helion lại khá dễ hiểu vì nó được thiết kế để giấu đi sự phức tạp của Triton.
Tôi sẽ dạy theo kiểu từ nền tảng đến chuyên sâu.
Muốn hiểu Helion trước tiên phải hiểu GPU Kernel.
Giả sử ta có phép cộng hai vector.
C = A + B
Trong PyTorch chỉ cần:
C = A + B
PyTorch sẽ gọi kernel CUDA đã được viết sẵn.
Nếu bạn muốn tự viết kernel GPU thì có 3 lựa chọn.
| Mức | Công nghệ | Độ khó |
|---|---|---|
| Cao nhất | CUDA C++ | Rất khó |
| Trung bình | Triton | Khó |
| Cao hơn PyTorch | Helion | Dễ hơn Triton |
Có thể hình dung:
PyTorch
↓
Helion
↓
Triton
↓
CUDA
↓
GPU
Helion thực chất không chạy trực tiếp trên GPU.
Nó sinh ra Triton.
Triton lại sinh ra CUDA/PTX.
Ví dụ viết Matrix Multiplication bằng Triton.
Bạn phải tự tính:
Ví dụ:
pid = tl.program_id(0)
pid_m = ...
pid_n = ...
offs_m = ...
offs_n = ...
ptr = base + ...
Khoảng 60% code chỉ để tính địa chỉ bộ nhớ.
Không phải tính toán.
Helion muốn bạn không phải làm việc đó.
Thay vì bảo GPU:
“Hãy đọc địa chỉ số 10032”
Helion cho bạn nói:
“Lấy tile này.”
Ví dụ
for tile_m, tile_n in hl.tile([m, n]):
đọc giống Python.
Nhưng Helion sẽ tự biến thành
program_id
↓
tile index
↓
pointer
↓
memory address
↓
GPU
Tài liệu viết:
Python Embedded DSL
DSL nghĩa là:
Domain Specific Language.
Tức là
Một ngôn ngữ nhỏ nằm trong Python.
Giống như:
SQL nằm trong Java
TensorFlow Graph nằm trong Python
JAX DSL
Helion DSL
Bạn vẫn viết Python.
Chỉ một số lệnh
hl.tile()
hl.zeros()
hl.range()
mang ý nghĩa đặc biệt.
Ví dụ:
@helion.kernel()
def matmul(...):
...
Trong kernel có hai vùng.
m, k = x.size()
out = torch.empty(...)
Đây là Python bình thường.
Chạy trên CPU.
for tile in hl.tile(...):
Đây không còn là Python nữa.
Helion sẽ dịch đoạn này thành Triton.
Nói cách khác
CPU code
↓
Helion boundary
↓
GPU code
Đây là khái niệm quan trọng nhất.
Giả sử ma trận
1024 × 1024
GPU không xử lý toàn bộ.
Nó chia thành các ô.
Ví dụ
□□□□□□□□
□□□□□□□□
□□□□□□□□
□□□□□□□□
Mỗi ô gọi là
Tile
Ví dụ
64×64
thì
1024
↓
64
↓
16 tile
Theo chiều ngang.
Chiều dọc cũng vậy.
Tổng cộng
16 ×16 =256 tile
GPU sẽ chạy song song 256 tile.
Trong Helion
for tile_m, tile_n in hl.tile([m, n]):
không phải vòng lặp Python.
Mà nghĩa là
“Hãy chia ma trận thành các tile.”
Helion:
for tile_m, tile_n in hl.tile([m, n]):
acc = hl.zeros(...)
for tile_k in hl.tile(k):
acc = torch.addmm(...)
Nếu viết dưới dạng sơ đồ.
M
┌────┬────┐
│ │ │
├────┼────┤
│ │ │
└────┴────┘
↑
GPU xử lý từng tile
Bên trong mỗi tile
for tile_k
lại chia tiếp chiều K.
Giống GEMM chuẩn.
Trong Triton bạn phải tự chọn
Tile = 16
hay
Tile =32
hay
Tile =64
Helion không.
Bạn chỉ viết
hl.tile()
Nó tự thử.
Đây là trái tim của Helion.
Giả sử có các lựa chọn:
Tile
16
32
64
Warp
4
8
16
Indexing
pointer
block_ptr
descriptor
PID
flat
blocked
persistent
Có hàng nghìn khả năng.
Ví dụ
64
×
8 warp
×
block_ptr
×
persistent
Hay
32
×
4 warp
×
pointer
×
flat
Helion sẽ chạy thử tất cả.
Đo thời gian.
Lấy cái nhanh nhất.
Đây chính là
Autotuning
Vì Helion có thể thử
1500
2000
3000
cấu hình.
Ví dụ log
Generation 1
↓
Generation 2
↓
Generation 3
...
↓
Generation 20
Đây giống thuật toán tiến hóa.
Nó giữ cấu hình tốt.
Loại cấu hình chậm.
Sinh cấu hình mới.
Trong log có
DifferentialEvolutionSearch
Đây là thuật toán tối ưu.
Không thử ngẫu nhiên.
Mà giống tiến hóa.
Config A
Config B
↓
Lai
↓
Config C
↓
Đột biến
↓
Config D
Sau nhiều thế hệ.
Nó hội tụ.
Đây là cấu hình GPU.
Ví dụ
Config(
block_sizes
num_warps
num_stages
indexing
)
Đây không phải code.
Mà là cách GPU chạy.
Ví dụ
block_size
↓
64×64
Hay
warp
↓
8
Sau khi autotune.
Helion in ra
Config(...)
Bạn copy.
@helion.kernel(
config=...
)
Lần sau
Không cần thử nữa.
Khởi động ngay.
Đây là điểm nhiều người mới học dễ nhầm lẫn.
| Config | Settings |
|---|---|
| Quyết định cách GPU thực thi | Quyết định cách Helion biên dịch |
| Ảnh hưởng hiệu năng | Ảnh hưởng quá trình phát triển |
| Được autotune | Không autotune |
| block size, num_warps, indexing… | print_output_code, autotune_effort, logging… |
Ví dụ:
@helion.kernel(
autotune_effort="none",
print_output_code=True
)
Đây là Settings: dùng để phát triển và gỡ lỗi.
Còn:
@helion.kernel(
config=helion.Config(
block_sizes=[64,64,64],
num_warps=8
)
)
Đây là Config: quyết định hiệu năng của kernel trên GPU.
Có thể hình dung toàn bộ quá trình như sau:
Mã Helion
│
▼
Phân tích Python AST
│
▼
Sinh Triton
│
▼
TorchInductor
│
▼
PTX/CUDA
│
▼
GPU NVIDIA
Helion đóng vai trò là một lớp trừu tượng cao hơn Triton. Bạn mô tả bài toán ở mức “xử lý theo tile”, còn Helion tự sinh mã Triton và tự tối ưu.
Nếu nhìn theo mức độ trừu tượng:
| Công nghệ | Bạn phải quan tâm |
|---|---|
| CUDA | Thread, Block, Shared Memory, Warp, Pointer, Synchronization |
| Triton | Tile, Pointer, Stride, PID, Mask |
| Helion | Chỉ cần Tile và các phép toán PyTorch |
| PyTorch | Chỉ cần gọi API có sẵn |
Xu hướng của Helion tương tự như lịch sử phát triển của lập trình: thay vì yêu cầu lập trình viên điều khiển phần cứng ở mức thấp, nó chuyển gánh nặng tối ưu sang trình biên dịch và bộ autotuner.
Để hiểu Helion một cách bài bản, nên học theo thứ tự sau:
torch.addmm, torch.sum,…hl.tile(), autotuning, Config, Settings.Với nền tảng này, bạn sẽ không chỉ biết cách viết kernel bằng Helion, mà còn hiểu vì sao Helion có thể sinh ra các kernel Triton hiệu năng cao và cách đánh giá, tinh chỉnh chúng khi cần.
%pip install "torch==2.9.*" --index-url https://download.pytorch.org/whl/cu126
%pip install helion
"""
Helion Softmax Kernel Examples
==============================
This example demonstrates multiple Helion kernel implementations of the softmax function,
including a simple wrapper around PyTorch's softmax, and a numerically optimized two-pass version.
The example also includes a check function to compare these kernels against PyTorch's
built-in softmax for correctness.
"""
# %%
from __future__ import annotations
import torch
import helion
from helion._testing import run_example
import helion.language as hl
# %%
@helion.kernel(autotune_effort="quick")
def softmax(x: torch.Tensor) -> torch.Tensor:
"""
Simple Helion kernel wrapping PyTorch's softmax function.
Args:
x (torch.Tensor): Input tensor of shape [n, m].
Returns:
torch.Tensor: Softmax output tensor of the same shape.
"""
n, _m = x.size()
out = torch.empty_like(x)
for tile_n in hl.tile(n):
out[tile_n, :] = torch.nn.functional.softmax(x[tile_n, :], dim=1)
return out
# %%
def check(m: int, n: int) -> None:
"""
Runs correctness checks comparing Helion softmax kernels against PyTorch's softmax.
Args:
m (int): Number of rows in input tensor.
n (int): Number of columns in input tensor.
"""
x = torch.randn([m, n], device="cuda", dtype=torch.float16)
run_example(softmax, lambda x: torch.nn.functional.softmax(x, dim=1), (x,))
# %%
def main() -> None:
"""
Main function to run the softmax kernel correctness check with example input size.
"""
check(4096, 2560)
# %%
if __name__ == "__main__":
main()
"""
Helion Softmax Kernel Examples
==============================
This example demonstrates multiple Helion kernel implementations of the softmax function,
including a simple wrapper around PyTorch's softmax, and a numerically optimized two-pass version.
The example also includes a check function to compare these kernels against PyTorch's
built-in softmax for correctness.
"""
# %%
from __future__ import annotations
import torch
import helion
from helion._testing import run_example
import helion.language as hl
# %%
@helion.kernel(autotune_effort="quick")
def softmax_two_pass(x: torch.Tensor) -> torch.Tensor:
"""
Numerically optimized Helion kernel performing softmax in two passes.
Args:
x (torch.Tensor): Input tensor of shape [m, n].
Returns:
torch.Tensor: Softmax output tensor of the same shape.
"""
m, n = x.size()
out = torch.empty_like(x)
block_size_m = hl.register_block_size(m)
block_size_n = hl.register_block_size(n)
for tile_m in hl.tile(m, block_size=block_size_m):
mi = hl.full([tile_m], float("-inf"), dtype=torch.float32)
di = hl.zeros([tile_m], dtype=torch.float32)
for tile_n in hl.tile(n, block_size=block_size_n):
values = x[tile_m, tile_n]
local_amax = torch.amax(values, dim=1)
mi_next = torch.maximum(mi, local_amax)
di = di * torch.exp(mi - mi_next) + torch.exp(
values - mi_next[:, None]
).sum(dim=1)
mi = mi_next
for tile_n in hl.tile(n, block_size=block_size_n):
values = x[tile_m, tile_n]
out[tile_m, tile_n] = torch.exp(values - mi[:, None]) / di[:, None]
return out
# %%
def check(m: int, n: int) -> None:
"""
Runs correctness checks comparing Helion softmax kernels against PyTorch's softmax.
Args:
m (int): Number of rows in input tensor.
n (int): Number of columns in input tensor.
"""
x = torch.randn([m, n], device="cuda", dtype=torch.float16)
run_example(softmax_two_pass, lambda x: torch.nn.functional.softmax(x, dim=1), (x,))
# %%
def main() -> None:
"""
Main function to run the softmax kernel correctness check with example input size.
"""
check(4096, 2560)
# %%
if __name__ == "__main__":
main()
Đoạn mã gồm hai cách viết Softmax bằng Helion:
| Kernel | Cách làm |
|---|---|
softmax | Gọi trực tiếp torch.nn.functional.softmax trên từng nhóm hàng |
softmax_two_pass | Tự triển khai thuật toán Softmax ổn định số học theo hai lượt |
Cả hai đều được Helion biên dịch thành một GPU kernel Triton.
Với một hàng dữ liệu:
[
x = [x_1,x_2,\ldots,x_n]
]
Softmax được tính như sau:
\frac{e^{x_i}}{\sum_j e^{x_j}}
]
Ví dụ:
[
x=[1,2,3]
]
Kết quả gần đúng:
[
[0.09,\ 0.24,\ 0.67]
]
Tổng các phần tử luôn bằng 1.
Trong đoạn mã, Softmax được tính theo chiều:
dim=1
Nghĩa là tính Softmax trên từng hàng của tensor hai chiều.
%pip install "torch==2.9.*" --index-url https://download.pytorch.org/whl/cu126
%pip install helion
Dòng đầu cài PyTorch 2.9 có hỗ trợ CUDA 12.6.
Dòng thứ hai cài Helion.
Cú pháp %pip thường được dùng trong:
Nếu chạy trong file Python thông thường thì cài bằng terminal:
pip install "torch==2.9.*" --index-url https://download.pytorch.org/whl/cu126
pip install helion
from __future__ import annotations
import torch
import helion
from helion._testing import run_example
import helion.language as hl
| Thành phần | Vai trò |
|---|---|
torch | Tạo tensor và sử dụng các phép toán PyTorch |
helion | Định nghĩa và biên dịch Helion kernel |
run_example | Kiểm tra kết quả và có thể benchmark kernel |
helion.language as hl | Cung cấp các toán tử DSL như hl.tile, hl.zeros |
annotations | Trì hoãn việc đánh giá type hint |
Lưu ý rằng helion._testing có dấu gạch dưới. Đây thường là API nội bộ phục vụ kiểm thử, không nhất thiết ổn định như API công khai.
@helion.kernel(autotune_effort="quick")
def softmax(x: torch.Tensor) -> torch.Tensor:
Decorator:
@helion.kernel(...)
biến hàm Python thành một Helion kernel.
Tham số:
autotune_effort="quick"
yêu cầu Helion thực hiện autotuning ở mức nhanh.
Có thể hiểu tương đối như sau:
| Mức autotune | Đặc điểm |
|---|---|
none | Không tìm cấu hình tối ưu |
quick | Thử một số cấu hình, thời gian ngắn hơn |
| Mức cao hơn | Tìm kiếm rộng hơn, mất nhiều thời gian hơn |
Autotuner có thể lựa chọn:
n, _m = x.size()
Nếu x có shape:
[4096, 2560]
thì:
n = 4096
_m = 2560
Tuy nhiên tên biến ở đây hơi dễ gây nhầm.
Theo cách đặt tên phổ biến, nên viết:
m, n = x.size()
Trong kernel hiện tại:
n thực chất là số hàng;_m thực chất là số cột.Dấu _ ở đầu _m cho biết biến này không được sử dụng tiếp.
Viết rõ hơn có thể là:
num_rows, _num_cols = x.size()
out = torch.empty_like(x)
Tạo tensor out:
x;Nếu x có:
shape = [4096, 2560]
dtype = float16
device = cuda
thì out cũng có các thuộc tính tương tự.
Dùng empty_like nhanh hơn việc khởi tạo toàn bộ tensor bằng 0 vì kernel sẽ ghi đè tất cả các giá trị.
for tile_n in hl.tile(n):
Đây không phải vòng lặp Python thông thường.
Nó có nghĩa:
Chia chiều gồm
nhàng thành các tile và ánh xạ các tile đó lên các chương trình GPU.
Giả sử có 4096 hàng và autotuner chọn tile 64 hàng:
[
4096 / 64 = 64\text{ tile}
]
GPU có thể xử lý các tile này song song.
tile_n không phải một số nguyên đơn lẻ. Nó đại diện cho một tập chỉ số hàng, chẳng hạn:
tile 0: hàng 0 đến 63
tile 1: hàng 64 đến 127
tile 2: hàng 128 đến 191
...
out[tile_n, :] = torch.nn.functional.softmax(
x[tile_n, :],
dim=1
)
Trong đó:
x[tile_n, :]
có nghĩa:
tile_n;Ví dụ tile chứa 64 hàng, tensor tạm có thể có shape:
[64, 2560]
Sau đó:
softmax(..., dim=1)
tính Softmax trên 2560 phần tử của từng hàng.
Kết quả được ghi vào đúng vị trí:
out[tile_n, :]
Có thể hình dung:
Tensor x: 4096 × 2560
┌──────────────────────────┐
│ tile hàng số 1 │ → softmax từng hàng
├──────────────────────────┤
│ tile hàng số 2 │ → softmax từng hàng
├──────────────────────────┤
│ tile hàng số 3 │ → softmax từng hàng
├──────────────────────────┤
│ ... │
└──────────────────────────┘
return out
Phần bên ngoài vòng lặp như:
x.size()
torch.empty_like()
return
được dùng để chuẩn bị kernel.
Phần bên trong:
for tile_n in hl.tile(n):
được Helion chuyển thành mã GPU.
def check(m: int, n: int) -> None:
Hàm nhận:
m: số hàng;n: số cột.x = torch.randn(
[m, n],
device="cuda",
dtype=torch.float16
)
Ví dụ:
check(4096, 2560)
sẽ tạo tensor:
shape = [4096, 2560]
device = GPU
dtype = float16
torch.randn tạo số ngẫu nhiên theo phân phối chuẩn.
run_example(
softmax,
lambda x: torch.nn.functional.softmax(x, dim=1),
(x,)
)
Ba tham số chính:
| Tham số | Ý nghĩa |
|---|---|
softmax | Kernel Helion cần kiểm tra |
lambda x: ... | Hàm PyTorch dùng làm kết quả tham chiếu |
(x,) | Tuple chứa đối số truyền vào |
Dấu phẩy trong:
(x,)
rất quan trọng. Đây là tuple có một phần tử.
Không có dấu phẩy:
(x)
thì chỉ là biểu thức x, không phải tuple.
Về bản chất, hàm kiểm tra làm gần giống:
actual = softmax(x)
expected = torch.nn.functional.softmax(x, dim=1)
assert actual gần bằng expected
maindef main() -> None:
check(4096, 2560)
Chạy kiểm thử với tensor có:
Tổng số phần tử:
[
4096 \times 2560 = 10{,}485{,}760
]
Vì kiểu float16 dùng 2 byte mỗi phần tử, riêng một tensor chiếm khoảng:
20{,}971{,}520\text{ byte}
]
Tức khoảng 20 MiB.
if __name__ == "__main__":
main()
Nếu file được chạy trực tiếp:
python softmax_example.py
thì main() được gọi.
Nếu file được import từ module khác thì main() không tự chạy.
Trong Jupyter Notebook, đoạn này thường vẫn chạy vì __name__ thường bằng "__main__".
softmax_two_passKernel thứ hai tự triển khai Softmax thay vì gọi trực tiếp:
torch.nn.functional.softmax
Mục tiêu là xử lý trường hợp số cột lớn mà không nhất thiết phải tải toàn bộ một hàng vào register cùng lúc.
@helion.kernel(autotune_effort="quick")
def softmax_two_pass(x: torch.Tensor) -> torch.Tensor:
Công thức trực tiếp:
[
\frac{e^{x_i}}{\sum_j e^{x_j}}
]
có thể bị tràn số.
Ví dụ:
[
e^{1000}
]
là một số quá lớn, dễ trở thành vô cực:
inf
Softmax ổn định được tính bằng cách trừ giá trị lớn nhất:
\frac{e^{x_i-m}}{\sum_j e^{x_j-m}}
]
với:
[
m=\max_j x_j
]
Vì:
[
x_i-m\leq 0
]
nên:
[
e^{x_i-m}\leq 1
]
Giảm đáng kể nguy cơ overflow.
softmax_two_passm, n = x.size()
out = torch.empty_like(x)
Ở kernel này tên biến đúng với quy ước:
| Biến | Ý nghĩa |
|---|---|
m | Số hàng |
n | Số cột |
block_size_m = hl.register_block_size(m)
block_size_n = hl.register_block_size(n)
Helion được yêu cầu tạo các tham số block size cho hai chiều:
m;n.Các giá trị này có thể được autotuner lựa chọn.
Có thể hình dung:
block_size_m ∈ {1, 2, 4, 8, ...}
block_size_n ∈ {32, 64, 128, 256, ...}
Tập giá trị thực tế phụ thuộc Helion, phần cứng và không gian tìm kiếm.
Khác với:
hl.tile(m)
để Helion tự xử lý hoàn toàn, ở đây code lấy block size ra thành biến nhằm dùng lại cùng một kích thước trong nhiều vòng lặp.
Đặc biệt, block_size_n được dùng ở cả lượt thứ nhất và lượt thứ hai.
for tile_m in hl.tile(m, block_size=block_size_m):
Kernel xử lý một nhóm hàng tại một thời điểm.
Giả sử:
m = 4096
block_size_m = 8
thì mỗi chương trình GPU xử lý 8 hàng.
Số chương trình cần thiết gần bằng:
[
4096/8=512
]
Các chương trình này có thể chạy song song.
mi và dimi = hl.full(
[tile_m],
float("-inf"),
dtype=torch.float32
)
di = hl.zeros(
[tile_m],
dtype=torch.float32
)
Mỗi hàng đang được xử lý có hai giá trị trạng thái:
| Biến | Ý nghĩa |
|---|---|
mi | Giá trị lớn nhất đã thấy trong các block cột trước |
di | Tổng mũ đã được hiệu chỉnh theo giá trị lớn nhất hiện tại |
mi được khởi tạo bằng:
[
-\infty
]
vì chưa đọc giá trị nào.
di được khởi tạo bằng:
[
0
]
Cả hai dùng float32 dù đầu vào là float16.
Điều này rất quan trọng để tăng độ chính xác khi:
for tile_n in hl.tile(n, block_size=block_size_n):
Chiều cột được chia thành các block.
Giả sử:
n = 2560
block_size_n = 256
thì mỗi hàng có 10 block cột:
0–255
256–511
512–767
...
2304–2559
values = x[tile_m, tile_n]
Nếu:
tile_m có 8 hàng
tile_n có 256 cột
thì values có shape logic:
[8, 256]
local_amax = torch.amax(values, dim=1)
Tìm giá trị lớn nhất trên từng hàng trong block cột hiện tại.
Nếu:
values shape = [8, 256]
thì:
local_amax shape = [8]
Mỗi hàng có một max cục bộ.
mi_next = torch.maximum(mi, local_amax)
So sánh:
mi: max của tất cả block đã xử lý;local_amax: max của block hiện tại.Kết quả:
\max(m_{\text{old}},m_{\text{local}})
]
diĐây là dòng quan trọng nhất:
di = di * torch.exp(mi - mi_next) + torch.exp(
values - mi_next[:, None]
).sum(dim=1)
Để hiểu, trước tiên hãy xét mục tiêu.
Sau khi đọc tất cả các cột, ta muốn:
[
d_i = \sum_j e^{x_{ij}-m_i}
]
trong đó:
[
m_i=\max_j x_{ij}
]
Nhưng dữ liệu được đọc theo từng block. Khi đọc block mới, giá trị max có thể thay đổi.
torch.exp(values - mi_next[:, None]).sum(dim=1)
Tính:
[
\sum_{x\in\text{block mới}}e^{x-m_{\text{next}}}
]
mi_next có shape:
[tile_m]
Còn values có shape:
[tile_m, tile_n]
Cần thêm một chiều:
mi_next[:, None]
để có shape:
[tile_m, 1]
Sau đó broadcasting:
values [tile_m, tile_n]
mi_next[:, None] [tile_m, 1]
─────────────────
kết quả [tile_m, tile_n]
di * torch.exp(mi - mi_next)
Trước đó, di được biểu diễn theo max cũ:
\sum_{\text{cũ}}e^{x-m_{\text{old}}}
]
Nhưng giờ max mới là (m_{\text{next}}).
Ta cần đổi hệ quy chiếu:
e^{x-m_{\text{old}}}
\cdot
e^{m_{\text{old}}-m_{\text{next}}}
]
Do đó:
d_{\text{old}}
\cdot
e^{m_{\text{old}}-m_{\text{next}}}
]
Đó chính là:
di * torch.exp(mi - mi_next)
d_{\text{old}}
e^{m_{\text{old}}-m_{\text{next}}}
+
\sum_{\text{block mới}}
e^{x-m_{\text{next}}}
]
Đây thường được gọi là online softmax recurrence.
Nó cho phép tính max và tổng mũ trong cùng một lượt đọc dữ liệu, thay vì phải:
Tuy nhiên toàn bộ kernel vẫn được gọi là “two-pass” vì còn một lượt thứ hai để ghi kết quả cuối cùng.
mi = mi_next
Sau khi cập nhật tổng, lưu max mới để xử lý block tiếp theo.
Giả sử một hàng được chia làm hai block.
Block thứ nhất:
[
[1,2]
]
Ta có:
[
m_{\text{old}}=2
]
e^{-1}+1
]
Xử lý block thứ hai:
[
[3,4]
]
Max mới:
[
m_{\text{next}}=4
]
Phần tổng cũ phải được chuyển từ hệ quy chiếu 2 sang hệ quy chiếu 4:
d_{\text{old}}e^{2-4}
]
Phần block mới:
[
e^{3-4}+e^{4-4}
]
Tổng cuối:
(e^{-1}+1)e^{-2}+e^{-1}+1
]
Rút gọn:
e^{-3}+e^{-2}+e^{-1}+1
]
Đúng bằng:
[
\sum_{x\in[1,2,3,4]}e^{x-4}
]
Sau vòng lặp đầu:
for tile_n in hl.tile(...):
ta có, với mỗi hàng:
mi = giá trị lớn nhất của toàn bộ hàng
di = tổng exp(x - mi) của toàn bộ hàng
Tức là đã có đầy đủ:
Nhưng chưa ghi kết quả đầu ra.
for tile_n in hl.tile(n, block_size=block_size_n):
Kernel duyệt lại từng block cột.
values = x[tile_m, tile_n]
Đọc lại dữ liệu đầu vào.
Sau đó:
out[tile_m, tile_n] = (
torch.exp(values - mi[:, None])
/ di[:, None]
)
Đây chính là công thức:
\frac{e^{x-m}}{d}
]
với:
[
d=\sum_j e^{x_j-m}
]
Broadcasting:
values [block_rows, block_cols]
mi[:, None] [block_rows, 1]
di[:, None] [block_rows, 1]
Mỗi hàng dùng một giá trị max và một mẫu số riêng.
Đọc x
→ tìm max
→ tính tổng exp đã chuẩn hóa
Đọc lại x
→ tính exp(x - max) / tổng
→ ghi out
Sơ đồ:
Lượt 1
x ───────────────────────────────────┐
│
max từng hàng + mẫu số │
▼
mi và di
│
Lượt 2 │
x ───────────────────────────────────┤
▼
exp(x-mi) / di
│
▼
out
| Tiêu chí | softmax | softmax_two_pass |
|---|---|---|
| Mức trừu tượng | Cao | Thấp hơn |
| Cách viết | Gọi PyTorch Softmax | Tự triển khai công thức |
| Chia tile theo hàng | Có | Có |
| Chia tile theo cột | Helion tự xử lý | Viết rõ trong code |
| Ổn định số học | Do PyTorch/Inductor xử lý | Thể hiện rõ bằng mi, di |
| Hỗ trợ hàng rất rộng | Phụ thuộc compiler | Chủ động loop theo block cột |
| Dễ đọc | Dễ hơn | Khó hơn |
| Kiểm soát thuật toán | Ít hơn | Nhiều hơn |
Một hàng có 2560 phần tử float16.
Dung lượng thô:
[
2560\times2=5120\text{ byte}
]
Nhưng trong quá trình tính Softmax, nhiều giá trị thường được chuyển sang float32.
Nếu cố giữ toàn bộ dữ liệu trung gian của nhiều hàng trong register, áp lực register có thể cao.
Kernel hai lượt cho phép xử lý chiều cột theo từng phần:
256 cột
→ 256 cột
→ 256 cột
→ ...
Thay vì bắt buộc xử lý toàn bộ 2560 cột cùng lúc.
Lợi ích tiềm năng:
Đổi lại, đầu vào phải được đọc hai lần.
Mặc dù trong softmax_two_pass có ba vòng lặp:
for tile_m in hl.tile(...):
for tile_n in hl.tile(...):
...
for tile_n in hl.tile(...):
...
không có nghĩa là tạo ba lần gọi kernel GPU.
Ý nghĩa gần đúng:
| Vòng lặp | Cách Helion ánh xạ |
|---|---|
tile_m bên ngoài | Ánh xạ lên grid/program GPU |
tile_n lượt 1 | Vòng lặp nội bộ trong mỗi program |
tile_n lượt 2 | Vòng lặp nội bộ thứ hai trong cùng program |
Một program GPU xử lý một tile hàng:
Program 0 → một nhóm hàng
Program 1 → nhóm hàng tiếp theo
Program 2 → nhóm hàng tiếp theo
...
Trong mỗi program:
lặp qua các block cột lần 1
lặp qua các block cột lần 2
Trong nội dung bạn gửi, phần sau xuất hiện hai lần:
from __future__ import annotations
import torch
import helion
from helion._testing import run_example
import helion.language as hl
Tiêu đề chuỗi tài liệu cũng bị lặp:
"""
Helion Softmax Kernel Examples
...
"""
Đây không phải lỗi nghiêm trọng, nhưng có thể gộp thành một file rõ ràng hơn.
from __future__ import annotations
import torch
import helion
import helion.language as hl
from helion._testing import run_example
@helion.kernel(autotune_effort="quick")
def softmax(x: torch.Tensor) -> torch.Tensor:
num_rows, _num_cols = x.size()
out = torch.empty_like(x)
for row_tile in hl.tile(num_rows):
out[row_tile, :] = torch.nn.functional.softmax(
x[row_tile, :],
dim=1,
)
return out
@helion.kernel(autotune_effort="quick")
def softmax_two_pass(x: torch.Tensor) -> torch.Tensor:
num_rows, num_cols = x.size()
out = torch.empty_like(x)
row_block_size = hl.register_block_size(num_rows)
col_block_size = hl.register_block_size(num_cols)
for row_tile in hl.tile(
num_rows,
block_size=row_block_size,
):
running_max = hl.full(
[row_tile],
float("-inf"),
dtype=torch.float32,
)
running_sum = hl.zeros(
[row_tile],
dtype=torch.float32,
)
# Lượt 1: tìm max và mẫu số.
for col_tile in hl.tile(
num_cols,
block_size=col_block_size,
):
values = x[row_tile, col_tile]
local_max = torch.amax(values, dim=1)
next_max = torch.maximum(running_max, local_max)
running_sum = (
running_sum
* torch.exp(running_max - next_max)
+ torch.exp(
values - next_max[:, None]
).sum(dim=1)
)
running_max = next_max
# Lượt 2: chuẩn hóa và ghi đầu ra.
for col_tile in hl.tile(
num_cols,
block_size=col_block_size,
):
values = x[row_tile, col_tile]
out[row_tile, col_tile] = (
torch.exp(values - running_max[:, None])
/ running_sum[:, None]
)
return out
def check(kernel, rows: int, cols: int) -> None:
x = torch.randn(
[rows, cols],
device="cuda",
dtype=torch.float16,
)
run_example(
kernel,
lambda value: torch.nn.functional.softmax(value, dim=1),
(x,),
)
def main() -> None:
check(softmax, 4096, 2560)
check(softmax_two_pass, 4096, 2560)
if __name__ == "__main__":
main()
Input x
↓
Chia theo nhóm hàng
↓
PyTorch softmax trên từng nhóm hàng
↓
Output
Input x
↓
Chia theo nhóm hàng
↓
Với mỗi nhóm hàng:
├─ Chia chiều cột thành block
├─ Lượt 1: tính running max và running sum
├─ Lượt 2: tính xác suất Softmax
└─ Ghi output
Ý tưởng quan trọng nhất của softmax_two_pass là công thức cập nhật online:
\max(m_{\text{old}},m_{\text{block}})
]
d_{\text{old}}e^{m_{\text{old}}-m_{\text{new}}}
+
\sum_{\text{block}}e^{x-m_{\text{new}}}
]
Công thức này giúp tính Softmax ổn định số học khi chiều cột lớn mà không cần giữ toàn bộ hàng trong register cùng lúc.
Trong bối cảnh của cơ quan Thuế, NVDTPC là tên viết tắt của:
Nghiệp vụ – Dự toán – Pháp chế
Đây là một phòng chuyên môn của Thuế tỉnh (trước đây là Cục Thuế tỉnh), có chức năng tham mưu và thực hiện các nhiệm vụ như:
Việc nhiều văn bản hiện nay có ký hiệu như:
cho thấy NVDTPC là mã của Phòng Nghiệp vụ – Dự toán – Pháp chế trong cơ cấu tổ chức của cơ quan Thuế. (Văn bản pháp luật)
Vì vậy, dòng ghi:
Nguyễn Thị Phương Linh – Phòng NVDTPC
được hiểu là:
Nguyễn Thị Phương Linh – Phòng Nghiệp vụ – Dự toán – Pháp chế (Thuế tỉnh ABC).
Trên website của Tổng cục Thuế tại https://www.gdt.gov.vn/wps/portal, dòng chữ “Đã kết nối EMC” không liên quan đến hóa đơn điện tử (eInvoice) hay máy tính tiền, mà có nghĩa là:
Website đã được kết nối với Hệ thống EMC (Hệ thống giám sát, đo lường mức độ cung cấp và sử dụng dịch vụ Chính phủ số) của cơ quan quản lý nhà nước.
EMC là hệ thống do Bộ Thông tin và Truyền thông (nay thuộc cơ quan quản lý về chuyển đổi số của Chính phủ) triển khai để:
Khi bạn thấy biểu tượng hoặc dòng chữ “Đã kết nối EMC”, điều đó chỉ có nghĩa là:
Đây là một thông báo dành cho việc quản trị và tuân thủ quy định của cơ quan nhà nước, không ảnh hưởng đến người dân hoặc doanh nghiệp khi sử dụng website.
Không.
Dòng chữ này không có nghĩa là:
Nó chỉ phản ánh trạng thái kỹ thuật của chính website gdt.gov.vn.
Vì vậy, nếu bạn nhìn thấy dòng “Đã kết nối EMC” ở cuối trang hoặc trên giao diện của cổng thông tin Tổng cục Thuế thì có thể hiểu đơn giản là:
Website của Tổng cục Thuế đã kết nối với Hệ thống EMC để phục vụ công tác giám sát và đo lường hiệu quả cung cấp dịch vụ công của cơ quan nhà nước.
Giao diện

và

và

và

và

và

và

và

và

và

Lỗi G12.18 khi xác thực sinh trắc học trên eTax Mobile hiện chưa thấy cơ quan thuế công bố bảng giải thích mã lỗi chi tiết. Trường hợp này thường liên quan đến việc dữ liệu người đại diện, VNeID hoặc phiên xác thực giữa eTax Mobile và hệ thống định danh chưa khớp.
Bạn xử lý lần lượt như sau:
Nếu vẫn lỗi, bạn nên chụp đầy đủ màn hình có mã G12.18, kèm:
Sau đó gửi cho bộ phận hỗ trợ thuế điện tử hoặc cơ quan thuế quản lý trực tiếp để họ kiểm tra trạng thái đồng bộ dữ liệu. Không gửi ảnh CCCD, mật khẩu hay mã OTP vào các nhóm hỗ trợ không chính thức.

https://www.gdt.gov.vn/wps/portal/Home/etax-mobile
01/TB-ĐKĐT

Nguồn: https://www.microsoft.com/en-us/corporate-responsibility/topics/open-weight/
Ngày 24 tháng 7 năm 2026
Vào những năm 1980, những người tiên phong trong lĩnh vực phần mềm mã nguồn mở đã thách thức quan niệm phổ biến rằng phần mềm chỉ có thể phát triển nếu các công ty kiểm soát chặt chẽ mã nguồn của mình. Phong trào này đã thúc đẩy một hệ sinh thái minh bạch, nơi các nhà phát triển trên toàn thế giới có thể nghiên cứu, sửa đổi và cải tiến phần mềm. Phần mềm được phát triển bởi cộng đồng mã nguồn mở hiện đang hỗ trợ hầu hết internet và là nền tảng của các hệ thống được sử dụng bởi các công ty công nghệ lớn nhất thế giới, cũng như quân đội Hoa Kỳ và các cơ quan liên bang tiến hành nghiên cứu khoa học, an ninh mạng và các nhiệm vụ quan trọng khác. Mã nguồn mở không chỉ làm giảm chi phí phần mềm; nó còn tạo ra một nền tảng kiến thức chung mà trên đó các thế hệ kỹ sư và doanh nhân Mỹ đã xây dựng nên chủ quyền của các tổ chức của họ.
Hoa Kỳ hiện đang đối mặt với một lựa chọn tương tự với trí tuệ nhân tạo. Khả năng dẫn đầu về AI của chúng ta sẽ được đánh giá không phải bởi một mô hình AI tiên tiến duy nhất, mà bởi việc Hoa Kỳ có xây dựng được một hệ sinh thái mạnh mẽ, cởi mở, lan tỏa đến mọi lĩnh vực hay không. Điều này rất cần thiết để tạo ra cơ hội đổi mới và thịnh vượng trên khắp đất nước. Nó đòi hỏi phải mở rộng khả năng tiếp cận AI, khuyến khích cạnh tranh, xây dựng các lớp ứng dụng mạnh mẽ và trao cho người dân Mỹ quyền kiểm soát lớn hơn đối với công nghệ mà họ đang sử dụng. Các mô hình mở – các mô hình AI mà bất kỳ ai cũng có thể tải xuống, kiểm tra, sửa đổi và chạy trên cơ sở hạ tầng của riêng họ – là một phần quan trọng của nền tảng đó bởi vì chúng giúp AI tiên tiến dễ tiếp cận hơn, dễ thích ứng hơn và phổ biến rộng rãi hơn.
Chính sách trọng số mở (open weights) giúp mở rộng khả năng tiếp cận nền kinh tế AI. Các công ty khởi nghiệp, doanh nghiệp lâu đời, trường đại học và các tổ chức công cộng có thể xây dựng dựa trên các mô hình tiên tiến mà không cần phải đào tạo từ đầu hoặc trả giá cao cho các mô hình tiên tiến cho mỗi tác vụ. Chính sách trọng số mở cho phép mọi tổ chức lựa chọn mô hình phù hợp với công việc phù hợp với chi phí hợp lý, dành khả năng ở quy mô tiên tiến cho các vấn đề thực sự cấp thiết và vận hành các mô hình chuyên biệt hiệu quả ở mọi nơi khác. Chính sự kỷ luật đó sẽ giúp AI bền vững về mặt kinh tế khi việc sử dụng nó mở rộng đến hàng tỷ tác vụ hàng ngày. Nước Mỹ sẽ chiến thắng kỷ nguyên AI bằng cách phổ biến nó vào quy trình làm việc của các nhà máy, bệnh viện, trang trại, lớp học và các doanh nghiệp nhỏ.
Việc cho phép sử dụng rộng rãi các thuật toán xử lý dữ liệu cũng thúc đẩy cạnh tranh, và chính cạnh tranh là yếu tố giúp phân bổ rộng rãi lợi ích của AI thay vì tập trung vào một số ít người. Bằng cách cho phép nhiều tổ chức xây dựng, điều chỉnh và triển khai các mô hình tiên tiến, việc cho phép sử dụng rộng rãi các thuật toán xử lý dữ liệu tạo ra sự cạnh tranh không chỉ giữa các nhà phát triển mô hình mà còn giữa các chip đám mây, ứng dụng và dịch vụ. Sự cạnh tranh đó thúc đẩy đổi mới, giảm chi phí và phân phối rộng rãi lợi ích của AI trong toàn bộ nền kinh tế.
Mô hình trọng số mở cũng mang lại cho khách hàng quyền kiểm soát lớn hơn. Khi các tổ chức đầu tư vào AI, họ muốn biết rằng họ sẽ không bị ràng buộc vào một nhà cung cấp duy nhất hoặc mất đi kiến thức và khả năng mà họ xây dựng được theo thời gian. Mô hình trọng số mở giúp cung cấp sự đảm bảo đó bằng cách cho phép các tổ chức kiểm soát dữ liệu của riêng họ, đánh giá và điều chỉnh các mô hình cho phù hợp với nhu cầu của riêng họ, và triển khai chúng ở bất cứ nơi nào yêu cầu kinh doanh của họ đòi hỏi. Và khi các tổ chức tạo ra giá trị với AI, trọng số mở cho phép họ sở hữu giá trị đó thông qua các mô hình tự cải tiến, khả năng chuyên biệt và kiến thức tích lũy thúc đẩy chủ quyền và sự thịnh vượng của Hoa Kỳ.
Chắc chắn, các mô hình trọng số mở tiềm ẩn những rủi ro thực sự và riêng biệt. Sau khi được công bố, các mô hình trọng số này nằm ngoài tầm kiểm soát của nhà phát triển ban đầu, và các phiên bản đã chỉnh sửa rất khó truy vết hoặc đảo ngược. Nhưng phản ứng đúng đắn đối với rủi ro này không phải là cấm các mô hình trọng số mở. Trong một thế giới mà các kẻ tấn công an ninh mạng sử dụng trí tuệ nhân tạo tiên tiến, những người bảo vệ cần quyền truy cập vào các mô hình có khả năng tương đương để họ có thể phát hiện, mô phỏng và ứng phó với các mối đe dọa mới nổi. Các mô hình mở mở rộng khả năng phòng thủ, tăng tính minh bạch và cho phép phát hiện và khắc phục các lỗ hổng trên nhiều nhóm.
Trên thực tế, tính mở có thể là một trong những con đường quan trọng nhất để đảm bảo an toàn và bảo mật cho AI. Việc chỉ dựa vào các mô hình đóng không tự nhiên là an toàn: chúng có thể bị xâm phạm, lạm dụng hoặc gặp lỗi mà người ngoài không thể phát hiện. Và việc tập trung các khả năng AI tiên tiến vào một số ít mô hình đóng sẽ làm tăng thêm rủi ro đó. Nó dẫn đến một số ít điểm yếu duy nhất, làm suy yếu cạnh tranh và để lại công nghệ quan trọng trong tay một vài nhà cung cấp. Mặt khác, các mô hình mở cho phép một cộng đồng rộng lớn các nhà nghiên cứu và nhà phát triển kiểm tra hành vi của chúng, xác định các lỗ hổng, phát triển các biện pháp bảo vệ và cải thiện chúng theo thời gian. Giống như phần mềm mã nguồn mở đã chứng minh rằng tính minh bạch có thể an toàn hơn tính che giấu, an toàn của AI có thể phụ thuộc vào việc trao cho nhiều người hơn khả năng kiểm tra và củng cố các mô hình mà xã hội đang dựa vào. Nó cho phép đánh giá và so sánh nghiêm ngặt, tấn công giả lập (red teaming) và các biện pháp bảo vệ gắn liền với những thiệt hại thực tế và đã được chứng minh, thay vì giả định rằng các hệ thống đóng an toàn hơn theo mặc định.
Một hệ sinh thái AI mạnh mẽ không phải là điều chắc chắn sẽ xảy ra. Các nhà hoạch định chính sách có một cơ hội quan trọng để hành động. Điều này bao gồm việc mở rộng khả năng tiếp cận điện toán cho các công ty khởi nghiệp và các nhà nghiên cứu, đầu tư vào các tài sản đào tạo dùng chung (bộ dữ liệu, công cụ, khung đánh giá) và duy trì tính đa dạng của lĩnh vực này bằng cách tránh những hạn chế sớm đối với các mô hình mở làm kìm hãm cạnh tranh hoặc đẩy đổi mới ra nước ngoài. Các biện pháp này cũng cần xem xét cách các lớp ứng dụng mạnh mẽ có thể mở rộng việc sử dụng AI trong toàn bộ nền kinh tế.
Trong việc định hình hệ sinh thái này, các nhà hoạch định chính sách cần cẩn trọng để không nhầm lẫn các kỹ thuật phát triển mô hình hợp pháp với hành vi chiếm đoạt bất hợp pháp. Việc chắt lọc, hay thực hành sử dụng đầu ra của một mô hình để giúp huấn luyện hoặc cải thiện một mô hình khác, là một kỹ thuật được sử dụng rộng rãi để cải thiện, đánh giá và xác thực mô hình. Nó phản ánh một truyền thống lâu đời về việc học hỏi, xây dựng và cải tiến các công nghệ hiện có, một truyền thống đã giúp thúc đẩy sự đổi mới kể từ khi phong trào phần mềm mã nguồn mở ra đời. Ngược lại, những nỗ lực bất hợp pháp nhằm khai thác giá trị từ các mô hình đóng gây ra những lo ngại chính đáng. Những lo ngại đó cần được giải quyết thông qua các khuôn khổ pháp lý và thương mại có mục tiêu cụ thể hơn là các hạn chế toàn diện đối với các kỹ thuật đóng vai trò quan trọng trong đổi mới trí tuệ nhân tạo.
Kỷ nguyên trí tuệ nhân tạo (AI) có thể là kỷ nguyên thịnh vượng. Với những lựa chọn đúng đắn, AI mở có thể mở rộng cơ hội, tăng cường cạnh tranh, củng cố vị thế dẫn đầu về công nghệ của Mỹ, giảm thiểu rủi ro và đảm bảo rằng lợi ích của công nghệ phi thường này được chia sẻ rộng rãi trong toàn bộ nền kinh tế. Tương lai đó đáng để xây dựng, và Hoa Kỳ nên dẫn đầu trong việc xây dựng nó.
Mạng lưới các nhà đổi mới Mỹ ● Andreessen Horowitz ● Arcee AI ● Arena ● Black Forest Labs ● Box ● Cisco ● Cohere ● CrowdStrike ● Dell Technologies ● DoorDash ● Emergence Capital ● Fireworks AI ● GitHub ● Hugging Face ● IBM ● The Linux Foundation ● Mariana Minerals ● Meta ● Microsoft ● Mistral ● Mozilla ● Nous Research ● NVIDIA ● OpenAI ● OpenClaw ● Palantir ● Palo Alto Networks ● Perplexity ● Prime Intellect ● Reflection ● Replit ● ServiceNow ● Telnyx ● Y Combinator
Thông báo về phương pháp trích khấu hao TSCĐ/ Thông báo
về việc chuyển lợi nhuận ra nước ngoài từ 22/7/2026 không phải nộp file giấy nữa.
Dùng mã thủ tục hành chính công 1.008593 để thực hiện việc này.
Group zalo: https://zalo.me/g/ekqioxaksfhiqwbd6dl1
Group Facebook hỗ trợ kê khai thuế, kế toán, phần mềm kế toán https://www.facebook.com/groups/238330069184113
Nhóm tư vấn, hỗ trợ kế toán, thuế BKIT.VN
Xem video agent hoạt động với OpenAI (ChatGPT) Codex (beta) tại https://ketoan.bkit.vn/2026/07/07/agent-skill-cho-phan-mem-ke-toan-bkit/ (bấm nút play video ở cuối trang). Streaming trực tiếp từ hệ thống https://file.bkit.vn không qua Youtube.
Mua Windows 11 không crack tại https://shop.bkit.vn/
Hóa đơn VAT + Cài đặt và kích hoạt (rất phức tạp, BKIT làm miễn phí) + Hợp đồng kinh tế + Biên bản giao nhận
Từ 10/9, doanh nghiệp chậm đóng BHTN có thể bị phạt đến 75 triệu đồng và vẫn phải truy đóng
Từ Nghị định số 283/2026/NĐ-CP của Chính phủ, xin nhắc nhở anh chị em kế toán, đặc biệt khối doanh nghiệp FDI, cần lưu ý/kiểm soát chặt chẽ tiến độ trích nộp bảo hiểm thất nghiệp. Nếu lỡ vướng phải, doanh nghiệp sẽ đối mặt với các rủi ro sau:
Địa chỉ quản lý bản quyền Windows 11 pro

Mua bản quyền phần mềm Windows 11 pro tại https://bkit.vn/info/tuan-thu-ban-quyen