当前位置: 首页 > news >正文

【AI编译器】triton学习:矩阵乘优化

Matrix Multiplication¶

主要内容:

块级矩阵乘法

多维指针算术

重新编排程序以提升L2缓存命

自动性能调整

Motivations¶

矩阵乘法是当今高性能计算系统的一个关键组件,在大多数情况下被用于构建硬件。由于该操作特别复杂,因此通常由软件提供商来进行实现,而不是使用者自己手动编写代码。这些库称为“内核库”(例如cuBLAS),可能会存在一定的版权限制,并且通常无法随意修改以适用于深度学习工作负载中的特殊需求(例如融合式活动函数)。本教程将带你了解如何使用 Triton 自行编写稳定、可扩展的矩阵乘法函数,以实现高性能计算系统所需的功能。

大概来说,我们将要写的内核实现下面这种固定步长算法,对一个(M、K)矩阵和一个(K、N)矩阵进行乘法运算:

# Do in parallel
for m in range(0, M, BLOCK_SIZE_M):# Do in parallelfor n in range(0, N, BLOCK_SIZE_N):acc = zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=float32)for k in range(0, K, BLOCK_SIZE_K):a = A[m : m+BLOCK_SIZE_M, k : k+BLOCK_SIZE_K]b = B[k : k+BLOCK_SIZE_K, n : n+BLOCK_SIZE_N]acc += dot(a, b)C[m : m+BLOCK_SIZE_M, n : n+BLOCK_SIZE_N] = acc

在每个嵌套程序的每次递归过程中,Triton 实例都会产生一个专门的计算机程序来执行。

Compute Kernel¶

上面算法的实现在Triton上其实并不困难。只是内部迭代过程中,对于要从A和B缓存区读取数据的记忆位置计算这一步骤会有些复杂。为了实现该操作,我们需要使用多维指针运算方式。

指针数学运算符号¶

如果 X 是一个行向矩阵,那么X[i,j]的存储位址为 &X[i,j] = X + istride_xi + jstride_xj。因此,可以将A[m:m+BLOCK_SIZE_M, k:k+BLOCK_SIZE_K]和 B[k:k+BLOCK_SIZE_K, n:n+BLOCK_SIZE_N]的块中存储地址表示成以下形式的虚拟代码:

&A[m : m+BLOCK_SIZE_M, k:k+BLOCK_SIZE_K] =  a_ptr + (m : m+BLOCK_SIZE_M)[:, None]*A.stride(0) + (k : k+BLOCK_SIZE_K)[None, :]*A.stride(1);
&B[k : k+BLOCK_SIZE_K, n:n+BLOCK_SIZE_N] =  b_ptr + (k : k+BLOCK_SIZE_K)[:, None]*B.stride(0) + (n : n+BLOCK_SIZE_N)[None, :]*B.stride(1);

这意味着在Triton软件中,我们可以将A和B两个向量的指针分段设置为0(即i=0)。且需要注意的是,当M与数据握的大小BLOCK_SIZE_M不是相匹配的时候,我们可以通过添加一个额外模式来处理这种情况,例如,在数据中往底部加上一些无用的值。此后我们将为K维度使用遮蔽式载入操作进行处理。

offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
offs_k = tl.arange(0, BLOCK_SIZE_K)
a_ptrs = a_ptr + (offs_am[:, None]*stride_am + offs_k [None, :]*stride_ak)
b_ptrs = b_ptr + (offs_k [:, None]*stride_bk + offs_bn[None, :]*stride_bn)

然后在内部循环中进行以下更新:

a_ptrs += BLOCK_SIZE_K * stride_ak;
b_ptrs += BLOCK_SIZE_K * stride_bk;

二级缓存的优化措施

如上所述,每个程序实例都会计算出一组长度为[BLOCK_SIZE_M,BLOCK_SIZE_N]的C语言代码块。这种方式非常重要,因为执行顺序可能导致该程序中L2缓存的命中率不同,而且令人遗憾的是,如果我们使用矩阵增量式顺序执行代码,则其性能将会受到影响。

pid = tl.program_id(axis=0)
grid_n = tl.cdiv(N, BLOCK_SIZE_N)
pid_m = pid // grid_n
pid_n = pid % grid_n

这根本不够用。

解决这个问题的一种方法是在排列块时采用数据重复使用的最佳组合。 我们可以“将所有 GROUP_M 行都集成到同一块中”,然后再按照对应的列进行排序:

