Skip to content

Standardize math functions in SYCL kernels: std::/::f → sycl::, __f → sycl::native::#3203

Draft
Copilot wants to merge 3 commits intomainfrom
copilot/update-sycl-math-functions
Draft

Standardize math functions in SYCL kernels: std::/::f → sycl::, __f → sycl::native::#3203
Copilot wants to merge 3 commits intomainfrom
copilot/update-sycl-math-functions

Conversation

Copy link
Copy Markdown
Contributor

Copilot AI commented Mar 27, 2026

SYCL kernels were inconsistently mixing sycl::, std::, ::, and sycl::native:: math functions. The rule: high-precision calls (expf, std::exp, ::expf) → sycl::exp; CUDA fast intrinsics (__expf) → sycl::native::exp.

Changes

  • src/comm/XPUMathCompat.h: Replace ::expf/::tanhf/::exp/::tanh with sycl::exp/sycl::tanh

  • 58 SYCL kernel files: Replace all std:: math functions with sycl:: equivalents in device code:

    • Transcendentals: exp, log, log1p, log2, log10, expm1, exp2, sqrt, pow
    • Trig: sin, cos, tan, sinh, cosh, tanh, asin, acos, atan, atan2
    • Misc: erf, erfc, lgamma, fabs, fmax, fmin, hypot, floor, ceil, round, trunc, fmod
  • Complex type handling: Added specializations that preserve std:: for complex number math where SYCL has no equivalent (e.g., std::log1p(complex), complex sqrt). Affects SqrtFunctor, ExpFunctor in UnaryKernels.cpp; LogspaceFunctor in RangeFactoriesKernel.cpp; Rsqrt<c10::complex<T>> in ForeachUnaryKernels.cpp.

  • sycl::native::: No existing __expf-style intrinsic calls found — no changes required for this category.

Kept as std:: (intentional):

  • Host-side code (computed before sycl_kernel_submit, e.g., AdaptiveAveragePooling, AttentionKernels scale factor)
  • std::modf / std::frexp — SYCL variants require multi_ptr, different API
  • compat_pow in SharedReduceOps.h — dispatched with complex accumulator types
// Before
return x_acc * std::tanh(std::log1p(std::exp(x_acc)));

// After
return x_acc * sycl::tanh(sycl::log1p(sycl::exp(x_acc)));
// XPUMathCompat.h — Before
__MATH_FUNCTIONS_DECL__ float exp(float x) { return ::expf(x); }

// After
__MATH_FUNCTIONS_DECL__ float exp(float x) { return sycl::exp(x); }

✨ Let Copilot coding agent set things up for you — coding agent works faster and does higher quality work when set up for your repo.

Copilot AI and others added 2 commits March 27, 2026 14:56
…sion, keep std:: for complex types and host code

Agent-Logs-Url: https://github.com/intel/torch-xpu-ops/sessions/1c6785b9-c285-4ddc-bfd2-ed4892d7387a

Co-authored-by: jianyizh <79236274+jianyizh@users.noreply.github.com>
…, improve readability

Agent-Logs-Url: https://github.com/intel/torch-xpu-ops/sessions/1c6785b9-c285-4ddc-bfd2-ed4892d7387a

Co-authored-by: jianyizh <79236274+jianyizh@users.noreply.github.com>
Copilot AI changed the title [WIP] Refactor math functions in SYCL kernels Standardize math functions in SYCL kernels: std::/::f → sycl::, __f → sycl::native:: Mar 27, 2026
Copilot AI requested a review from jianyizh March 27, 2026 15:01
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.

2 participants