Skip to content

[MPS Dev] 2026.07.01 - #5

Open
Manfredss wants to merge 48 commits into
mainfrom
dev_2026_07_01
Open

[MPS Dev] 2026.07.01#5
Manfredss wants to merge 48 commits into
mainfrom
dev_2026_07_01

Conversation

@Manfredss

@Manfredss Manfredss commented Jul 1, 2026

Copy link
Copy Markdown
Owner

PR Category

User Experience

PR Types

New features

Description

Expands the Apple MPS (MetalPerformanceShadersGraph) backend from 51 → 79
forward kernels
by porting the missing "siblings" of the op categories already
implemented — the ops that map cleanly onto MPSGraph primitives. All kernels are
forward-only and float32-only, one self-contained .mm file each, following the
existing MPS kernel template.

What's added (28 kernels)

Category Ops
Unary math asinh acosh atanh expm1 log1p trunc
Activations (no attr) tanh_shrink softsign logsigmoid hardswish relu6 swish
Activations (1 attr) elu celu mish softshrink hard_shrink
Activations (2 attr) hardtanh hardsigmoid thresholded_relu softplus
Reduction activation log_softmax
Binary elementwise atan2 floor_divide remainder fmax fmin heaviside

Implementation notes

  • Each kernel builds an MPSGraph, runs it via runWithMTLCommandQueue:…, and
    registers float only, matching the existing MPS kernels. New files are picked
    up automatically by the mps/*.mm glob in paddle/phi/kernels/CMakeLists.txt
    (no build-file changes).
  • Semantics are matched to the CPU reference functors (funcs/activation_functor.h,
    funcs/elementwise_functor.h) so MPS output equals CPU output:
    • floor_divide = floor(x/y), remainder = x - floor(x/y)*y (sign of divisor)
    • logsigmoid / softplus / mish use the numerically-stable functor forms
    • hardswish bakes threshold/scale/offset = 6/6/3; relu6 clamps to 6; swish β=1
    • elu = $α(e^x−1)$ for x<0; celu = $α(e^{x/α}−1)$ for x<0
  • log_softmax relies on MPSGraph reductions keeping the reduced axis (size 1),
    so the max/sum subtractions broadcast without an explicit reshape.
  • New MPSGraph APIs used: selectWithPredicateTensor:…,
    atan2WithPrimaryTensor:secondaryTensor:, and the inverse-hyperbolics
    asinh/acosh/atanhWithTensor:.

Testing

Adds test/test_mps_math_activation_kernels.py, which exercises every new op on
the MPS backend against the CPU backend (same public paddle API) plus a set of
numpy known-value checks. Guarded by paddle.is_compiled_with_mps() /
paddle.mps.is_available(), so it skips cleanly where MPS is unavailable.

Build + run on macOS:

make sure you have cmake and gcc installed on your Mac. For more information, check this guidance:
macOS 下使用 make 从源码编译

rm -rf build && mkdir build && cd build
cmake .. -DPY_VERSION=3.10 -DWITH_GPU=OFF -DWITH_MPS=ON
make -j$(sysctl -n hw.ncpu)
python test/test_mps_math_activation_kernels.py

Scope

Forward, float32 only. Reductions (prod/all/any/argmax/argmin/
logsumexp), selection/util (where/clip/cast/isnan/bitwise), and shape
ops remain as a follow-up.

🤖 Generated with Claude Code

是否引起精度变化

ShigureNyako and others added 30 commits June 11, 2026 11:43
---------

Co-authored-by: Codex <noreply@openai.com>
* Add uint32 max enforce helper

* fix bugs

* fix bugs
* acc paddle.sqrt

* fix

* add FLAGS_use_accuracy_compatible_kernel

* fix

* fix

---------

Co-authored-by: fanbohao <fanbohao@stu.pku.edu.cn>
…()` and `paddle.autograd.backward()` (#79267)

* support create_graph in paddle.Tensor.backward() and paddle.autograd.backward()

* add param alias and combine tests
…and_masks (#79275)

* [API Compatibility] Align torch.nn.attention.flex_attention.or_masks/and_masks

* [API Compatibility] Fix

* [API Compatibility] Fix
…e -part (#79277)

* align smooth_l1_loss by adding it to compat module

* remove additional comments

* add coverage
Co-authored-by: Codex <noreply@openai.com>
…9246)

* fix int64 = int32

agent check step 2

* fix bugs

* fix bugs

* fix bugs

* fix bugs

* fix bugs

* Address review comments

* fix bugs

* fix bugs

* fix bugs

* fix codestyle

* revert for_range.h

* fix premature int narrowing in gpu kernels

* Update generate_proposals_kernel.cu

* Update cross_entropy.cu
* add key_validation

* fix

* fix

* fix test

* fix

* fix
* [CINN] Update LLVM dependency to 13.0.1

* [CINN] Fix LLVM zlib target dependency

* [CI] Install libtinfo5 in Ubuntu 22 image

* [CINN] Drop removed LLVM LambdaResolver include

* [CINN] Adapt LLVM backend to LLVM 13 APIs

* [CINN] Fix remaining LLVM 13 API compile errors

* [CINN] Fix LLVM 13 target initialization linkage

* [CINN] Avoid LLVM RTTI references with LLVM 13

* [CINN] Address LLVM 13 follow-up fixes

* [LLVM] Update LLVM download URLs and initialization for native target support

* [CI] Remove obsolete LLVM tinfo dependencies

* [CINN] Remove obsolete LLVM dependency fixes

* [CINN] Restore LLVM object cache

* [CINN] Restore LLVM optimizer target machine API
* return tensor for paddle.nn.init.normal_

* add paddle.nn.init.sparse_

* add error raise test

* add EN doc

* revert to old implementation to test ci

* add test skip

* update return type

* add return value for normal_

* return tensor for all paddle.nn.init functions
…finite.check and torch.distributions.categorical.Categorical (#79279)

* [API Compatibility] Align torch.distributions.constraints.positive_definite.check and torch.distributions.categorical.Categorical

* [API Compatibility] Fix

* fix

* [API Compatibility] Fix

* fix

* fix

* fix

* fix

* fix

* fix
* [API Compatibility] Align torch BatchNorm

* [API Compatibility] Fix

* [API Compatibility] Fix

* [API Compatibility] Fix

* [API Compatibility] Fix

* [API Compatibility] Fix

* fix

* fix

* fix

* fix

* fix
* feat(ap): add DCU support for AP backend

* ap get device type api

* ap unified dcu/gpu template

* ap unified dcu/gpu template

* tmp commit

* general ap template

* fix ap cutlass template bug

* autotune for dcu

* fix ap NVGPU autotune bug

* [AP]: merge matmul_variadic_tpl

* [AP]: move cutlass patch file

* update gitignore

* [AP]: cutlass_matmul ap namespace

* [AP]: add ck patch

* add license

* Implement CompileCommand for dcu.

* [AP] Adapt DCU backend from CK to hytlass and fix DTK 25.04.2 build.

* Optimize unittest.

* Remove ck_patch.

* Add hytlass patch.

* Rename namespace from cutlass to cutlass_patch.

* Reorganize codes.

* Add hip backend.

* Set cutlass_dir through environ.

* Fix stream, setting for ap_path and ap_flags, and conversion of int.

* Fix setting of FLAGS_prim_all.

---------

Co-authored-by: _Aln0 <jiangyonglin@baidu.com>
Co-authored-by: _Aln0 <2629184561@qq.com>
* Fix depthwise conv int32 overflow

fix bugs

fix bugs

fix bugs

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

fix bugs

fix bugs

* revert enfotce.h

* fix bugs

* fix bugs

* fix bugs

* fix bugs

* refine int64 helper variable names

* fix bugs

* fix bugs

* fix bugs

* fix bugs

* fix bugs
* support gcc15

* feat update coverage ci and add install gcc15

* ci rm git build

* fix Coverage git

* update docker file

* clean docker py3.9

* Fix coverage Docker GCC 15 PPA setup

* fix fa gcc 15 build error

* clean include

* Revert coverage CI dockerfile changes
---------

Co-authored-by: Nyako Shigure <shigure_nyako@outlook.com>
Co-authored-by: Nyakku Shigure <sigure.qaq@gmail.com>
ForFishes and others added 18 commits June 25, 2026 14:29
* Enable C++20 toolchain baseline

* Fix C++20 CI baseline regressions

* Lower GCC C++20 baseline to 11

* Force CUDA C++20 CI exposure

* Patch Crypto++ for C++20 constructor syntax

* Fix AP type constructor for C++20

* Fix C++20 accumulate lambda in cum maxmin

* Fix PIR attribute checks for C++20

* Fix Windows C++20 compile errors

* Use std numbers for angle constant

* Fix CUDA angle constant for C++20

* Fix distributed PIR null checks for C++20

* Fix DDim visitor return type for nvcc

* Update Windows C++20 CI toolchain

* Keep Windows builds on C++17 baseline

* Document Windows C++20 follow-up

* Guard std numbers usage for Windows C++17

* Link static PIR deps for CUDA tests

* Fallback to C++17 for older GCC

* Warn on GCC C++17 fallback

* Link nv tests with static registry deps

* Fix C++17 numbers fallback and nv test deps

* Fix C++20 coverage build errors

* Link nv tests with registry deps

* Link phi core after nv test registry deps

* Link var helper with type info

* Group var helper with nv test type info deps

* Keep var helper type info link order

* Narrow AP attr TypeId probing

* Group paddle gtest static deps

* Avoid null check ambiguity for PIR type

* Fix C++20 accumulate callbacks in inference tests

* Remove global paddle gtest static deps

* Preserve AP attribute matching with TypeId deps

* Localize static deps for GPU tests

* Fix threadpool patch application
* Align cuBLAS workspace size for SM10 GPUs

* fix bugs

* Align cuBLASLt workspace size with cuBLAS

* Update gpu_context.cc
…ten/kl_div/GRU/Softmax/L1Loss/avg_pool2d/LogSoftmax/hann_window/det/pinverse/addcdiv_/real/inv/ParameterList/scatter_reduce_/xavier_uniform/ge_/sign_/pinv/take/resize_as_/huber_loss/float/set_default_tensor_type/fmod/pack_padded_sequence/pad_packed_sequence/ctc_loss/nan_to_num/fmod_/fill_diagonal_/assert_allclose/absolute/randint_like Edit By AI Agent (#79312)

* [API Compatibility] nll_loss/bernoulli_/weight_norm/resize_/imag/Flatten/kl_div/GRU/Softmax/L1Loss/avg_pool2d/LogSoftmax/hann_window/det/pinverse/addcdiv_/real/inv/ParameterList/scatter_reduce_/xavier_uniform/ge_/sign_/pinv/take/resize_as_/huber_loss/float/set_default_tensor_type/fmod/pack_padded_sequence/pad_packed_sequence/ctc_loss/nan_to_num/fmod_/fill_diagonal_/assert_allclose/absolute/randint_like Edit By AI Agent

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

* remove example code temporarily

* fix some bug

* add coverage

* fix example code

---------

Co-authored-by: zhouwei25 <zhouwei25@baidu.com>
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* support sublayer and layer share weight with SharedLayerDesc

* add ut

---------

Co-authored-by: AlAuAu <458134681@qq.com>
…as for `DataLoader` (#79385)

* support optimizer for LambdaDecay

* add alias for paddle.io.DataLoader

* set default for shuffle
* [API Compatibility] API signature mechanism optimization

* fix

* Refactor _paddle_docs decorator usage
Co-authored-by: adam_xiaoyao <adam_x@126.com>
* [PIR] Fix compatible dims check

* mark VerifyCompatibleDims as IR_API to fix windows build

* [Coverage] keep paddle_test gcda for changed source

---------

Co-authored-by: SigureMo <sigure.qaq@gmail.com>
* Add tensor metadata to api tracer output

* fix: hide api tracer tensor place on xpu

* fix

* [CI] add approval check for int variables in kernel code

* empty

* Fix API tracer XPU metadata handling

* Stabilize API tracer metadata tests

* Use stable non-contiguous tensor in API tracer test

* Allow backend-specific tensor contiguity in tracer test

---------

Co-authored-by: zrr1999 <zrr1999@users.noreply.github.com>
@Manfredss Manfredss changed the title Dev 2026 07 01 [MPS Dev] 2026.07.01 Jul 1, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.