# Program ID
pid = tl.program_id(axis=0)
# Number of program ids along the M axis
num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)
# Number of programs ids along the N axis
num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)
# Number of programs in group
num_pid_in_group = GROUP_SIZE_M * num_pid_n
# Id of the group this program is in
group_id = pid // num_pid_in_group
# Row-id of the first program in the group
first_pid_m = group_id * GROUP_SIZE_M
# If `num_pid_m` isn't divisible by `GROUP_SIZE_M`, the last group is smaller
group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
# *Within groups*, programs are ordered in a column-major order
# Row-id of the program in the *launch grid*
pid_m = first_pid_m + ((pid % num_pid_in_group) % group_size_m)
# Col-id of the program in the *launch grid*
pid_n = (pid % num_pid_in_group) // group_size_m

以下的矩阵乘法是每个矩阵块为九个,共有九个矩阵乘法操作。通过比较如果我们按行向量序排列输出的话,则需要在SRAM中加载90个元素来计算第一层的9个输出值,而若是以固定单元格为基础进行分组操作,只需加载54个元素。



实际上,这会使我们矩阵乘法的算法执行效率提高超过10%(例如:A100在220至245TFLOPS之间)。

Final Result¶

import torchimport triton
import triton.language as tldef is_cuda():return triton.runtime.driver.active.get_current_target().backend == "cuda"def is_hip_mi200():target = triton.runtime.driver.active.get_current_target()return target.backend == 'hip' and target.arch == 'gfx90a'def get_cuda_autotune_config():return [triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8}, num_stages=3,num_warps=8),triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=4,num_warps=4),triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=4,num_warps=4),triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=4,num_warps=4),triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=4,num_warps=4),triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 32, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=4,num_warps=4),triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 32, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=5,num_warps=2),triton.Config({'BLOCK_SIZE_M': 32, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=5,num_warps=2),# Good config for fp8 inputs.triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 8}, num_stages=3,num_warps=8),triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 8}, num_stages=3,num_warps=8),triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 8}, num_stages=4,num_warps=4),triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 8}, num_stages=4,num_warps=4),triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 8}, num_stages=4,num_warps=4),triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8}, num_stages=4,num_warps=4),triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8}, num_stages=4,num_warps=4),triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 32, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8}, num_stages=4,num_warps=4)]def get_hip_autotune_config():return [triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 16, 'GROUP_SIZE_M': 1, 'waves_per_eu': 2},num_warps=4, num_stages=0),triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 16, 'GROUP_SIZE_M': 4, 'waves_per_eu': 2},num_warps=8, num_stages=0),triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 1, 'waves_per_eu': 2},num_warps=8, num_stages=0),triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8, 'waves_per_eu': 3},num_warps=4, num_stages=0),triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 1, 'waves_per_eu': 8},num_warps=4, num_stages=0),]def get_autotune_config():if is_cuda():return get_cuda_autotune_config()else:return get_hip_autotune_config()# `triton.jit`'ed functions can be auto-tuned by using the `triton.autotune` decorator, which consumes:
#   - A list of `triton.Config` objects that define different configurations of
#       meta-parameters (e.g., `BLOCK_SIZE_M`) and compilation options (e.g., `num_warps`) to try
#   - An auto-tuning *key* whose change in values will trigger evaluation of all the
#       provided configs
@triton.autotune(configs=get_autotune_config(),key=['M', 'N', 'K'],
)
@triton.jit
def matmul_kernel(# Pointers to matricesa_ptr, b_ptr, c_ptr,# Matrix dimensionsM, N, K,# The stride variables represent how much to increase the ptr by when moving by 1# element in a particular dimension. E.g. `stride_am` is how much to increase `a_ptr`# by to get the element one row down (A has M rows).stride_am, stride_ak,  #stride_bk, stride_bn,  #stride_cm, stride_cn,# Meta-parametersBLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr,  #GROUP_SIZE_M: tl.constexpr,  #ACTIVATION: tl.constexpr  #
):"""Kernel for computing the matmul C = A x B.A has shape (M, K), B has shape (K, N) and C has shape (M, N)"""# -----------------------------------------------------------# Map program ids `pid` to the block of C it should compute.# This is done in a grouped ordering to promote L2 data reuse.# See above `L2 Cache Optimizations` section for details.pid = tl.program_id(axis=0)num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)num_pid_in_group = GROUP_SIZE_M * num_pid_ngroup_id = pid // num_pid_in_groupfirst_pid_m = group_id * GROUP_SIZE_Mgroup_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)pid_m = first_pid_m + ((pid % num_pid_in_group) % group_size_m)pid_n = (pid % num_pid_in_group) // group_size_m# ----------------------------------------------------------# Create pointers for the first blocks of A and B.# We will advance this pointer as we move in the K direction# and accumulate# `a_ptrs` is a block of [BLOCK_SIZE_M, BLOCK_SIZE_K] pointers# `b_ptrs` is a block of [BLOCK_SIZE_K, BLOCK_SIZE_N] pointers# See above `Pointer Arithmetic` section for detailsoffs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % Moffs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % Noffs_k = tl.arange(0, BLOCK_SIZE_K)a_ptrs = a_ptr + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)b_ptrs = b_ptr + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)# -----------------------------------------------------------# Iterate to compute a block of the C matrix.# We accumulate into a `[BLOCK_SIZE_M, BLOCK_SIZE_N]` block# of fp32 values for higher accuracy.# `accumulator` will be converted back to fp16 after the loop.accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):# Load the next block of A and B, generate a mask by checking the K dimension.# If it is out of bounds, set it to 0.a = tl.load(a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K, other=0.0)b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0)# We accumulate along the K dimension.accumulator = tl.dot(a, b, accumulator)# Advance the ptrs to the next K block.a_ptrs += BLOCK_SIZE_K * stride_akb_ptrs += BLOCK_SIZE_K * stride_bk# You can fuse arbitrary activation functions here# while the accumulator is still in FP32!if ACTIVATION == "leaky_relu":accumulator = leaky_relu(accumulator)c = accumulator.to(tl.float16)# -----------------------------------------------------------# Write back the block of the output matrix C with masks.offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)c_ptrs = c_ptr + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)tl.store(c_ptrs, c, mask=c_mask)# We can fuse `leaky_relu` by providing it as an `ACTIVATION` meta-parameter in `matmul_kernel`.
@triton.jit
def leaky_relu(x):return tl.where(x >= 0, x, 0.01 * x)

我们现在可以创建一个操作简单的函数包装器,具有两个输入张量参数(), 该函数:(1)检查输入张量是否符合任何形状要求;(2)分配输出;(3)调用上面所示的核心计算。

def matmul(a, b, activation=""):# Check constraints.assert a.shape[1] == b.shape[0], "Incompatible dimensions"assert a.is_contiguous(), "Matrix A must be contiguous"M, K = a.shapeK, N = b.shape# Allocates output.c = torch.empty((M, N), device=a.device, dtype=torch.float16)# 1D launch kernel where each block gets its own program.grid = lambda META: (triton.cdiv(M, META['BLOCK_SIZE_M']) * triton.cdiv(N, META['BLOCK_SIZE_N']), )matmul_kernel[grid](a, b, c,  #M, N, K,  #a.stride(0), a.stride(1),  #b.stride(0), b.stride(1),  #c.stride(0), c.stride(1),  #ACTIVATION=activation  #)return c

Unit Test¶

我们可以对自定义的矩阵乘法运算与原生执行操作的Torch进行测试(即cuBLAS)。

torch.manual_seed(0)
a = torch.randn((512, 512), device='cuda', dtype=torch.float16)
b = torch.randn((512, 512), device='cuda', dtype=torch.float16)
triton_output = matmul(a, b)
torch_output = torch.matmul(a, b)
print(f"triton_output_with_fp16_inputs={triton_output}")
print(f"torch_output_with_fp16_inputs={torch_output}")
# Bigger tolerance for AMD MI200 devices.
# MI200 devices use reduced precision fp16 and bf16 and flush input and
# output denormal values to zero. Detailed info is at: https://pytorch.org/docs/stable/notes/numerical_accuracy.html#reduced-precision-fp16-and-bf16-gemms-and-convolutions-on-amd-instinct-mi200-devices
rtol = 1e-2 if is_hip_mi200() else 0
if torch.allclose(triton_output, torch_output, atol=1e-2, rtol=rtol):print("✅ Triton and Torch match")
else:print("❌ Triton and Torch differ")TORCH_HAS_FP8 = hasattr(torch, "float8_e5m2")
if TORCH_HAS_FP8 and is_cuda():torch.manual_seed(0)a = torch.randn((512, 512), device="cuda", dtype=torch.float16)b = torch.randn((512, 512), device="cuda", dtype=torch.float16)a = a.to(torch.float8_e5m2)# pre-transpose b for efficiency.b = b.Tb = b.to(torch.float8_e5m2)triton_output = matmul(a, b)torch_output = torch.matmul(a.to(torch.float16), b.to(torch.float16))print(f"triton_output_with_fp8_inputs={triton_output}")print(f"torch_output_with_fp8_inputs={torch_output}")if torch.allclose(triton_output, torch_output, atol=0.125, rtol=0):print("✅ Triton and Torch match")else:print("❌ Triton and Torch differ")
triton_output_with_fp16_inputs=tensor([[-10.9531,  -4.7109,  15.6953,  ..., -28.4062,   4.3320, -26.4219],[ 26.8438,  10.0469,  -5.4297,  ..., -11.2969,  -8.5312,  30.7500],[-13.2578,  15.8516,  18.0781,  ..., -21.7656,  -8.6406,  10.2031],...,[ 40.2812,  18.6094, -25.6094,  ...,  -2.7598,  -3.2441,  41.0000],[ -6.1211, -16.8281,   4.4844,  ..., -21.0312,  24.7031,  15.0234],[-17.0938, -19.0000,  -0.3831,  ...,  21.5469, -30.2344, -13.2188]],device='cuda:0', dtype=torch.float16)
torch_output_with_fp16_inputs=tensor([[-10.9531,  -4.7109,  15.6953,  ..., -28.4062,   4.3320, -26.4219],[ 26.8438,  10.0469,  -5.4297,  ..., -11.2969,  -8.5312,  30.7500],[-13.2578,  15.8516,  18.0781,  ..., -21.7656,  -8.6406,  10.2031],...,[ 40.2812,  18.6094, -25.6094,  ...,  -2.7598,  -3.2441,  41.0000],[ -6.1211, -16.8281,   4.4844,  ..., -21.0312,  24.7031,  15.0234],[-17.0938, -19.0000,  -0.3831,  ...,  21.5469, -30.2344, -13.2188]],device='cuda:0', dtype=torch.float16)
✅ Triton and Torch match
triton_output_with_fp8_inputs=tensor([[-21.4375,  13.1719,   6.0352,  ...,  28.7031,   8.6719, -40.7500],[ 10.0000,  37.0000,  -5.5664,  ...,  20.9844,  46.8125,  30.8281],[ 19.5625,  -3.0078, -20.0469,  ...,  -2.1309,  -8.0625,  12.5625],...,[-18.1562, -34.1562, -27.4219,  ..., -27.3906, -24.0938, -12.3516],[ -3.3945,  -8.6250, -23.6562,  ...,  -4.1094,  -3.5332, -16.0781],[-23.9688,  -3.2637, -33.6875,  ...,  17.3125, -36.6250,  25.8594]],device='cuda:0', dtype=torch.float16)
torch_output_with_fp8_inputs=tensor([[-21.4375,  13.1719,   6.0352,  ...,  28.7031,   8.6719, -40.7500],[ 10.0000,  37.0000,  -5.5664,  ...,  20.9844,  46.8125,  30.8281],[ 19.5625,  -3.0078, -20.0469,  ...,  -2.1309,  -8.0625,  12.5625],...,[-18.1562, -34.1562, -27.4219,  ..., -27.3906, -24.0938, -12.3516],[ -3.3945,  -8.6250, -23.6562,  ...,  -4.1094,  -3.5332, -16.0781],[-23.9688,  -3.2637, -33.6875,  ...,  17.3125, -36.6250,  25.8594]],device='cuda:0', dtype=torch.float16)
✅ Triton and Torch match

Benchmark¶

性能指标¶

我们可以对现有的内核与cuBLAS或rocBLAS进行比较,这里以对方阵为例,但也能够根据你自定义需求对其他矩阵形状进行性能测试。

ref_lib = 'cuBLAS' if is_cuda() else 'rocBLAS'configs = []
for fp8_inputs in [False, True]:if fp8_inputs and (not TORCH_HAS_FP8 or not is_cuda()):continueconfigs.append(triton.testing.Benchmark(x_names=["M", "N", "K"],  # Argument names to use as an x-axis for the plotx_vals=[128 * i for i in range(2, 33)],  # Different possible values for `x_name`line_arg="provider",  # Argument name whose value corresponds to a different line in the plot# Possible values for `line_arg`# Don't compare to cublas for fp8 cases as torch.matmul doesn't support fp8 at the moment.line_vals=["triton"] if fp8_inputs else [ref_lib.lower(), "triton"],  # Label name for the linesline_names=["Triton"] if fp8_inputs else [ref_lib, "Triton"],  # Line stylesstyles=[("green", "-"), ("blue", "-")],ylabel="TFLOPS",  # Label name for the y-axisplot_name="matmul-performance-" +("fp16" if not fp8_inputs else "fp8"),  # Name for the plot, used also as a file name for saving the plot.args={"fp8_inputs": fp8_inputs},))@triton.testing.perf_report(configs)
def benchmark(M, N, K, provider, fp8_inputs):a = torch.randn((M, K), device='cuda', dtype=torch.float16)b = torch.randn((K, N), device='cuda', dtype=torch.float16)if TORCH_HAS_FP8 and fp8_inputs:a = a.to(torch.float8_e5m2)b = b.Tb = b.to(torch.float8_e5m2)quantiles = [0.5, 0.2, 0.8]if provider == ref_lib.lower():ms, min_ms, max_ms = triton.testing.do_bench(lambda: torch.matmul(a, b), quantiles=quantiles)if provider == 'triton':ms, min_ms, max_ms = triton.testing.do_bench(lambda: matmul(a, b), quantiles=quantiles)perf = lambda ms: 2 * M * N * K * 1e-12 / (ms * 1e-3)return perf(ms), perf(max_ms), perf(min_ms)benchmark.run(show_plots=True, print_data=True)




matmul-performance-fp16:M       N       K      cuBLAS      Triton
0    256.0   256.0   256.0    4.096000    4.096000
1    384.0   384.0   384.0   12.288000   12.288000
2    512.0   512.0   512.0   26.214401   26.214401
3    640.0   640.0   640.0   42.666665   42.666665
4    768.0   768.0   768.0   63.195428   68.056616
5    896.0   896.0   896.0   78.051553   93.661869
6   1024.0  1024.0  1024.0  110.376426   99.864382
7   1152.0  1152.0  1152.0  135.726544  129.825388
8   1280.0  1280.0  1280.0  163.840004  163.840004
9   1408.0  1408.0  1408.0  155.765024  132.970149
10  1536.0  1536.0  1536.0  176.947204  157.286398
11  1664.0  1664.0  1664.0  179.978245  176.449258
12  1792.0  1792.0  1792.0  172.914215  204.353162
13  1920.0  1920.0  1920.0  200.347822  168.585369
14  2048.0  2048.0  2048.0  226.719125  190.650180
15  2176.0  2176.0  2176.0  211.827867  211.827867
16  2304.0  2304.0  2304.0  229.691080  228.592087
17  2432.0  2432.0  2432.0  203.583068  199.251522
18  2560.0  2560.0  2560.0  224.438347  218.453323
19  2688.0  2688.0  2688.0  200.704002  198.602388
20  2816.0  2816.0  2816.0  211.719459  210.696652
21  2944.0  2944.0  2944.0  218.579083  220.513412
22  3072.0  3072.0  3072.0  208.173173  208.173173
23  3200.0  3200.0  3200.0  214.046818  219.178074
24  3328.0  3328.0  3328.0  208.067338  208.973281
25  3456.0  3456.0  3456.0  217.308808  219.080343
26  3584.0  3584.0  3584.0  216.142772  211.565625
27  3712.0  3712.0  3712.0  209.428397  213.000737
28  3840.0  3840.0  3840.0  210.250955  205.179974
29  3968.0  3968.0  3968.0  213.142249  215.971570
30  4096.0  4096.0  4096.0  221.847481  215.784121
matmul-performance-fp8:M       N       K      Triton
0    256.0   256.0   256.0    3.276800
1    384.0   384.0   384.0   10.053818
2    512.0   512.0   512.0   20.164923
3    640.0   640.0   640.0   34.133334
4    768.0   768.0   768.0   42.130286
5    896.0   896.0   896.0   58.538665
6   1024.0  1024.0  1024.0   61.680940
7   1152.0  1152.0  1152.0   80.702267
8   1280.0  1280.0  1280.0  102.400003
9   1408.0  1408.0  1408.0   82.602666
10  1536.0  1536.0  1536.0   99.688560
11  1664.0  1664.0  1664.0  116.868992
12  1792.0  1792.0  1792.0  135.414749
13  1920.0  1920.0  1920.0  100.905113
14  2048.0  2048.0  2048.0  114.912434
15  2176.0  2176.0  2176.0  121.226797
16  2304.0  2304.0  2304.0  134.201527
17  2432.0  2432.0  2432.0  134.423269
18  2560.0  2560.0  2560.0  146.941707
19  2688.0  2688.0  2688.0  118.171514
20  2816.0  2816.0  2816.0  129.036114
21  2944.0  2944.0  2944.0  139.988852
22  3072.0  3072.0  3072.0  144.446699
23  3200.0  3200.0  3200.0  139.433550
24  3328.0  3328.0  3328.0  131.131689
25  3456.0  3456.0  3456.0  139.725414
26  3584.0  3584.0  3584.0  149.113421
27  3712.0  3712.0  3712.0  142.506914
28  3840.0  3840.0  3840.0  138.413021
29  3968.0  3968.0  3968.0  147.194128
30  4096.0  4096.0  4096.0  156.430916

相关文章:

【AI编译器】triton学习:矩阵乘优化

Matrix Multiplication 主要内容&#xff1a; 块级矩阵乘法 多维指针算术 重新编排程序以提升L2缓存命 自动性能调整 Motivations 矩阵乘法是当今高性能计算系统的一个关键组件&#xff0c;在大多数情况下被用于构建硬件。由于该操作特别复杂&#xff0c;因此通常由软件提…...

动静分离网络

动静分离网络的主要目的是分别处理视频帧中的静止区域和运动区域&#xff0c;以便对不同区域采用不同的去噪策略。这里提供一个实现思路&#xff0c;通过两个分支网络分别处理静止区域和运动区域&#xff0c;然后将两者的输出融合起来。 实现步骤 帧差图生成&#xff1a;计算…...

Python商务数据分析知识专栏(三)——Python数据分析的应用①Matplotlib数据可视化基础

Python商务数据分析知识专栏&#xff08;三&#xff09;——Python数据分析的应用①Matplotlib数据可视化基础 Matplotlib数据可视化基础1.掌握绘图基本语法与常用绘图2.分析特征间关系3.分析特征内部数据分布与分散情况 Matplotlib数据可视化基础 1.掌握绘图基本语法与常用绘…...

DataV大屏组件库

DataV官方文档 DataV组件库基于Vue &#xff08;React版 (opens new window)&#xff09; &#xff0c;主要用于构建大屏&#xff08;全屏&#xff09;数据展示页面即数据可视化&#xff0c;具有多种类型组件可供使用&#xff1a; 源码下载...

paraview跨节点并行渲染

参考&#xff1a; https://cloud.tencent.com/developer/ask/sof/101483588 ParaView 支持使用其内置的网络拓扑来进行跨节点的并行渲染。以下是一个简单的步骤来设置和运行跨节点的并行渲染&#xff1a; 确保你的计算环境支持多节点计算&#xff0c;比如通过SSH、MPI或其他集…...

Java中相等比较详解

本文对Java中的相等判断进行详细解释&#xff0c;包括&#xff0c;equals和compareTo等。 一、 运算符 1. 用途 基本数据类型&#xff1a;用于比较两个基本数据类型的值是否相等。 引用类型&#xff1a;用于比较两个对象引用是否指向同一个对象。 2. 示例 // 基本数据类型比…...

HBuilder X 小白日记01

1.创建项目 2.右击项目&#xff0c;可创建html文件 3.保存CtrlS&#xff0c;运行一下 我们写的内容&#xff0c;一般是写在body里面 注释的快捷键&#xff1a;Ctrl/ h标签 <h1> 定义重要等级最高的(最大)的标题。<h6> 定义最小的标题。 H标签起侧重、强调的作用…...

使用Protocol Buffers优化数据传输

使用Protocol Buffers优化数据传输 大家好&#xff0c;我是免费搭建查券返利机器人省钱赚佣金就用微赚淘客系统3.0的小编&#xff0c;也是冬天不穿秋裤&#xff0c;天冷也要风度的程序猿&#xff01; 什么是Protocol Buffers&#xff1f; Protocol Buffers&#xff08;简称P…...

如何把mkv转成mp4?介绍一下将mkv转成MP4的几种方法

如何把mkv转成mp4&#xff1f;如果你有一个MKV格式的视频文件&#xff0c;但是需要将其转换为MP4格式以便更广泛地在各种设备和平台上播放和共享&#xff0c;你可以通过进行简单的文件格式转换来实现。转换MKV到MP4格式可以提供更好的兼容性&#xff0c;并确保你的视频文件能够…...

PHP语言学习02

好久不见&#xff0c;学如逆水行舟&#xff0c;不进则退&#xff0c;真是这样。。。突然感觉自己有点废。。。 <?php phpinfo(); ?> 新生第一个代码。 要想看到运行结果&#xff0c;打开浏览器&#xff08;127.0.0.1/start/demo01.php&#xff09; 其中&#xff0c…...

PX2资料及问题记录

PX2的一些资料 官方论坛&#xff1a;https://devtalk.nvidia.com/default/board/182/drive-px2/ 官方网站&#xff1a;https://www.nvidia.com/en-us/self-driving-cars/ap2x/ 开发网站&#xff1a;https://developer.nvidia.com/drive/downloads docker docker run --devic…...

Jenkins容器的部署

本文主要是记录如何在Centos7上安装docker,以及在docker里面配置tomcat、mysql、jenkins等环境。 一、安装docker 1.1 准备工作 centos7、VMware17Pro 1.2 通过yum在线安装dokcer yum -y install docker1.3 启动docker服务 systemctl start docker.service1.4 查看docke…...

QT 自绘树形控件

资源来自:https://gitee.com/qt-open-source-collection/NavListView/blob/master/navlistview.h 1、解决的问题:一处编译报错;空白区域绘制背景;点击页面崩溃 2、源码: #ifndef NAVLISTVIEW_H #define NAVLISTVIEW_H/*** 作者:feiyangqingyun(QQ:517216493) 2016-10-1…...

axios之CancelToken取消请求

从 v0.22.0 开始&#xff0c;Axios 支持以 fetch API 方式—— AbortController 取消请求 此 API 从 v0.22.0 开始已被弃用&#xff0c;不应在新项目中使用 官网链接 1. 背景 最近项目中遇到一个场景&#xff0c;当连续触发一个请求时&#xff0c;如果是同一个接口&#xf…...

Unity | API鉴权用到的函数汇总

目录 一、HMAC-SHA1 二、UriEncode 三、Date 四、Content-MD5 五、参数操作 六、阿里云API鉴权 一、HMAC-SHA1 使用 RFC 2104 中定义的 HMAC-SHA1 方法生成带有密钥的哈希值&#xff1a; private static string CalculateSignature(string secret, string data){byte[] k…...

【python】socket通信代码解析

目录 一、socket通信原理 1.1 服务器端 1.2 客户端 二、socket通信主要应用场景 2.1 简单的服务器和客户端通信 2.2 并发服务器 2.3 UDP通信 2.4 文件传输 2.5 HTTP服务器 2.6 邮件发送与接收 2.7 FTP客户端 2.8 P2P文件共享 2.9 网络游戏 三、python中Socket编…...

FastGPT 手动部署错误:MongooseServerSelectionError: getaddrinfo EAI_AGAIN mongo

在运行 FastGPT 时&#xff0c;mongodb 报如下错误&#xff1a; MongooseServerSelectionError: getaddrinfo EAI_AGAIN mongo 这是因为 mongo 没有解析出来&#xff0c;在 hosts 文件中添加如下信息&#xff1a; 127.0.0.1 mongo 重新运行 FastGPT 即可。 参考链接&#xff…...

用英文介绍芝加哥(1):Making Modern Chicago Part 1 Building a Boomtown

Making Modern Chicago | Part 1: Building a Boomtown Link: https://www.youtube.com/watch?vpNdX0Dm-J8Y&listPLmSQiOQJmbZ7TU39cyx7gizM9i8nOuZXy&index4 Summary Summary of Chicago’s History and Development Urban Planning and Growth Chicago, often r…...

【启明智显分享】低成本RISC-V工业级HMI方案推荐

伴随着工业4.0的迅猛发展&#xff0c;工业HMI以方便、快捷的特点逐渐成为工业的日常应用&#xff0c;成为备受追捧的全新多媒体交互设备。 什么是工业HMI&#xff1f;工业HMI是用于工业自动化系统中的人机交互界面&#xff0c;通常由触摸屏、按钮、指示灯、显示器等组成&#…...

深入探索STM32的SPI功能:W25Q64 Flash存储器全攻略

摘要 随着嵌入式系统对存储需求的增长&#xff0c;选择合适的存储设备变得尤为重要。W25Q64 Flash存储器以其8Mbit的存储容量和SPI接口的便捷性&#xff0c;成为STM32微控制器项目中的优选存储方案。本文将深入探索STM32的SPI功能&#xff0c;提供W25Q64 Flash存储器的全面集成…...

MPNet:旋转机械轻量化故障诊断模型详解python代码复现

目录 一、问题背景与挑战 二、MPNet核心架构 2.1 多分支特征融合模块(MBFM) 2.2 残差注意力金字塔模块(RAPM) 2.2.1 空间金字塔注意力(SPA) 2.2.2 金字塔残差块(PRBlock) 2.3 分类器设计 三、关键技术突破 3.1 多尺度特征融合 3.2 轻量化设计策略 3.3 抗噪声…...

Linux链表操作全解析

Linux C语言链表深度解析与实战技巧 一、链表基础概念与内核链表优势1.1 为什么使用链表&#xff1f;1.2 Linux 内核链表与用户态链表的区别 二、内核链表结构与宏解析常用宏/函数 三、内核链表的优点四、用户态链表示例五、双向循环链表在内核中的实现优势5.1 插入效率5.2 安全…...

ServerTrust 并非唯一

NSURLAuthenticationMethodServerTrust 只是 authenticationMethod 的冰山一角 要理解 NSURLAuthenticationMethodServerTrust, 首先要明白它只是 authenticationMethod 的选项之一, 并非唯一 1 先厘清概念 点说明authenticationMethodURLAuthenticationChallenge.protectionS…...

Caliper 配置文件解析:config.yaml

Caliper 是一个区块链性能基准测试工具,用于评估不同区块链平台的性能。下面我将详细解释你提供的 fisco-bcos.json 文件结构,并说明它与 config.yaml 文件的关系。 fisco-bcos.json 文件解析 这个文件是针对 FISCO-BCOS 区块链网络的 Caliper 配置文件,主要包含以下几个部…...

Redis的发布订阅模式与专业的 MQ(如 Kafka, RabbitMQ)相比,优缺点是什么?适用于哪些场景?

Redis 的发布订阅&#xff08;Pub/Sub&#xff09;模式与专业的 MQ&#xff08;Message Queue&#xff09;如 Kafka、RabbitMQ 进行比较&#xff0c;核心的权衡点在于&#xff1a;简单与速度 vs. 可靠与功能。 下面我们详细展开对比。 Redis Pub/Sub 的核心特点 它是一个发后…...

并发编程 - go版

1.并发编程基础概念 进程和线程 A. 进程是程序在操作系统中的一次执行过程&#xff0c;系统进行资源分配和调度的一个独立单位。B. 线程是进程的一个执行实体,是CPU调度和分派的基本单位,它是比进程更小的能独立运行的基本单位。C.一个进程可以创建和撤销多个线程;同一个进程中…...

day36-多路IO复用

一、基本概念 &#xff08;服务器多客户端模型&#xff09; 定义&#xff1a;单线程或单进程同时监测若干个文件描述符是否可以执行IO操作的能力 作用&#xff1a;应用程序通常需要处理来自多条事件流中的事件&#xff0c;比如我现在用的电脑&#xff0c;需要同时处理键盘鼠标…...

uniapp 集成腾讯云 IM 富媒体消息(地理位置/文件)

UniApp 集成腾讯云 IM 富媒体消息全攻略&#xff08;地理位置/文件&#xff09; 一、功能实现原理 腾讯云 IM 通过 消息扩展机制 支持富媒体类型&#xff0c;核心实现方式&#xff1a; 标准消息类型&#xff1a;直接使用 SDK 内置类型&#xff08;文件、图片等&#xff09;自…...

jdbc查询mysql数据库时,出现id顺序错误的情况

我在repository中的查询语句如下所示&#xff0c;即传入一个List<intager>的数据&#xff0c;返回这些id的问题列表。但是由于数据库查询时ID列表的顺序与预期不一致&#xff0c;会导致返回的id是从小到大排列的&#xff0c;但我不希望这样。 Query("SELECT NEW com…...

32单片机——基本定时器

STM32F103有众多的定时器&#xff0c;其中包括2个基本定时器&#xff08;TIM6和TIM7&#xff09;、4个通用定时器&#xff08;TIM2~TIM5&#xff09;、2个高级控制定时器&#xff08;TIM1和TIM8&#xff09;&#xff0c;这些定时器彼此完全独立&#xff0c;不共享任何资源 1、定…...