[SYCL] Add test to cover std::complex<float/double> mul/div by jinge90 · Pull Request #21622 · intel/llvm

jinge90 · 2026-03-25T08:39:07Z

libdevice provides __mulsc3, __divsc3, __muldc3, __divdc3 to support complex number multiplication and division in 'no-fast-math' mode on Linux platform only. These 4 builtins are not invoked by user code but inserted by compiler when handling complex mul/div expression. This PR adds e2e test for these builtins explicitly, the testing logic is ported from compiler-rt test suite for normal CPU platform:
https://github.com/llvm/llvm-project/blob/main/compiler-rt/test/builtins/Unit/divsc3_test.c
https://github.com/llvm/llvm-project/blob/main/compiler-rt/test/builtins/Unit/divdc3_test.c
https://github.com/llvm/llvm-project/blob/main/compiler-rt/test/builtins/Unit/mulsc3_test.c
https://github.com/llvm/llvm-project/blob/main/compiler-rt/test/builtins/Unit/muldc3_test.c

Signed-off-by: jinge90 <ge.jin@intel.com>

jinge90 · 2026-03-26T13:34:04Z

Hi, @intel/llvm-reviewers-runtime and @sergey-semenov
Could you take a look at this PR?
Thanks very much.

Copilot

Pull request overview

Adds new SYCL end-to-end coverage for device-side std::complex<float/double> multiplication and division, intended to exercise the libdevice/compiler-emitted __mul{sd}c3 / __div{sd}c3 builtins under “no-fast-math” behavior.

Changes:

Introduce shared complex mul/div test utilities and result validation helpers.
Add a comprehensive complex input matrix (incl. NaN/Inf/±0) and run mul/div tests for float and double.
Wire the new tests into the existing std_complex_math_* e2e test executables (non-Windows).

Reviewed changes

Copilot reviewed 3 out of 3 changed files in this pull request and generated 2 comments.

File	Description
sycl/test-e2e/DeviceLib/std_complex_math_test.cpp	Adds `std::complex<float>` mul/div e2e coverage using shared helpers and a large edge-case input set.
sycl/test-e2e/DeviceLib/std_complex_math_fp64_test.cpp	Adds `std::complex<double>` mul/div e2e coverage using the shared helpers and input set.
sycl/test-e2e/DeviceLib/complex_utils.hpp	New helper header providing classification/compare logic and SYCL kernels for mul/div matrix evaluation.

Copilot · 2026-03-26T17:12:58Z

sycl/test-e2e/DeviceLib/std_complex_math_fp64_test.cpp


+#include "complex_utils.hpp"
 #include "math_utils.hpp"


This test is intended to exercise the complex mul/div builtins emitted in "no-fast-math" mode, but the RUN line here doesn’t pass -fno-fast-math (unlike std_complex_math_test.cpp and other DeviceLib math tests). Please define/use %{mathflags} and compile this test with -fno-fast-math to ensure the intended codegen is covered.

Copilot · 2026-03-26T17:12:58Z

sycl/test-e2e/DeviceLib/complex_utils.hpp

