2

I have the following line of code in a sycl application under a parallel_for:

  queue.submit(
        [&map, &output](cl::sycl::handler& cgh)
        {
            auto d_image = output.get_access<access::mode::read_write>(cgh);
            auto d_map_probabilities =
                map.probabilities.get_access<access::mode::read>(cgh);
            cgh.parallel_for<class CreateImage>(
                map.probabilities.get_range(),
                [=](id<3> work_position)
                {
                    // ... Code!

                    float3 relative{};
                    float elevation =
                        atan2(relative.z(), length(relative.xy())) * 30.0f * M_2_PI_F;
                }
          }
    );

For which i get the error:

error: no matching function for call to 'length'
                        atan2(relative.z(), length(relative.xy())) * 30.0f * M_2_PI_F;
                                            ^~~~~~
/opt/sycl/bin/../include/sycl/CL/sycl/builtins.hpp:1032:7: note: candidate template ignored: requirement 'detail::is_contained<sycl::detail::SwizzleOp<sycl::vec<float, 3>, sycl::detail::GetOp<float>, sycl::detail::GetOp<float>, sycl::detail::GetOp, 0, 1>, sycl::detail::type_list<sycl::detail::type_list<float>, sycl::detail::type_list<sycl::vec<float, 1>, sycl::vec<float, 2>, sycl::vec<float, 3>, sycl::vec<float, 4>>>>::value' was not satisfied [with T = sycl::detail::SwizzleOp<sycl::vec<float, 3>, sycl::detail::GetOp<float>, sycl::detail::GetOp<float>, sycl::detail::GetOp, 0, 1>]
float length(T p) __NOEXC {
      ^
/opt/sycl/bin/../include/sycl/CL/sycl/builtins.hpp:1039:8: note: candidate template ignored: requirement 'detail::is_contained<sycl::detail::SwizzleOp<sycl::vec<float, 3>, sycl::detail::GetOp<float>, sycl::detail::GetOp<float>, sycl::detail::GetOp, 0, 1>, sycl::detail::type_list<sycl::detail::type_list<double>, sycl::detail::type_list<sycl::vec<double, 1>, sycl::vec<double, 2>, sycl::vec<double, 3>, sycl::vec<double, 4>>>>::value' was not satisfied [with T = sycl::detail::SwizzleOp<sycl::vec<float, 3>, sycl::detail::GetOp<float>, sycl::detail::GetOp<float>, sycl::detail::GetOp, 0, 1>]
double length(T p) __NOEXC {
       ^
/opt/sycl/bin/../include/sycl/CL/sycl/builtins.hpp:1046:6: note: candidate template ignored: requirement 'detail::is_contained<sycl::detail::SwizzleOp<sycl::vec<float, 3>, sycl::detail::GetOp<float>, sycl::detail::GetOp<float>, sycl::detail::GetOp, 0, 1>, sycl::detail::type_list<sycl::detail::type_list<sycl::detail::half_impl::half>, sycl::detail::type_list<sycl::vec<sycl::detail::half_impl::half, 1>, sycl::vec<sycl::detail::half_impl::half, 2>, sycl::vec<sycl::detail::half_impl::half, 3>, sycl::vec<sycl::detail::half_impl::half, 4>>>>::value' was not satisfied [with T = sycl::detail::SwizzleOp<sycl::vec<float, 3>, sycl::detail::GetOp<float>, sycl::detail::GetOp<float>, sycl::detail::GetOp, 0, 1>]
half length(T p) __NOEXC {
     ^

Basically it comes down to this. The operation relative.xy() produces the type:

sycl::detail::SwizzleOp<sycl::vec<float, 3>, sycl::detail::GetOp<float>, sycl::detail::GetOp<float>, sycl::detail::GetOp, 0, 1>

which is not implicitly convertible to float2 which is required by the length function. It should be though, this was fine to do in opencl. This can be solved by the workaround:

float2 rel_part = relative.xy();
float elevation = atan2(relative.z(), length(rel_part)) * 30.0f * M_2_PI_F;

Or

float elevation = atan2(relative.z(), length(float2{relative.xy()})) * 30.0f * M_2_PI_F;

But this defeats the purpose of a nice swizzle definition. Is there a way around this?

Fantastic Mr Fox
  • 32,495
  • 27
  • 95
  • 175
  • I think this is a DPC++ CTAD bug. This works: `length(relative.xy())`. I can raise this on the DPC++ issue tracker with relevant implementation details, if you like? – Joe Todd Jul 06 '22 at 10:00
  • https://github.com/intel/llvm/issues/6402 – Joe Todd Jul 06 '22 at 10:12
  • According to section [4.14.2.1](https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#_vec_interface) from SYCL 2020 spec, `XYZW_SWIZZLE()` returns a value of `__swizzled_vec__` type. And according to 4.17.1 and 4.17.8 sections, `length()` doesn't have overload for this type. So, according to the specification, it necessary to materialize swizzle to `vec_t` first. It possible to use already mentioned variants or `length(relative.xy().as())` or even `length(+relative.xy())`. But I fully agree, that it is more convenient to use `length(relative.xy())` in the code. – Egor Jul 06 '22 at 18:26
  • @JoeTodd Nice. That was quick. – Fantastic Mr Fox Jul 06 '22 at 22:43

0 Answers0