Skip to content

SYCL Compiler Support, main branch (2024.11.05.) #301

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
krasznaa merged 3 commits into from
Nov 7, 2024
Merged

SYCL Compiler Support, main branch (2024.11.05.) #301

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
47b0439
Switched to using <sycl/sycl.hpp>.
krasznaa Nov 5, 2024
415a39d
Introduced explicit support for AdaptiveCpp 24.06.0.
krasznaa Nov 5, 2024
9df45ce
Removed the last mentions of <CL/sycl.hpp> and cl::sycl.
krasznaa Nov 6, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,8 @@ set( CMAKE_ARCHIVE_OUTPUT_DIRECTORY
# Flags controlling the meta-build system.
option( VECMEM_USE_SYSTEM_LIBS "Use system libraries by default" FALSE )
option( VECMEM_BUILD_TESTING "Build the unit tests of VecMem" TRUE )
option( VECMEM_TEST_UBSAN "Use the undefined behavior sanitizer for the tests"
TRUE )

# Include the VecMem CMake code.
list( APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake" )
Expand Down
19 changes: 11 additions & 8 deletions cmake/sycl/CMakeDetermineSYCLCompiler.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,16 +24,15 @@ if( NOT "$ENV{SYCLCXX}" STREQUAL "" )
endif()

# Determine the type and version of the SYCL compiler.
execute_process( COMMAND "${CMAKE_SYCL_COMPILER_INIT}" "--version"
OUTPUT_VARIABLE _syclVersionOutput
ERROR_VARIABLE _syclVersionError
RESULT_VARIABLE _syclVersionResult )
if( NOT ${_syclVersionResult} EQUAL 0 )
execute_process( COMMAND "${CMAKE_SYCL_COMPILER_INIT}"
"--hipsycl-version"
foreach( _version_cmdl "--acpp-version" "--version" "--hipsycl-version" )
execute_process( COMMAND "${CMAKE_SYCL_COMPILER_INIT}" "${_version_cmdl}"
OUTPUT_VARIABLE _syclVersionOutput
ERROR_VARIABLE _syclVersionError
RESULT_VARIABLE _syclVersionResult )
endif()
if( ${_syclVersionResult} EQUAL 0 )
break()
endif()
endforeach()
if( ${_syclVersionResult} EQUAL 0 )
if( "${_syclVersionOutput}" MATCHES "ComputeCpp" )
set( CMAKE_SYCL_COMPILER_ID "ComputeCpp" CACHE STRING
Expand All @@ -47,6 +46,10 @@ if( NOT "$ENV{SYCLCXX}" STREQUAL "" )
set( CMAKE_SYCL_COMPILER_ID "IntelLLVM" CACHE STRING
"Identifier for the SYCL compiler in use" )
set( _syclVersionRegex "clang version ([0-9\.]+)" )
elseif( "${_syclVersionOutput}" MATCHES "AdaptiveCpp" )
set( CMAKE_SYCL_COMPILER_ID "AdaptiveCpp" CACHE STRING
"Identifier for the SYCL compiler in use" )
set( _syclVersionRegex "AdaptiveCpp version: ([0-9\.]+)" )
elseif( "${_syclVersionOutput}" MATCHES "hipSYCL" )
set( CMAKE_SYCL_COMPILER_ID "hipSYCL" CACHE STRING
"Identifier for the SYCL compiler in use" )
Expand Down
4 changes: 2 additions & 2 deletions cmake/sycl/CMakeTestSYCLCompiler.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
# VecMem project, part of the ACTS project (R&D line)
#
# (c) 2021 CERN for the benefit of the ACTS project
# (c) 2021-2024 CERN for the benefit of the ACTS project
#
# Mozilla Public License Version 2.0

Expand All @@ -17,7 +17,7 @@ endif()
# Try to use the HIP compiler.
file( WRITE
"${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/main.sycl"
"#include <CL/sycl.hpp>\n"
"#include <sycl/sycl.hpp>\n"
Comment thread
krasznaa marked this conversation as resolved.
"int main() {\n"
"#if (!defined(CL_SYCL_LANGUAGE_VERSION)) &&"
" (!defined(SYCL_LANGUAGE_VERSION))\n"
Expand Down
31 changes: 31 additions & 0 deletions cmake/sycl/Platform/Linux-AdaptiveCpp-SYCL.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,31 @@
# VecMem project, part of the ACTS project (R&D line)
#
# (c) 2022-2024 CERN for the benefit of the ACTS project
#
# Mozilla Public License Version 2.0

# Use the standard GNU compiler options for hipSYCL.
include( Platform/Linux-GNU )
__linux_compiler_gnu( SYCL )
Comment thread
krasznaa marked this conversation as resolved.
include( Compiler/GNU )
__compiler_gnu( SYCL )

# Set up the dependency file generation for this platform. Note that SYCL
# compilation only works with Makefile and Ninja generators, so no check is made
Comment thread
krasznaa marked this conversation as resolved.
# here for the current generator.
set( CMAKE_SYCL_DEPENDS_USE_COMPILER TRUE )
set( CMAKE_SYCL_DEPFILE_FORMAT gcc )

# Set an archive (static library) creation command explicitly for this platform.
set( CMAKE_SYCL_CREATE_STATIC_LIBRARY
"<CMAKE_AR> qc <TARGET> <LINK_FLAGS> <OBJECTS>" )

# Set the flags controlling the C++ standard used by the SYCL compiler.
set( CMAKE_SYCL17_STANDARD_COMPILE_OPTION "-std=c++17" )
set( CMAKE_SYCL17_EXTENSION_COMPILE_OPTION "-std=c++17" )

set( CMAKE_SYCL20_STANDARD_COMPILE_OPTION "-std=c++20" )
set( CMAKE_SYCL20_EXTENSION_COMPILE_OPTION "-std=c++20" )

set( CMAKE_SYCL23_STANDARD_COMPILE_OPTION "-std=c++23" )
set( CMAKE_SYCL23_EXTENSION_COMPILE_OPTION "-std=c++23" )
20 changes: 10 additions & 10 deletions core/cmake/vecmem-setup-core.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,39 +44,39 @@ function( vecmem_setup_core libName )

# Test which SYCL printf function(s) is/are available.
vecmem_check_sycl_source_compiles( "
#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>
#ifdef __SYCL_DEVICE_ONLY__
# define VECMEM_MSG_ATTRIBUTES __attribute__((opencl_constant))
#else
# define VECMEM_MSG_ATTRIBUTES
#endif
int main() {
const VECMEM_MSG_ATTRIBUTES char __msg[] = \"Test message %i\";
cl::sycl::ext::oneapi::experimental::printf(__msg, 20);
::sycl::ext::oneapi::experimental::printf(__msg, 20);
return 0;
}
" VECMEM_HAVE_SYCL_EXT_ONEAPI_PRINTF )
vecmem_check_sycl_source_compiles( "
#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>
#ifdef __SYCL_DEVICE_ONLY__
# define VECMEM_MSG_ATTRIBUTES __attribute__((opencl_constant))
#else
# define VECMEM_MSG_ATTRIBUTES
#endif
int main() {
const VECMEM_MSG_ATTRIBUTES char __msg[] = \"Test message %i\";
cl::sycl::ONEAPI::experimental::printf(__msg, 20);
::sycl::ONEAPI::experimental::printf(__msg, 20);
return 0;
}
" VECMEM_HAVE_SYCL_ONEAPI_PRINTF )

# Set up the appropriate flag based on these checks.
if( VECMEM_HAVE_SYCL_EXT_ONEAPI_PRINTF )
target_compile_definitions( ${libName} INTERFACE
$<BUILD_INTERFACE:VECMEM_SYCL_PRINTF_FUNCTION=cl::sycl::ext::oneapi::experimental::printf> )
$<BUILD_INTERFACE:VECMEM_SYCL_PRINTF_FUNCTION=::sycl::ext::oneapi::experimental::printf> )
elseif( VECMEM_HAVE_SYCL_ONEAPI_PRINTF )
target_compile_definitions( ${libName} INTERFACE
$<BUILD_INTERFACE:VECMEM_SYCL_PRINTF_FUNCTION=cl::sycl::ONEAPI::experimental::printf> )
$<BUILD_INTERFACE:VECMEM_SYCL_PRINTF_FUNCTION=::sycl::ONEAPI::experimental::printf> )
else()
message( WARNING "No valid printf function found for SYCL."
" Enabling debug messages will likely not work in device code." )
Expand All @@ -87,12 +87,12 @@ function( vecmem_setup_core libName )

# Test whether sycl::atomic_ref is available.
vecmem_check_sycl_source_compiles( "
#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>
int main() {
int dummy = 0;
cl::sycl::atomic_ref<int, sycl::memory_order::relaxed,
cl::sycl::memory_scope::device,
cl::sycl::access::address_space::global_space>
::sycl::atomic_ref<int, sycl::memory_order::relaxed,
::sycl::memory_scope::device,
::sycl::access::address_space::global_space>
atomic_dummy(dummy);
atomic_dummy.store(3);
atomic_dummy.fetch_add(1);
Expand Down
8 changes: 4 additions & 4 deletions core/include/vecmem/containers/impl/device_vector.ipp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
/* VecMem project, part of the ACTS project (R&D line)
*
* (c) 2021-2022 CERN for the benefit of the ACTS project
* (c) 2021-2024 CERN for the benefit of the ACTS project
*
* Mozilla Public License Version 2.0
*/
Expand Down Expand Up @@ -197,7 +197,7 @@ VECMEM_HOST_AND_DEVICE auto device_vector<TYPE>::emplace_back(Args&&... args)
// Increment the size of the vector at first. So that we would "claim" the
// index from other threads.
device_atomic_ref<size_type> asize(*m_size);
const size_type index = asize.fetch_add(1);
const size_type index = asize.fetch_add(1u);
assert(index < m_capacity);

// Instantiate the new value.
Expand All @@ -217,7 +217,7 @@ VECMEM_HOST_AND_DEVICE auto device_vector<TYPE>::push_back(
// Increment the size of the vector at first. So that we would "claim" the
// index from other threads.
device_atomic_ref<size_type> asize(*m_size);
const size_type index = asize.fetch_add(1);
const size_type index = asize.fetch_add(1u);
assert(index < m_capacity);

// Instantiate the new value.
Expand Down Expand Up @@ -315,7 +315,7 @@ VECMEM_HOST_AND_DEVICE auto device_vector<TYPE>::pop_back() -> size_type {

// Decrement the size of the vector, and remember this new size.
device_atomic_ref<size_type> asize(*m_size);
const size_type new_size = asize.fetch_sub(1) - 1;
const size_type new_size = asize.fetch_sub(1u) - 1;

// Remove the last element.
destruct(new_size);
Expand Down
2 changes: 1 addition & 1 deletion core/include/vecmem/edm/impl/device.ipp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,7 +73,7 @@ device<schema<VARTYPES...>, INTERFACE>::push_back_default() -> size_type {
// Increment the size of the container at first. So that we would "claim"
// the index from other threads.
device_atomic_ref<size_type> asize(*m_size);
const size_type index = asize.fetch_add(1);
const size_type index = asize.fetch_add(1u);
assert(index < m_capacity);

// Construct the new elements in all of the vector variables.
Expand Down
29 changes: 13 additions & 16 deletions core/include/vecmem/memory/details/sycl_builtin_device_atomic_ref.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,7 +11,7 @@
#include "vecmem/memory/device_address_space.hpp"

// SYCL include(s).
#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

namespace vecmem {
namespace sycl {
Expand All @@ -27,23 +27,20 @@ struct builtin_address_space {};
/// Specialization for global device memory
template <>
struct builtin_address_space<device_address_space::global> {
static constexpr cl::sycl::memory_order ord =
cl::sycl::memory_order::relaxed;
static constexpr cl::sycl::memory_scope scp =
cl::sycl::memory_scope::device;
static constexpr cl::sycl::access::address_space add =
cl::sycl::access::address_space::global_space;
static constexpr ::sycl::memory_order ord = ::sycl::memory_order::relaxed;
static constexpr ::sycl::memory_scope scp = ::sycl::memory_scope::device;
static constexpr ::sycl::access::address_space add =
::sycl::access::address_space::global_space;
};

/// Specialization for local device memory
template <>
struct builtin_address_space<device_address_space::local> {
static constexpr cl::sycl::memory_order ord =
cl::sycl::memory_order::relaxed;
static constexpr cl::sycl::memory_scope scp =
cl::sycl::memory_scope::work_group;
static constexpr cl::sycl::access::address_space add =
cl::sycl::access::address_space::local_space;
static constexpr ::sycl::memory_order ord = ::sycl::memory_order::relaxed;
static constexpr ::sycl::memory_scope scp =
::sycl::memory_scope::work_group;
static constexpr ::sycl::access::address_space add =
::sycl::access::address_space::local_space;
};

} // namespace details
Expand All @@ -52,9 +49,9 @@ struct builtin_address_space<device_address_space::local> {
template <typename T,
device_address_space address = device_address_space::global>
using builtin_device_atomic_ref =
cl::sycl::atomic_ref<T, details::builtin_address_space<address>::ord,
details::builtin_address_space<address>::scp,
details::builtin_address_space<address>::add>;
::sycl::atomic_ref<T, details::builtin_address_space<address>::ord,
Comment thread
krasznaa marked this conversation as resolved.
details::builtin_address_space<address>::scp,
details::builtin_address_space<address>::add>;

} // namespace sycl
} // namespace vecmem
21 changes: 10 additions & 11 deletions core/include/vecmem/memory/impl/atomic.ipp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,22 +14,21 @@

// SYCL include(s).
#if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>
#endif

/// Helpers for explicit calls to the SYCL atomic functions
#if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
#define __VECMEM_SYCL_ATOMIC_CALL0(FNAME, PTR) \
cl::sycl::atomic_##FNAME<value_type>( \
cl::sycl::atomic<value_type>(cl::sycl::global_ptr<value_type>(PTR)))
#define __VECMEM_SYCL_ATOMIC_CALL1(FNAME, PTR, ARG1) \
cl::sycl::atomic_##FNAME<value_type>( \
cl::sycl::atomic<value_type>(cl::sycl::global_ptr<value_type>(PTR)), \
ARG1)
#define __VECMEM_SYCL_ATOMIC_CALL2(FNAME, PTR, ARG1, ARG2) \
cl::sycl::atomic_##FNAME<value_type>( \
cl::sycl::atomic<value_type>(cl::sycl::global_ptr<value_type>(PTR)), \
ARG1, ARG2)
::sycl::atomic_##FNAME<value_type>( \
::sycl::atomic<value_type>(::sycl::global_ptr<value_type>(PTR)))
#define __VECMEM_SYCL_ATOMIC_CALL1(FNAME, PTR, ARG1) \
Comment thread
krasznaa marked this conversation as resolved.
::sycl::atomic_##FNAME<value_type>( \
::sycl::atomic<value_type>(::sycl::global_ptr<value_type>(PTR)), ARG1)
#define __VECMEM_SYCL_ATOMIC_CALL2(FNAME, PTR, ARG1, ARG2) \
::sycl::atomic_##FNAME<value_type>( \
::sycl::atomic<value_type>(::sycl::global_ptr<value_type>(PTR)), ARG1, \
ARG2)
#endif

namespace vecmem {
Expand Down
38 changes: 19 additions & 19 deletions core/include/vecmem/memory/impl/sycl_custom_device_atomic_ref.ipp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,7 +8,7 @@
#pragma once

// SYCL include(s).
#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

namespace vecmem {
namespace sycl {
Expand All @@ -19,43 +19,43 @@ struct custom_address_space {};

template <>
struct custom_address_space<device_address_space::global> {
static constexpr cl::sycl::access::address_space add =
cl::sycl::access::address_space::global_space;
static constexpr ::sycl::access::address_space add =
::sycl::access::address_space::global_space;

template <typename T>
using ptr_t = cl::sycl::global_ptr<T>;
using ptr_t = ::sycl::global_ptr<T>;
};

template <>
struct custom_address_space<device_address_space::local> {
static constexpr cl::sycl::access::address_space add =
cl::sycl::access::address_space::local_space;
static constexpr ::sycl::access::address_space add =
::sycl::access::address_space::local_space;
template <typename T>
using ptr_t = cl::sycl::local_ptr<T>;
using ptr_t = ::sycl::local_ptr<T>;
};

} // namespace details

#define __VECMEM_SYCL_ATOMIC_CALL0(FNAME, PTR) \
cl::sycl::atomic_##FNAME<value_type, \
details::custom_address_space<address>::add>( \
cl::sycl::atomic<value_type, \
details::custom_address_space<address>::add>( \
::sycl::atomic_##FNAME<value_type, \
details::custom_address_space<address>::add>( \
::sycl::atomic<value_type, \
details::custom_address_space<address>::add>( \
typename details::custom_address_space<address>::template ptr_t< \
value_type>(PTR)))
#define __VECMEM_SYCL_ATOMIC_CALL1(FNAME, PTR, ARG1) \
cl::sycl::atomic_##FNAME<value_type, \
details::custom_address_space<address>::add>( \
cl::sycl::atomic<value_type, \
details::custom_address_space<address>::add>( \
::sycl::atomic_##FNAME<value_type, \
details::custom_address_space<address>::add>( \
::sycl::atomic<value_type, \
details::custom_address_space<address>::add>( \
typename details::custom_address_space<address>::template ptr_t< \
value_type>(PTR)), \
ARG1)
#define __VECMEM_SYCL_ATOMIC_CALL2(FNAME, PTR, ARG1, ARG2) \
cl::sycl::atomic_##FNAME<value_type, \
details::custom_address_space<address>::add>( \
cl::sycl::atomic<value_type, \
details::custom_address_space<address>::add>( \
::sycl::atomic_##FNAME<value_type, \
details::custom_address_space<address>::add>( \
::sycl::atomic<value_type, \
details::custom_address_space<address>::add>( \
typename details::custom_address_space<address>::template ptr_t< \
value_type>(PTR)), \
ARG1, ARG2)
Expand Down
2 changes: 1 addition & 1 deletion core/include/vecmem/memory/memory_order.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@
#if (defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)) && \
defined(VECMEM_HAVE_SYCL_ATOMIC_REF)
// SYCL include(s).
#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>
#endif

namespace vecmem {
Expand Down
6 changes: 3 additions & 3 deletions sycl/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,7 +41,7 @@ set_target_properties( vecmem_sycl PROPERTIES
CXX_VISIBILITY_PRESET "hidden"
SYCL_VISIBILITY_PRESET "hidden" )
vecmem_check_sycl_source_compiles( "
#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>
int main() { return 0; }
"
VECMEM_HAVE_SYCL_VISIBILITY_MS_COMPAT
Expand Down Expand Up @@ -94,9 +94,9 @@ endif()
# Check if sycl::queue::memset is available, and set a compiler option
# accordingly.
vecmem_check_sycl_source_compiles( "
#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>
int main() {
cl::sycl::queue queue;
::sycl::queue queue;
queue.memset(nullptr, 0, 100);
return 0;
}
Expand Down
Loading