+#pragma once
+#include <complex>
+#include <sycl/detail/core.hpp>
+enum { zero, non_zero, inf, NaN, non_zero_nan };
+template <typename T> int complex_classify(std::complex<T> x) {
+  if (x == std::complex<T>(0, 0))
+    return zero;
+  if (std::isinf(x.real()) || std::isinf(x.imag()))
+    return inf;


complex_utils.hpp uses std::array, std::isnan, and std::isinf but only includes <complex>. Please add the missing standard headers (e.g., <array> and <cmath>) so this helper is self-contained and doesn’t rely on include order in the including test files.

Copilot

Pull request overview

Copilot reviewed 3 out of 3 changed files in this pull request and generated 7 comments.

Copilot · 2026-03-31T08:30:39Z

sycl/test-e2e/DeviceLib/std_complex_math_test.cpp

 // DEFINE: %{mathflags} = %if cl_options %{/clang:-fno-fast-math%} %else %{-fno-fast-math%}
 // RUN: %{build} %{mathflags} -o %t1.out
 // RUN: %{run} %t1.out


The PR description says the mul/div builtins being exercised are Linux-only, but this test only excludes Windows (#ifndef _WIN32). That means the new mul/div coverage will still run on other non-Windows platforms (e.g. Darwin) where those builtins may not be available. Consider adding an appropriate LIT guard (e.g. // REQUIRES: linux or // UNSUPPORTED: system-darwin) so the test only runs where the targeted device-lib support exists.

Copilot · 2026-03-31T08:30:40Z

sycl/test-e2e/DeviceLib/std_complex_math_fp64_test.cpp

 // REQUIRES: aspect-fp64
-// RUN: %{build} -o %t1.out
+// DEFINE: %{mathflags} = %if cl_options %{/clang:-fno-fast-math%} %else %{-fno-fast-math%}
+// RUN: %{build} %{mathflags} -o %t1.out
 // RUN: %{run} %t1.out


Same issue here: the new mul/div builtins coverage is described as Linux-only, but the test is only gated with #ifndef _WIN32. Please add a LIT platform guard (e.g. // REQUIRES: linux or // UNSUPPORTED: system-darwin) so it doesn’t run on non-Linux non-Windows hosts by accident.

sycl/test-e2e/DeviceLib/std_complex_math_test.cpp

Copilot · 2026-03-31T08:30:40Z

sycl/test-e2e/DeviceLib/std_complex_math_fp64_test.cpp

+complex<double> complex_input[] = {{1.e-6, 1.e-6},
+                                   {-1.e-6, 1.e-6},
+                                   {-1.e-6, -1.e-6},
+                                   {1.e-6, -1.e-6},


Same as in the float test: complex_input doesn’t need external linkage. Consider making it static (or move to an anonymous namespace) since it’s only consumed within this test TU.

Copilot · 2026-03-31T08:30:40Z

sycl/test-e2e/DeviceLib/complex_utils.hpp

+                            std::complex<T> *complex_input) {
+  constexpr size_t OutputSize = InputSize * InputSize;
+  sycl::range<1> numOfMulInput{InputSize};
+  sycl::range<1> numOfMulOutput{OutputSize};
+  std::array<std::complex<T>, OutputSize> complex_mul_result;
+  {
+    sycl::buffer<std::complex<T>, 1> buffer_complex_mul(complex_input,
+                                                        numOfMulInput);


device_complex_test_mul only reads from complex_input. Consider taking const std::complex<T>* (and creating the input buffer from a const pointer where possible) to make the API intent clear and prevent accidental mutation.

Suggested change

std::complex<T> *complex_input) {

constexpr size_t OutputSize = InputSize * InputSize;

sycl::range<1> numOfMulInput{InputSize};

sycl::range<1> numOfMulOutput{OutputSize};

std::array<std::complex<T>, OutputSize> complex_mul_result;

{

sycl::buffer<std::complex<T>, 1> buffer_complex_mul(complex_input,

numOfMulInput);

const std::complex<T> *complex_input) {

constexpr size_t OutputSize = InputSize * InputSize;

sycl::range<1> numOfMulInput{InputSize};

sycl::range<1> numOfMulOutput{OutputSize};

std::array<std::complex<T>, OutputSize> complex_mul_result;

{

sycl::buffer<const std::complex<T>, 1> buffer_complex_mul(complex_input,

numOfMulInput);

Copilot · 2026-03-31T08:30:41Z

sycl/test-e2e/DeviceLib/complex_utils.hpp

+                            std::complex<T> *complex_input) {
+  constexpr size_t OutputSize = InputSize * InputSize;
+  sycl::range<1> numOfDivInput{InputSize};
+  sycl::range<1> numOfDivOutput{OutputSize};
+  std::array<std::complex<T>, OutputSize> complex_div_result;
+  {
+    sycl::buffer<std::complex<T>, 1> buffer_complex_div(complex_input,
+                                                        numOfDivInput);


device_complex_test_div only reads from complex_input. Consider taking const std::complex<T>* (and using a read-only buffer/accessor) to better express const-correctness.

Suggested change

std::complex<T> *complex_input) {

constexpr size_t OutputSize = InputSize * InputSize;

sycl::range<1> numOfDivInput{InputSize};

sycl::range<1> numOfDivOutput{OutputSize};

std::array<std::complex<T>, OutputSize> complex_div_result;

{

sycl::buffer<std::complex<T>, 1> buffer_complex_div(complex_input,

numOfDivInput);

const std::complex<T> *complex_input) {

constexpr size_t OutputSize = InputSize * InputSize;

sycl::range<1> numOfDivInput{InputSize};

sycl::range<1> numOfDivOutput{OutputSize};

std::array<std::complex<T>, OutputSize> complex_div_result;

{

sycl::buffer<const std::complex<T>, 1> buffer_complex_div(complex_input,

numOfDivInput);

Copilot · 2026-03-31T08:30:41Z

sycl/test-e2e/DeviceLib/std_complex_math_test.cpp

+complex<float> complex_input[] = {{1.e-6, 1.e-6},
+                                  {-1.e-6, 1.e-6},
+                                  {-1.e-6, -1.e-6},
+                                  {1.e-6, -1.e-6},
+                                  {1.e+6, 1.e-6},
+                                  {-1.e+6, 1.e-6},
+                                  {-1.e+6, -1.e-6},
+                                  {1.e+6, -1.e-6},
+
+                                  {1.e-6, 1.e+6},
+                                  {-1.e-6, 1.e+6},
+                                  {-1.e-6, -1.e+6},
+                                  {1.e-6, -1.e+6},
+
+                                  {1.e+6, 1.e+6},
+                                  {-1.e+6, 1.e+6},
+                                  {-1.e+6, -1.e+6},
+                                  {1.e+6, -1.e+6},
+
+                                  {NAN, NAN},
+                                  {-INFINITY, NAN},
+                                  {-2, NAN},
+                                  {-1, NAN},
+                                  {-0.5, NAN},
+                                  {-0., NAN},


In this complex<float> initializer list, literals like 1.e-6 are double. With braced initialization ({...}) this triggers C++ narrowing rules, and values like 1.e-6 are not exactly representable as float, so this is likely to be a hard compile error (“narrowing conversion from 'double' to 'float'”). Use f-suffixed literals (e.g. 1.e-6f) and/or switch these elements to complex<float>(...) (parentheses) to avoid narrowing errors.

Suggested change

complex<float> complex_input[] = {{1.e-6, 1.e-6},

{-1.e-6, 1.e-6},

{-1.e-6, -1.e-6},

{1.e-6, -1.e-6},

{1.e+6, 1.e-6},

{-1.e+6, 1.e-6},

{-1.e+6, -1.e-6},

{1.e+6, -1.e-6},

{1.e-6, 1.e+6},

{-1.e-6, 1.e+6},

{-1.e-6, -1.e+6},

{1.e-6, -1.e+6},

{1.e+6, 1.e+6},

{-1.e+6, 1.e+6},

{-1.e+6, -1.e+6},

{1.e+6, -1.e+6},

{NAN, NAN},

{-INFINITY, NAN},

{-2, NAN},

{-1, NAN},

{-0.5, NAN},

{-0., NAN},

complex<float> complex_input[] = {{1.e-6f, 1.e-6f},

{-1.e-6f, 1.e-6f},

{-1.e-6f, -1.e-6f},

{1.e-6f, -1.e-6f},

{1.e+6f, 1.e-6f},

{-1.e+6f, 1.e-6f},

{-1.e+6f, -1.e-6f},

{1.e+6f, -1.e-6f},

{1.e-6f, 1.e+6f},

{-1.e-6f, 1.e+6f},

{-1.e-6f, -1.e+6f},

{1.e-6f, -1.e+6f},

{1.e+6f, 1.e+6f},

{-1.e+6f, 1.e+6f},

{-1.e+6f, -1.e+6f},

{1.e+6f, -1.e+6f},

{NAN, NAN},

{-INFINITY, NAN},

{-2.f, NAN},

{-1.f, NAN},

{-0.5f, NAN},

{-0.f, NAN},

…_divsc3

Signed-off-by: jinge90 <ge.jin@intel.com>

Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>

Copilot

Pull request overview

Copilot reviewed 3 out of 3 changed files in this pull request and generated 4 comments.

sycl/test-e2e/DeviceLib/std_complex_math_test.cpp

Copilot · 2026-04-01T06:26:55Z

sycl/test-e2e/DeviceLib/complex_utils.hpp

+    deviceQueue.submit([&](sycl::handler &cgh) {
+      auto complex_mul_access =
+          buffer_complex_mul.template get_access<sycl::access::mode::read>(cgh);
+      auto complex_mul_res_access =
+          buffer_complex_mul_res.template get_access<sycl::access::mode::write>(
+              cgh);
+      cgh.single_task<class DeviceComplexMulTest>([=]() {
+        size_t i, j;
+        for (i = 0; i < InputSize; ++i) {


device_complex_test_mul is a function template, but the kernel is always named DeviceComplexMulTest. If this helper ever gets instantiated more than once in the same translation unit (e.g., for both float and double, or different InputSize), the identical kernel name type can cause a SYCL kernel name collision/ODR violation. Consider making the kernel name depend on the template parameters (e.g., a templated kernel-name type) or using an unnamed kernel form if supported by the project/toolchain.

Copilot · 2026-04-01T06:26:55Z

sycl/test-e2e/DeviceLib/complex_utils.hpp

+  size_t i, j;
+  for (i = 0; i < InputSize; ++i)
+    for (j = 0; j < InputSize; ++j) {
+      if (complex_compare_mul(complex_input[i], complex_input[j],
+                              complex_mul_result[i * InputSize + j])) {
+        return 1;
+      }
+    }
+  return 0;


Both device_complex_test_mul/device_complex_test_div return 1 on the first mismatch and don’t report which (i,j) input pair failed or what the computed/expected values were. That makes failures hard to diagnose in CI. Consider accumulating a failure count (like the other tests in this directory) and/or printing the failing indices and values before returning.

Copilot · 2026-04-01T06:26:55Z

sycl/test-e2e/DeviceLib/complex_utils.hpp

+    deviceQueue.submit([&](sycl::handler &cgh) {
+      auto complex_div_access =
+          buffer_complex_div.template get_access<sycl::access::mode::read>(cgh);
+      auto complex_div_res_access =
+          buffer_complex_div_res.template get_access<sycl::access::mode::write>(
+              cgh);
+      cgh.single_task<class DeviceComplexDivTest>([=]() {
+        size_t i, j;
+        for (i = 0; i < InputSize; ++i) {


device_complex_test_div is a function template, but the kernel is always named DeviceComplexDivTest. If this helper is instantiated more than once in the same translation unit (different T/InputSize), the repeated kernel name type can cause a SYCL kernel name collision/ODR violation. Consider parameterizing the kernel-name type on the template arguments (or using an unnamed kernel form).

Signed-off-by: jinge90 <ge.jin@intel.com>

Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>

jinge90 · 2026-04-01T07:36:17Z

Hi, @intel/llvm-reviewers-runtime and @sergey-semenov
This PR ported compiler-rt test for __mulsc3/dc3 and __divsc3/dc3 builtin to sycl for Linux platform, could you help review?
Thanks very much.

jinge90 · 2026-04-02T01:33:03Z

Hi, @uditagarwal97
Could you help review this PR, Sergey is on vacation.
Thanks very much.

jinge90 added 20 commits November 24, 2025 16:52

remove dummy bin

2d50750

Signed-off-by: jinge90 <ge.jin@intel.com>

Merge remote-tracking branch 'upstream/sycl' into sycl

9c76946

Merge remote-tracking branch 'upstream/sycl' into sycl

fe628a9

Merge remote-tracking branch 'upstream/sycl' into sycl

7a7e62b

merge remote-tracking branch 'upstream/sycl' into sycl

4de8fa5

Merge remote-tracking branch 'upstream/sycl' into sycl

e52bb98

Merge remote-tracking branch 'upstream/sycl' into sycl

4b63322

Merge remote-tracking branch 'upstream/sycl' into sycl

544d912

Merge remote-tracking branch 'upstream/sycl' into sycl

2ad6143

Merge remote-tracking branch 'upstream/sycl' into sycl

cf391d9

Merge remote-tracking branch 'upstream/sycl' into sycl

e0f2a15

Merge remote-tracking branch 'upstream/sycl' into sycl

aa7a97c

Merge remote-tracking branch 'upstream/sycl' into sycl

9244680

Merge remote-tracking branch 'upstream/sycl' into sycl

b4b5793

Merge remote-tracking branch 'upstream/sycl' into sycl

50f8f75

Merge remote-tracking branch 'upstream/sycl' into sycl

ebc7166

Merge remote-tracking branch 'upstream/sycl' into sycl

159748c

Merge remote-tracking branch 'upstream/sycl' into sycl

17362e0

Merge remote-tracking branch 'upstream/sycl' into sycl

9e38387

Merge remote-tracking branch 'origin/sycl' into sycl

dd1e30b

jinge90 requested a review from a team as a code owner March 25, 2026 08:39

jinge90 requested a review from sergey-semenov March 25, 2026 08:39

jinge90 marked this pull request as draft March 25, 2026 08:39

jinge90 added 2 commits March 25, 2026 17:08

Merge remote-tracking branch 'upstream/sycl' into sycl

01fa557

[SYCL] Add E2E test to cover std::complex<T> mul/div

08ff3e3

Signed-off-by: jinge90 <ge.jin@intel.com>

jinge90 marked this pull request as ready for review March 26, 2026 02:50

jinge90 added 2 commits March 26, 2026 11:15

make the mulsc3/divsc3 test to run on Linux

bc9a08e

Signed-off-by: jinge90 <ge.jin@intel.com>

fix typo

2102954

Signed-off-by: jinge90 <ge.jin@intel.com>

uditagarwal97 requested a review from Copilot March 26, 2026 17:06

Copilot started reviewing on behalf of uditagarwal97 March 26, 2026 17:08 View session

Copilot AI reviewed Mar 26, 2026

View reviewed changes

jinge90 requested a review from Copilot March 31, 2026 08:22

Copilot started reviewing on behalf of jinge90 March 31, 2026 08:24 View session

Copilot AI reviewed Mar 31, 2026

View reviewed changes

jinge90 and others added 4 commits March 31, 2026 16:46

Merge remote-tracking branch 'upstream/sycl' into add_test_for_mulsc3…

a7d4c12

…_divsc3

fix header

5f7adec

Signed-off-by: jinge90 <ge.jin@intel.com>

include cmath header

1a15fd8

Signed-off-by: jinge90 <ge.jin@intel.com>

Update sycl/test-e2e/DeviceLib/std_complex_math_test.cpp

fc65987

Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>

jinge90 requested a review from Copilot April 1, 2026 06:19

Copilot started reviewing on behalf of jinge90 April 1, 2026 06:21 View session

Copilot AI reviewed Apr 1, 2026

View reviewed changes

jinge90 and others added 2 commits April 1, 2026 14:44

address copilot comments

a30bd8d

Signed-off-by: jinge90 <ge.jin@intel.com>

Update sycl/test-e2e/DeviceLib/std_complex_math_test.cpp

b5938c9

Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>

jinge90 requested a review from uditagarwal97 April 2, 2026 01:31

Conversation

jinge90 commented Mar 25, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

jinge90 commented Mar 26, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Copilot AI Mar 26, 2026

Choose a reason for hiding this comment

Uh oh!

jinge90 Mar 31, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 26, 2026

Choose a reason for hiding this comment

Uh oh!

jinge90 Mar 31, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Copilot AI Mar 31, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 31, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Copilot AI Mar 31, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 31, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 31, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 31, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Uh oh!

Copilot AI Apr 1, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Apr 1, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Apr 1, 2026

Choose a reason for hiding this comment

Uh oh!

jinge90 commented Apr 1, 2026

Uh oh!

jinge90 commented Apr 2, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

jinge90 commented Mar 25, 2026 •

edited

Loading