Std::complex multiplication

bdallen · June 8, 2020, 3:50pm

Is there a convenient / portable way to use std::complex with sycl 1.2.1, including at least the four basic arithmetic operations?

I have noticed that addition and subtraction of std::complex works fine in both ComputeCpp and Intel DPCPP SYCL implementations. Multiplication and division fail at runtime when running on an Intel GPU or CPU (using latest intel-compute-runtime from git). ComputeCpp warns that __muldc3 is undefined during compilation and terminates with ‘cl::sycl::compile_program_error’. dpcpp has no compiler warning, also fails at runtime with compile_program_errors with error: undefined reference to __muldc3()’`. Interestingly hipsycl on AMD gfx803 works fine for this case. With DPCPP, if I link the libsycl-cmath-fp64.o and libsycl-complex-fp64.o, then it works on GPU. What is the recommendation for doing this in computecpp?

Here is my test code:

github.com

bd4/sycl-test/blob/master/complex.cxx

#include <CL/sycl.hpp>
#include <complex>
#include <iostream>

#define OP *

using namespace cl::sycl;

int main(int argc, char **argv) {
    constexpr int N = 16;
    using dcomplex = std::complex<double>;
    dcomplex h_a[N];
    dcomplex h_b[N];
    dcomplex h_c[N];
    int i;

    for (i = 0; i < N; i++) {
        h_a[i] = dcomplex(i, -i);
        h_b[i] = dcomplex(i,  i);
    }

This file has been truncated. show original

I noticed that the ComputeCpp mandelbrot example uses separate im/re variables and it’s own complex multiply algorithm in the kernel. Fine for simple kernels, but for porting libraries already using thrust::complex / std::complex to SYCL, it would be useful to have a portable way to do this.

duncan · June 8, 2020, 5:39pm

Hi @bdallen,

Great question! I knew I had some code lying around from a couple of years ago when I was first messing around with this, so I looked it up - sadly I’ve not left any particularly enlightening comments for myself. I have a sample which is similar to yours - using std::complex<T> on-device with for some simple operations.

I’ve since had a quick search and I’m starting to remember what’s going on. The compiler replaces part of the complex multiply and divide operations with a function call to a support library that handles cases like NaN and infinite inputs. I’m not sure that it’s necessarily a great idea to do this in general, but on the assumption that your code avoids NaN and infinite numbers, you can use the compiler flag -ffast-math to avoid emitting this support code and I think at this stage it basically just works.

I think a more specificaion-based answer would suggest that it won’t work, but personally I think it’s interesting to look at code that we might be able to run, because the technology allows us to do so.

I hope this helps,
Duncan.

bdallen · June 9, 2020, 3:45pm

Thanks for the tip! Unfortunately I was not able to get it working with ComputeCpp-CE 2.0.0. I tried building with -ffast-math, -cl-fast-relaxed-math, in different combinations, with the same result - still warning that __muldc3 is undefined. Do you remember which version you were testing with and what device/cpu backend you were using?

duncan · June 9, 2020, 6:38pm

Hi @bdallen,

I tried my old code from when I was last messing with this, which appears to be working fine, and I also copied in your code (I hope you don’t mind!). I get the output:

[GPU  ] Intel(R) HD Graphics {Intel(R) Corporation}
0: (0,0) OP (0,0) = (0,0)
1: (1,-1) OP (1,1) = (2,0)
2: (2,-2) OP (2,2) = (8,0)
3: (3,-3) OP (3,3) = (18,0)
4: (4,-4) OP (4,4) = (32,0)
5: (5,-5) OP (5,5) = (50,0)
6: (6,-6) OP (6,6) = (72,0)
7: (7,-7) OP (7,7) = (98,0)
8: (8,-8) OP (8,8) = (128,0)
9: (9,-9) OP (9,9) = (162,0)
10: (10,-10) OP (10,10) = (200,0)
11: (11,-11) OP (11,11) = (242,0)
12: (12,-12) OP (12,12) = (288,0)
13: (13,-13) OP (13,13) = (338,0)
14: (14,-14) OP (14,14) = (392,0)
15: (15,-15) OP (15,15) = (450,0)

I assume this is basically the correct output? I’m using ComputeCpp 2.0.0 on an Intel GPU of some kind (though I should warn I have a very old driver installed, mostly down to laziness). I get the same results if I run on the Intel CPU driver.

How are you compiling these samples? I am using the CMake we distribute to compile these samples, and setting the flag -ffast-math by the CMake variable COMPUTECPP_USER_FLAGS, which is designed for this purpose.

I hope this helps!

bdallen · June 10, 2020, 12:12pm

Thanks, it does work if I build with cmake using the FindComputeCpp etc from the sdk, and the add_sycl_to_target macro/function with COMPUTECPP_USER_FLAGS=-ffast-math. I was using the -sycl-driver command to build which is not working:

compute++ -sycl-driver -std=c++14 -ffast-math -I/opt/ComputeCpp-CE/include -L/opt/ComputeCpp-CE/lib -lComputeCpp -o build-computecpp/complex complex.cxx

I was able to get the Makefile build to work as well, using a separate -sycl pass with -ffast-math instead of -sycl-driver. I will experiment with it more.

duncan · June 16, 2020, 10:23am

Ah, that kind of makes sense. It’s really hijacking a host flag and I imagine in -sycl-driver mode -ffast-math gets applied to the host code only instead of the device code. Does that mean that -cl-fast-relaxed-math doesn’t do anything even in -sycl-driver mode?

zjin · June 30, 2020, 4:41pm

This is what I see:

(working) compute++ = -O3 -sycl-driver -std=c++14 -ffast-math complex.cpp -lComputeCpp
(not working) compute++ = -O3 -sycl-driver -std=c++14 complex.cpp -lComputeCpp
(not working) compute++ = -sycl-driver -std=c++14 -ffast-math complex.cpp -lComputeCpp

Thank you for your verification

bdallen · July 10, 2020, 1:01pm

Thanks zjin, I get the same results with ComputeCpp 2.0.0. -sycl-driver works if both -O3 and -ffast-math are enabled.

bdallen · July 10, 2020, 1:02pm

I started a discussion about standardizing complex numbers in SYCL here: https://github.com/intel/llvm/issues/2061, in case anyone here is interested.