Hi,
I have a SYCL kernel that uses several math functions (sin, cos, sqrt, etc.). When I compile for CPUs targeting SPIR it compiles without warnings. However, when I target PTX64 to run it on a NVIDIA GPU I get a bunch of warnings about generated intrinsics being illegal in SPIR unless I remove all the calls to the math functions. Despite the warnings, the kernel runs as expected, so can/should I ignore these warnings?
Thanks!
Jolly
Warnings
warning: [Computecpp:CC0035]: Intrinsic llvm.fabs.f64 has been generated in function _ZZZ17InvariantMassSYCLPK15PtEtaPhiEVectorS1_mENK3$_0clERN2cl4sycl7handlerEENKUlNS4_4itemILi1ELb1EEEE_clE__ns1S8_ which is illegal in SPIR and may result in a
compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fabs.f64 has been generated in function _ZN2cl4sycl4sinhIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fabs.f64 has been generated in function _ZN2cl4sycl6detail4sinhE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fabs.f64 has been generated in function _ZN2cl4sycl3sinIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fabs.f64 has been generated in function _ZN2cl4sycl6detail3sinE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fabs.f64 has been generated in function _ZN2cl4sycl3cosIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fabs.f64 has been generated in function _ZN2cl4sycl6detail3cosE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fma.f64 has been generated in function _ZZZ17InvariantMassSYCLPK15PtEtaPhiEVectorS1_mENK3$_0clERN2cl4sycl7handlerEENKUlNS4_4itemILi1ELb1EEEE_clE__ns1S8_ which is illegal in SPIR and may result in a
warning: [Computecpp:CC0035]: Intrinsic llvm.fma.f64 has been generated in function _ZN2cl4sycl4sinhIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fma.f64 has been generated in function _ZN2cl4sycl6detail4sinhE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fma.f64 has been generated in function _ZN2cl4sycl3sinIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fma.f64 has been generated in function _ZN2cl4sycl6detail3sinE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fma.f64 has been generated in function _ZN2cl4sycl3cosIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.fma.f64 has been generated in function _ZN2cl4sycl6detail3cosE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.trunc.f64 has been generated in function _ZZZ17InvariantMassSYCLPK15PtEtaPhiEVectorS1_mENK3$_0clERN2cl4sycl7handlerEENKUlNS4_4itemILi1ELb1EEEE_clE__ns1S8_ which is illegal in SPIR and may result in a
warning: [Computecpp:CC0035]: Intrinsic llvm.trunc.f64 has been generated in function _ZN2cl4sycl3sinIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.trunc.f64 has been generated in function _ZN2cl4sycl6detail3sinE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.trunc.f64 has been generated in function _ZN2cl4sycl3cosIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.trunc.f64 has been generated in function _ZN2cl4sycl6detail3cosE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.floor.f64 has been generated in function _ZZZ17InvariantMassSYCLPK15PtEtaPhiEVectorS1_mENK3$_0clERN2cl4sycl7handlerEENKUlNS4_4itemILi1ELb1EEEE_clE__ns1S8_ which is illegal in SPIR and may result in a
warning: [Computecpp:CC0035]: Intrinsic llvm.floor.f64 has been generated in function _ZN2cl4sycl3sinIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.floor.f64 has been generated in function _ZN2cl4sycl6detail3sinE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.floor.f64 has been generated in function _ZN2cl4sycl3cosIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.floor.f64 has been generated in function _ZN2cl4sycl6detail3cosE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.minnum.f64 has been generated in function _ZZZ17InvariantMassSYCLPK15PtEtaPhiEVectorS1_mENK3$_0clERN2cl4sycl7handlerEENKUlNS4_4itemILi1ELb1EEEE_clE__ns1S8_ which is illegal in SPIR and may result in a
warning: [Computecpp:CC0035]: Intrinsic llvm.minnum.f64 has been generated in function _ZN2cl4sycl3sinIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.minnum.f64 has been generated in function _ZN2cl4sycl6detail3sinE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.minnum.f64 has been generated in function _ZN2cl4sycl3cosIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.minnum.f64 has been generated in function _ZN2cl4sycl6detail3cosE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.copysign.f64 has been generated in function _ZZZ17InvariantMassSYCLPK15PtEtaPhiEVectorS1_mENK3$_0clERN2cl4sycl7handlerEENKUlNS4_4itemILi1ELb1EEEE_clE__ns1S8_ which is illegal in SPIR and may result in a
warning: [Computecpp:CC0035]: Intrinsic llvm.copysign.f64 has been generated in function _ZN2cl4sycl4sinhIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.copysign.f64 has been generated in function _ZN2cl4sycl6detail4sinhE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.sqrt.f64 has been generated in function _ZZZ17InvariantMassSYCLPK15PtEtaPhiEVectorS1_mENK3$_0clERN2cl4sycl7handlerEENKUlNS4_4itemILi1ELb1EEEE_clE__ns1S8_ which is illegal in SPIR and may result in a
warning: [Computecpp:CC0035]: Intrinsic llvm.sqrt.f64 has been generated in function _ZN2cl4sycl4sqrtIddLi0EEE__ns1dd which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
warning: [Computecpp:CC0035]: Intrinsic llvm.sqrt.f64 has been generated in function _ZN2cl4sycl6detail4sqrtE__ns1d which is illegal in SPIR and may result in a compilation failure [-Wsycl-undef-func]
382 warnings generated.
Source File
#include <CL/sycl.hpp>
#include <iostream>
namespace sycl = cl::sycl;
auto exception_handler(sycl::exception_list exceptions)
{
for (std::exception_ptr const &e_ptr : exceptions) {
try {
std::rethrow_exception(e_ptr);
} catch (sycl::exception const &e) {
std::cout << "Caught asynchronous SYCL exception:\n" << e.what() << std::endl;
}
}
}
struct PtEtaPhiEVector {
double Pt, Eta, Phi, E;
PtEtaPhiEVector(double _pt, double _eta, double _phi, double _e) : Pt(_pt), Eta(_eta), Phi(_phi), E(_e) {}
};
class invariant_masses;
double *InvariantMassSYCL(const PtEtaPhiEVector *v1, const PtEtaPhiEVector *v2, size_t size)
{
double *invMasses = new double[size];
sycl::gpu_selector device_selector;
sycl::queue queue(device_selector, exception_handler);
std::cout << "Running InvariantMassSYCL on " << queue.get_device().get_info<sycl::info::device::name>() << "\n";
{
sycl::buffer<PtEtaPhiEVector, 1> v1_sycl(v1, sycl::range<1>(size));
sycl::buffer<PtEtaPhiEVector, 1> v2_sycl(v2, sycl::range<1>(size));
sycl::buffer<double, 1> im_sycl(invMasses, sycl::range<1>(size));
queue.submit([&](sycl::handler &cgh) {
auto v1_acc = v1_sycl.get_access<sycl::access::mode::read>(cgh);
auto v2_acc = v2_sycl.get_access<sycl::access::mode::read>(cgh);
auto im_acc = im_sycl.get_access<sycl::access::mode::discard_write>(cgh);
cgh.parallel_for<class invariant_masses>(sycl::range<1>(size), [=](sycl::item<1> item) {
size_t id = item.get_linear_id();
auto const local_v1 = v1_acc[id];
auto const local_v2 = v2_acc[id];
// Conversion from (pt, Eta, phi, mass) to (x, y, z, e) coordinate system
const auto x1 = local_v1.Pt * sycl::cos(local_v1.Phi);
const auto y1 = local_v1.Pt * sycl::sin(local_v1.Phi);
const auto z1 = local_v1.Pt * sycl::sinh(local_v1.Eta);
const auto e1 = local_v1.E;
const auto x2 = local_v2.Pt * sycl::cos(local_v2.Phi);
const auto y2 = local_v2.Pt * sycl::sin(local_v2.Phi);
const auto z2 = local_v2.Pt * sycl::sinh(local_v2.Eta);
const auto e2 = local_v2.E;
// Addition of particle four-vector elements
const auto e = e1 + e2;
const auto x = x1 + x2;
const auto y = y1 + y2;
const auto z = z1 + z2;
auto mm = e * e - x * x - y * y - z * z;
im_acc[id] = mm < 0 ? -sycl::sqrt(-mm) : sycl::sqrt(mm);
});
});
}
try {
queue.wait_and_throw();
} catch (sycl::exception const &e) {
std::cout << "Caught synchronous SYCL exception:\n" << e.what() << std::endl;
}
return invMasses;
}
int main(int argc, char const *argv[])
{
std::vector<double> e1 = {50, 50, 50, 50, 100};
std::vector<double> pt1 = {0, 5, 5, 10, 10};
std::vector<double> eta1 = {0.0, 0.0, -1.0, 0.5, 2.5};
std::vector<double> phi1 = {0.0, 0.0, 0.0, -0.5, -2.4};
std::vector<double> e2 = {40, 40, 40, 40, 30};
std::vector<double> pt2 = {0, 5, 5, 10, 2};
std::vector<double> eta2 = {0.0, 0.0, 0.5, 0.4, 1.2};
std::vector<double> phi2 = {0.0, 0.0, 0.0, 0.5, 2.4};
std::vector<PtEtaPhiEVector> p1, p2;
for (size_t i = 0; i < e1.size(); i++) {
p1.push_back(PtEtaPhiEVector(pt1[i], eta1[i], phi1[i], e1[i]));
p2.push_back(PtEtaPhiEVector(pt2[i], eta2[i], phi2[i], e2[i]));
}
const auto SYCLinvMasses = InvariantMassSYCL(p1.data(), p2.data(), e1.size());
std::vector<double> expectedInvMass = {90.000000, 89.442719, 89.382905, 87.778731, 112.949487};
for (int i = 0; i < e1.size(); i++) {
if (fabs(SYCLinvMasses[i] - expectedInvMass[i]) > 1e-6) {
printf("The difference between expectedInvMass[%d] and SYCLinvMasses[%d] is %f, which exceeds 1e-6, where"
"expectedInvMass[% d] evaluates to %f, SYCLinvMasses[% d] evaluates to % f",
i, i, fabs(SYCLinvMasses[i] - expectedInvMass[i]), i, expectedInvMass[i], i, SYCLinvMasses[i]);
}
}
return 0;
}
CMakeLists.txt
cmake_minimum_required(VERSION 3.4.3)
project(test_application)
set(SRC_FILES "reproducer.cpp")
set(SYCL_FILES "reproducer.cpp")
list(APPEND CMAKE_MODULE_PATH "${CMAKE_SOURCE_DIR}/cmake")
set(ComputeCpp_DIR CACHE STRING "NOT-FOUND")
if (NOT ComputeCpp_DIR)
message(FATAL_ERROR
"SYCL implementation root not provided, please specify "
"the path to the root of the chosen SYCL implementation using "
"ComputeCpp_DIR=<path/to/install/root>.")
endif()
find_package(ComputeCpp REQUIRED)
add_executable("${PROJECT_NAME}" "${SRC_FILES}" "${HEADER_FILES}")
target_include_directories("${PROJECT_NAME}" PRIVATE
"${ComputeCpp_INCLUDE_DIRS}"
"${CMAKE_SOURCE_DIR}/src")
target_compile_definitions("${PROJECT_NAME}" PRIVATE CL_TARGET_OPENCL_VERSION=300 )
list(APPEND COMPUTECPP_USER_FLAGS -O2 -std=c++17 -no-serial-memop -DCOMPUTE_CPP_BITCODE=ptx64 -sycl-target ptx64)
add_sycl_to_target(TARGET "${PROJECT_NAME}" SOURCES "${SYCL_FILES}")