diff --git a/CMakeLists.txt b/CMakeLists.txt index 88f9b732..2830e644 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -159,8 +159,7 @@ register_model(hip HIP HIPStream.cpp) register_model(cuda CUDA CUDAStream.cu) register_model(kokkos KOKKOS KokkosStream.cpp) register_model(sycl SYCL SYCLStream.cpp) -register_model(sycl2020-acc SYCL2020 SYCLStream2020.cpp) -register_model(sycl2020-usm SYCL2020 SYCLStream2020.cpp) +register_model(sycl2020 SYCL2020 SYCLStream2020.cpp) register_model(acc ACC ACCStream.cpp) # defining RAJA collides with the RAJA namespace so USE_RAJA register_model(raja USE_RAJA RAJAStream.cpp) diff --git a/src/std/STDStream.cpp b/src/std/STDStream.cpp index 678457ae..47b1ebce 100644 --- a/src/std/STDStream.cpp +++ b/src/std/STDStream.cpp @@ -8,7 +8,7 @@ #include #include -#if defined(DATA20) || defined(INDICES) +#if defined(DATA23) || defined(INDICES) #include #endif @@ -89,7 +89,7 @@ template void STDStream::copy() { // c[i] = a[i] -#if defined(DATA17) || defined(DATA20) +#if defined(DATA17) || defined(DATA23) std::copy(exe_policy, a, a + array_size, c); #elif INDICES std::for_each_n(exe_policy, counting_iter(0), array_size, [a=a,c=c](intptr_t i) { @@ -104,7 +104,7 @@ template void STDStream::mul() { // b[i] = scalar * c[i]; -#if defined(DATA17) || defined(DATA20) +#if defined(DATA17) || defined(DATA23) std::transform(exe_policy, c, c + array_size, b, [](T ci){ return startScalar*ci; }); #elif INDICES std::for_each_n(exe_policy, counting_iter(0), array_size, [b=b, c=c](intptr_t i) { @@ -119,7 +119,7 @@ template void STDStream::add() { // c[i] = a[i] + b[i]; -#if defined(DATA17) || defined(DATA20) +#if defined(DATA17) || defined(DATA23) std::transform(exe_policy, a, a + array_size, b, c, std::plus()); #elif INDICES std::for_each_n(exe_policy, counting_iter(0), array_size, [a=a, b=b, c=c](intptr_t i) { @@ -134,7 +134,7 @@ template void STDStream::triad() { // a[i] = b[i] + scalar * c[i]; -#if defined(DATA17) || defined(DATA20) +#if defined(DATA17) || defined(DATA23) std::transform(exe_policy, b, b + array_size, c, a, [scalar = startScalar](T bi, T ci){ return bi+scalar*ci; }); #elif INDICES std::for_each_n(exe_policy, counting_iter(0), array_size, [a=a, b=b, c=c](intptr_t i) { @@ -149,13 +149,13 @@ template void STDStream::nstream() { // a[i] += b[i] + scalar * c[i]; -#if defined(DATA17) || defined(DATA20) // Until we can require GCC 14.1 +#if defined(DATA17) || defined(DATA23) // Until we can require GCC 14.1 // Need to do in two round-trips with C++17 STL. // 1: a[i] += b[i] // 2: a[i] += scalar * c[i]; std::transform(exe_policy, a, a + array_size, b, a, [](T ai, T bi){ return ai + bi; }); std::transform(exe_policy, a, a + array_size, c, a, [](T ai, T ci){ return ai + startScalar*ci; }); -#elif DATA20 +#elif DATA23 // Requires GCC 14.1 (Ubuntu 24.04): auto as = std::ranges::subrange(a, a + array_size); auto bs = std::ranges::subrange(b, b + array_size); @@ -178,7 +178,7 @@ void STDStream::nstream() template T STDStream::dot() { -#if defined(DATA17) || defined(DATA20) +#if defined(DATA17) || defined(DATA23) // sum = 0; sum += a[i] * b[i]; return sum; return std::transform_reduce(exe_policy, a, a + array_size, b, T{0}); #elif INDICES diff --git a/src/std/STDStream.h b/src/std/STDStream.h index aa26eb33..254d68d7 100644 --- a/src/std/STDStream.h +++ b/src/std/STDStream.h @@ -12,8 +12,8 @@ #ifdef DATA17 #define STDIMPL "DATA17" -#elif DATA20 -#define STDIMPL "DATA20" +#elif DATA23 +#define STDIMPL "DATA23" #elif INDICES #define STDIMPL "INDICES" #else diff --git a/src/sycl2020-acc/SYCLStream2020.cpp b/src/sycl2020-acc/SYCLStream2020.cpp deleted file mode 100644 index d0f97e68..00000000 --- a/src/sycl2020-acc/SYCLStream2020.cpp +++ /dev/null @@ -1,286 +0,0 @@ - -// Copyright (c) 2015-16 Tom Deakin, Simon McIntosh-Smith, -// University of Bristol HPC -// -// For full license terms please see the LICENSE file distributed with this -// source code - -#include "SYCLStream2020.h" - -#include - -// Cache list of devices -bool cached = false; -std::vector devices; -void getDeviceList(void); - -template -SYCLStream::SYCLStream(BenchId bs, const intptr_t array_size, const int device_index, - T initA, T initB, T initC) - : array_size(array_size), - d_a {array_size}, - d_b {array_size}, - d_c {array_size}, - d_sum {1} -{ - if (!cached) - getDeviceList(); - - if (device_index >= devices.size()) - throw std::runtime_error("Invalid device index"); - - sycl::device dev = devices[device_index]; - - // Print out device information - std::cout << "Using SYCL device " << getDeviceName(device_index) << std::endl; - std::cout << "Driver: " << getDeviceDriver(device_index) << std::endl; - - // Check device can support FP64 if needed - if (sizeof(T) == sizeof(double)) - { - if (!dev.has(sycl::aspect::fp64)) - { - throw std::runtime_error("Device does not support double precision, please use --float"); - } - } - - queue = std::make_unique(dev, sycl::async_handler{[&](sycl::exception_list l) - { - bool error = false; - for(auto e: l) - { - try - { - std::rethrow_exception(e); - } - catch (sycl::exception e) - { - std::cout << e.what(); - error = true; - } - } - if(error) - { - throw std::runtime_error("SYCL errors detected"); - } - }}); - - // No longer need list of devices - devices.clear(); - cached = true; - - init_arrays(initA, initB, initC); -} - - -template -void SYCLStream::copy() -{ - queue->submit([&](sycl::handler &cgh) - { - sycl::accessor ka {d_a, cgh, sycl::read_only}; - sycl::accessor kc {d_c, cgh, sycl::write_only}; - cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) - { - kc[idx] = ka[idx]; - }); - }); - queue->wait(); -} - -template -void SYCLStream::mul() -{ - const T scalar = startScalar; - queue->submit([&](sycl::handler &cgh) - { - sycl::accessor kb {d_b, cgh, sycl::write_only}; - sycl::accessor kc {d_c, cgh, sycl::read_only}; - cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) - { - kb[idx] = scalar * kc[idx]; - }); - }); - queue->wait(); -} - -template -void SYCLStream::add() -{ - queue->submit([&](sycl::handler &cgh) - { - sycl::accessor ka {d_a, cgh, sycl::read_only}; - sycl::accessor kb {d_b, cgh, sycl::read_only}; - sycl::accessor kc {d_c, cgh, sycl::write_only}; - cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) - { - kc[idx] = ka[idx] + kb[idx]; - }); - }); - queue->wait(); -} - -template -void SYCLStream::triad() -{ - const T scalar = startScalar; - queue->submit([&](sycl::handler &cgh) - { - sycl::accessor ka {d_a, cgh, sycl::write_only}; - sycl::accessor kb {d_b, cgh, sycl::read_only}; - sycl::accessor kc {d_c, cgh, sycl::read_only}; - cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) - { - ka[idx] = kb[idx] + scalar * kc[idx]; - }); - }); - queue->wait(); -} - -template -void SYCLStream::nstream() -{ - const T scalar = startScalar; - - queue->submit([&](sycl::handler &cgh) - { - sycl::accessor ka {d_a, cgh}; - sycl::accessor kb {d_b, cgh, sycl::read_only}; - sycl::accessor kc {d_c, cgh, sycl::read_only}; - cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) - { - ka[idx] += kb[idx] + scalar * kc[idx]; - }); - }); - queue->wait(); -} - -template -T SYCLStream::dot() -{ - - queue->submit([&](sycl::handler &cgh) - { - sycl::accessor ka {d_a, cgh, sycl::read_only}; - sycl::accessor kb {d_b, cgh, sycl::read_only}; - - cgh.parallel_for(sycl::range<1>{array_size}, - // Reduction object, to perform summation - initialises the result to zero - // AdaptiveCpp doesn't sypport the initialize_to_identity property yet -#if defined(__HIPSYCL__) || defined(__OPENSYCL__) || defined(__ADAPTIVECPP__) - sycl::reduction(d_sum. template get_access(cgh), sycl::plus()), -#else - sycl::reduction(sum, sycl::plus(), sycl::property::reduction::initialize_to_identity{}), -#endif - [=](sycl::id<1> idx, auto& sum) - { - sum += ka[idx] * kb[idx]; - }); - }); - - // Get access on the host, and return a copy of the data (single number) - // This will block until the result is available, so no need to wait on the queue. - sycl::host_accessor result {d_sum, sycl::read_only}; - return result[0]; - -} - -template -void SYCLStream::init_arrays(T initA, T initB, T initC) -{ - queue->submit([&](sycl::handler &cgh) - { - sycl::accessor ka {d_a, cgh, sycl::write_only, sycl::no_init}; - sycl::accessor kb {d_b, cgh, sycl::write_only, sycl::no_init}; - sycl::accessor kc {d_c, cgh, sycl::write_only, sycl::no_init}; - - cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) - { - ka[idx] = initA; - kb[idx] = initB; - kc[idx] = initC; - }); - }); - - queue->wait(); -} - -template -void SYCLStream::get_arrays(T const*& a, T const*& b, T const*& c) -{ - sycl::host_accessor _a {d_a, sycl::read_only}; - sycl::host_accessor _b {d_b, sycl::read_only}; - sycl::host_accessor _c {d_c, sycl::read_only}; - a = &_a[0]; - b = &_b[0]; - c = &_c[0]; -} - -void getDeviceList(void) -{ - // Ask SYCL runtime for all devices in system - devices = sycl::device::get_devices(); - cached = true; -} - -void listDevices(void) -{ - getDeviceList(); - - // Print device names - if (devices.size() == 0) - { - std::cerr << "No devices found." << std::endl; - } - else - { - std::cout << std::endl; - std::cout << "Devices:" << std::endl; - for (int i = 0; i < devices.size(); i++) - { - std::cout << i << ": " << getDeviceName(i) << std::endl; - } - std::cout << std::endl; - } -} - -std::string getDeviceName(const int device) -{ - if (!cached) - getDeviceList(); - - std::string name; - - if (device < devices.size()) - { - name = devices[device].get_info(); - } - else - { - throw std::runtime_error("Error asking for name for non-existant device"); - } - - return name; -} - -std::string getDeviceDriver(const int device) -{ - if (!cached) - getDeviceList(); - - std::string driver; - - if (device < devices.size()) - { - driver = devices[device].get_info(); - } - else - { - throw std::runtime_error("Error asking for driver for non-existant device"); - } - - return driver; -} - -template class SYCLStream; -template class SYCLStream; diff --git a/src/sycl2020-acc/SYCLStream2020.h b/src/sycl2020-acc/SYCLStream2020.h deleted file mode 100644 index c0caae2e..00000000 --- a/src/sycl2020-acc/SYCLStream2020.h +++ /dev/null @@ -1,54 +0,0 @@ - -// Copyright (c) 2015-16 Tom Deakin, Simon McIntosh-Smith, -// University of Bristol HPC -// -// For full license terms please see the LICENSE file distributed with this -// source code - -#pragma once - -#include -#include - -#include "Stream.h" - -#include - -#define IMPLEMENTATION_STRING "SYCL2020 accessors" - -template -class SYCLStream : public Stream -{ - protected: - // Size of arrays - size_t array_size; - - // SYCL objects - // Queue is a pointer because we allow device selection - std::unique_ptr queue; - - // Buffers - sycl::buffer d_a; - sycl::buffer d_b; - sycl::buffer d_c; - sycl::buffer d_sum; - - public: - - SYCLStream(BenchId bs, const intptr_t array_size, const int device_id, - T initA, T initB, T initC); - ~SYCLStream() = default; - - void copy() override; - void add() override; - void mul() override; - void triad() override; - void nstream() override; - T dot() override; - - void get_arrays(T const*& a, T const*& b, T const*& c) override; - void init_arrays(T initA, T initB, T initC); -}; - -// Populate the devices list -void getDeviceList(void); diff --git a/src/sycl2020-acc/model.cmake b/src/sycl2020-acc/model.cmake deleted file mode 100644 index 9847b348..00000000 --- a/src/sycl2020-acc/model.cmake +++ /dev/null @@ -1,92 +0,0 @@ - -register_flag_optional(CMAKE_CXX_COMPILER - "Any CXX compiler that is supported by CMake detection, this is used for host compilation when required by the SYCL compiler" - "c++") - -register_flag_required(SYCL_COMPILER - "Compile using the specified SYCL compiler implementation - Supported values are - ONEAPI-ICPX - icpx as a standalone compiler - ONEAPI-Clang - oneAPI's Clang driver (enabled via `source /opt/intel/oneapi/setvars.sh --include-intel-llvm`) - DPCPP - dpc++ as a standalone compiler (https://github.com/intel/llvm) - HIPSYCL - hipSYCL compiler (https://github.com/illuhad/hipSYCL) - AdaptiveCpp - AdaptiveCpp compiler (https://github.com/adaptivecpp/adaptivecpp)") - -register_flag_optional(SYCL_COMPILER_DIR - "Absolute path to the selected SYCL compiler directory, most are packaged differently so set the path according to `SYCL_COMPILER`: - ONEAPI-ICPX - `icpx` must be used for OneAPI 2023 and later on releases (i.e `source /opt/intel/oneapi/setvars.sh` first) - ONEAPI-Clang - set to the directory that contains the Intel clang++ binary. - AdaptiveCpp|HIPSYCL|DPCPP - set to the root of the binary distribution that contains at least `bin/`, `include/`, and `lib/`" - "") - -macro(setup) - set(CMAKE_CXX_STANDARD 17) - - - if (${SYCL_COMPILER} STREQUAL "AdaptiveCpp") - set(adaptivecpp_DIR ${SYCL_COMPILER_DIR}/lib/cmake/adaptivecpp) - - if (NOT EXISTS "${AdaptiveCpp_DIR}") - message(WARNING "Falling back to AdaptiveCpp < 0.9.0 CMake structure") - set(AdaptiveCpp_DIR ${SYCL_COMPILER_DIR}/lib/cmake) - endif () - if (NOT EXISTS "${AdaptiveCpp_DIR}") - message(FATAL_ERROR "Can't find the appropriate CMake definitions for AdaptiveCpp") - endif () - - # register_definitions(_GLIBCXX_USE_CXX11_ABI=0) - find_package(AdaptiveCpp CONFIG REQUIRED) - message(STATUS "ok") - elseif (${SYCL_COMPILER} STREQUAL "HIPSYCL") - set(hipSYCL_DIR ${SYCL_COMPILER_DIR}/lib/cmake/hipSYCL) - - if (NOT EXISTS "${hipSYCL_DIR}") - message(WARNING "Falling back to hipSYCL < 0.9.0 CMake structure") - set(hipSYCL_DIR ${SYCL_COMPILER_DIR}/lib/cmake) - endif () - if (NOT EXISTS "${hipSYCL_DIR}") - message(FATAL_ERROR "Can't find the appropriate CMake definitions for hipSYCL") - endif () - - # register_definitions(_GLIBCXX_USE_CXX11_ABI=0) - find_package(hipSYCL CONFIG REQUIRED) - message(STATUS "ok") - elseif (${SYCL_COMPILER} STREQUAL "DPCPP") - set(CMAKE_CXX_COMPILER ${SYCL_COMPILER_DIR}/bin/clang++) - include_directories(${SYCL_COMPILER_DIR}/include/sycl) - register_append_cxx_flags(ANY -fsycl) - register_append_link_flags(-fsycl) - elseif (${SYCL_COMPILER} STREQUAL "ONEAPI-ICPX") - set(CMAKE_CXX_COMPILER icpx) - set(CMAKE_C_COMPILER icx) - register_append_cxx_flags(ANY -fsycl) - register_append_link_flags(-fsycl) - elseif (${SYCL_COMPILER} STREQUAL "ONEAPI-Clang") - set(CMAKE_CXX_COMPILER clang++) - set(CMAKE_C_COMPILER clang) - register_append_cxx_flags(ANY -fsycl) - register_append_link_flags(-fsycl) - else () - message(FATAL_ERROR "SYCL_COMPILER=${SYCL_COMPILER} is unsupported") - endif () - -endmacro() - - -macro(setup_target NAME) - if (${SYCL_COMPILER} STREQUAL "AdaptiveCpp") - # so AdaptiveCpp has this weird (and bad) CMake usage where they append their - # own custom integration header flags AFTER the target has been specified - # hence this macro here - add_sycl_to_target( - TARGET ${NAME} - SOURCES ${IMPL_SOURCES}) - elseif (${SYCL_COMPILER} STREQUAL "HIPSYCL") - # so hipSYCL has this weird (and bad) CMake usage where they append their - # own custom integration header flags AFTER the target has been specified - # hence this macro here - add_sycl_to_target( - TARGET ${NAME} - SOURCES ${IMPL_SOURCES}) - endif () -endmacro() diff --git a/src/sycl2020-usm/SYCLStream2020.cpp b/src/sycl2020-usm/SYCLStream2020.cpp deleted file mode 100644 index c8b863ad..00000000 --- a/src/sycl2020-usm/SYCLStream2020.cpp +++ /dev/null @@ -1,266 +0,0 @@ - -// Copyright (c) 2015-23 Tom Deakin, Simon McIntosh-Smith, and Tom Lin -// University of Bristol HPC -// -// For full license terms please see the LICENSE file distributed with this -// source code - -#include "SYCLStream2020.h" - -#include - -// Cache list of devices -bool cached = false; -std::vector devices; -void getDeviceList(void); - -template -SYCLStream::SYCLStream(BenchId bs, const intptr_t array_size, const int device_index, - T initA, T initB, T initC) - : array_size(array_size) -{ - if (!cached) - getDeviceList(); - - if (device_index >= devices.size()) - throw std::runtime_error("Invalid device index"); - - sycl::device dev = devices[device_index]; - - // Print out device information - std::cout << "Using SYCL device " << getDeviceName(device_index) << std::endl; - std::cout << "Driver: " << getDeviceDriver(device_index) << std::endl; - - // Check device can support FP64 if needed - if (sizeof(T) == sizeof(double)) - { - if (!dev.has(sycl::aspect::fp64)) - { - throw std::runtime_error("Device does not support double precision, please use --float"); - } - } - - queue = std::make_unique(dev, sycl::async_handler{[&](sycl::exception_list l) - { - bool error = false; - for(auto e: l) - { - try - { - std::rethrow_exception(e); - } - catch (sycl::exception e) - { - std::cout << e.what(); - error = true; - } - } - if(error) - { - throw std::runtime_error("SYCL errors detected"); - } - }}); - - a = sycl::malloc_shared(array_size, *queue); - b = sycl::malloc_shared(array_size, *queue); - c = sycl::malloc_shared(array_size, *queue); - sum = sycl::malloc_shared(1, *queue); - - // No longer need list of devices - devices.clear(); - cached = true; - - init_arrays(initA, initB, initC); -} - -template -SYCLStream::~SYCLStream() { - sycl::free(a, *queue); - sycl::free(b, *queue); - sycl::free(c, *queue); - sycl::free(sum, *queue); -} - -template -void SYCLStream::copy() -{ - queue->submit([&](sycl::handler &cgh) - { - cgh.parallel_for(sycl::range<1>{array_size}, [=, c = this->c, a = this->a](sycl::id<1> idx) - { - c[idx] = a[idx]; - }); - }); - queue->wait(); -} - -template -void SYCLStream::mul() -{ - const T scalar = startScalar; - queue->submit([&](sycl::handler &cgh) - { - cgh.parallel_for(sycl::range<1>{array_size}, [=, b = this->b, c = this->c](sycl::id<1> idx) - { - b[idx] = scalar * c[idx]; - }); - }); - queue->wait(); -} - -template -void SYCLStream::add() -{ - queue->submit([&](sycl::handler &cgh) - { - cgh.parallel_for(sycl::range<1>{array_size}, [=, c = this->c, a = this->a, b = this->b](sycl::id<1> idx) - { - c[idx] = a[idx] + b[idx]; - }); - }); - queue->wait(); -} - -template -void SYCLStream::triad() -{ - const T scalar = startScalar; - queue->submit([&](sycl::handler &cgh) - { - cgh.parallel_for(sycl::range<1>{array_size}, [=, a = this->a, b = this->b, c = this->c](sycl::id<1> idx) - { - a[idx] = b[idx] + scalar * c[idx]; - }); - }); - queue->wait(); -} - -template -void SYCLStream::nstream() -{ - const T scalar = startScalar; - - queue->submit([&](sycl::handler &cgh) - { - cgh.parallel_for(sycl::range<1>{array_size}, [=, a = this->a, b = this->b, c = this->c](sycl::id<1> idx) - { - a[idx] += b[idx] + scalar * c[idx]; - }); - }); - queue->wait(); -} - -template -T SYCLStream::dot() -{ - queue->submit([&](sycl::handler &cgh) - { - cgh.parallel_for(sycl::range<1>{array_size}, - // Reduction object, to perform summation - initialises the result to zero - // AdaptiveCpp doesn't sypport the initialize_to_identity property yet -#if defined(__HIPSYCL__) || defined(__OPENSYCL__) || defined(__ADAPTIVECPP__) - sycl::reduction(sum, sycl::plus()), -#else - sycl::reduction(sum, sycl::plus(), sycl::property::reduction::initialize_to_identity{}), -#endif - [a = this->a, b = this->b](sycl::id<1> idx, auto& sum) - { - sum += a[idx] * b[idx]; - }); - }); - queue->wait(); - return *sum; -} - -template -void SYCLStream::init_arrays(T initA, T initB, T initC) -{ - queue->submit([&](sycl::handler &cgh) - { - cgh.parallel_for(sycl::range<1>{array_size}, [=, a = this->a, b = this->b, c = this->c](sycl::id<1> idx) - { - a[idx] = initA; - b[idx] = initB; - c[idx] = initC; - }); - }); - - queue->wait(); -} - -template -void SYCLStream::get_arrays(T const*& h_a, T const*& h_b, T const*& h_c) -{ - h_a = a; - h_b = b; - h_c = c; -} - -void getDeviceList(void) -{ - // Ask SYCL runtime for all devices in system - devices = sycl::device::get_devices(); - cached = true; -} - -void listDevices(void) -{ - getDeviceList(); - - // Print device names - if (devices.size() == 0) - { - std::cerr << "No devices found." << std::endl; - } - else - { - std::cout << std::endl; - std::cout << "Devices:" << std::endl; - for (int i = 0; i < devices.size(); i++) - { - std::cout << i << ": " << getDeviceName(i) << std::endl; - } - std::cout << std::endl; - } -} - -std::string getDeviceName(const int device) -{ - if (!cached) - getDeviceList(); - - std::string name; - - if (device < devices.size()) - { - name = devices[device].get_info(); - } - else - { - throw std::runtime_error("Error asking for name for non-existant device"); - } - - return name; -} - -std::string getDeviceDriver(const int device) -{ - if (!cached) - getDeviceList(); - - std::string driver; - - if (device < devices.size()) - { - driver = devices[device].get_info(); - } - else - { - throw std::runtime_error("Error asking for driver for non-existant device"); - } - - return driver; -} - -template class SYCLStream; -template class SYCLStream; diff --git a/src/sycl2020-usm/SYCLStream2020.h b/src/sycl2020-usm/SYCLStream2020.h deleted file mode 100644 index c88c87a3..00000000 --- a/src/sycl2020-usm/SYCLStream2020.h +++ /dev/null @@ -1,54 +0,0 @@ - -// Copyright (c) 2015-16 Tom Deakin, Simon McIntosh-Smith, -// University of Bristol HPC -// -// For full license terms please see the LICENSE file distributed with this -// source code - -#pragma once - -#include -#include - -#include "Stream.h" - -#include - -#define IMPLEMENTATION_STRING "SYCL2020 USM" - -template -class SYCLStream : public Stream -{ - protected: - // Size of arrays - size_t array_size; - - // SYCL objects - // Queue is a pointer because we allow device selection - std::unique_ptr queue; - - // Buffers - T *a{}; - T *b{}; - T *c{}; - T *sum{}; - - public: - - SYCLStream(BenchId bs, const intptr_t array_size, const int device_index, - T initA, T initB, T initC); - ~SYCLStream(); - - void copy() override; - void add() override; - void mul() override; - void triad() override; - void nstream() override; - T dot() override; - - void get_arrays(T const*& a, T const*& b, T const*& c) override; - void init_arrays(T initA, T initB, T initC); -}; - -// Populate the devices list -void getDeviceList(void); diff --git a/src/sycl2020-usm/model.cmake b/src/sycl2020-usm/model.cmake deleted file mode 100644 index 72aa7c40..00000000 --- a/src/sycl2020-usm/model.cmake +++ /dev/null @@ -1,91 +0,0 @@ - -register_flag_optional(CMAKE_CXX_COMPILER - "Any CXX compiler that is supported by CMake detection, this is used for host compilation when required by the SYCL compiler" - "c++") - -register_flag_required(SYCL_COMPILER - "Compile using the specified SYCL compiler implementation - Supported values are - ONEAPI-ICPX - icpx as a standalone compiler - ONEAPI-Clang - oneAPI's Clang driver (enabled via `source /opt/intel/oneapi/setvars.sh --include-intel-llvm`) - DPCPP - dpc++ as a standalone compiler (https://github.com/intel/llvm) - HIPSYCL - hipSYCL compiler (https://github.com/illuhad/hipSYCL) - AdaptiveCpp - AdaptiveCpp compiler (https://github.com/adaptivecpp/adaptivecpp)") - -register_flag_optional(SYCL_COMPILER_DIR - "Absolute path to the selected SYCL compiler directory, most are packaged differently so set the path according to `SYCL_COMPILER`: - ONEAPI-ICPX - `icpx` must be used for OneAPI 2023 and later on releases (i.e `source /opt/intel/oneapi/setvars.sh` first) - ONEAPI-Clang - set to the directory that contains the Intel clang++ binary. - AdaptiveCpp|HIPSYCL|DPCPP - set to the root of the binary distribution that contains at least `bin/`, `include/`, and `lib/`" - "") - -macro(setup) - set(CMAKE_CXX_STANDARD 17) - - if (${SYCL_COMPILER} STREQUAL "AdaptiveCpp") - set(adaptivecpp_DIR ${SYCL_COMPILER_DIR}/lib/cmake/adaptivecpp) - - if (NOT EXISTS "${AdaptiveCpp_DIR}") - message(WARNING "Falling back to AdaptiveCpp < 0.9.0 CMake structure") - set(AdaptiveCpp_DIR ${SYCL_COMPILER_DIR}/lib/cmake) - endif () - if (NOT EXISTS "${AdaptiveCpp_DIR}") - message(FATAL_ERROR "Can't find the appropriate CMake definitions for AdaptiveCpp") - endif () - - # register_definitions(_GLIBCXX_USE_CXX11_ABI=0) - find_package(AdaptiveCpp CONFIG REQUIRED) - message(STATUS "ok") - elseif (${SYCL_COMPILER} STREQUAL "HIPSYCL") - set(hipSYCL_DIR ${SYCL_COMPILER_DIR}/lib/cmake/hipSYCL) - - if (NOT EXISTS "${hipSYCL_DIR}") - message(WARNING "Falling back to hipSYCL < 0.9.0 CMake structure") - set(hipSYCL_DIR ${SYCL_COMPILER_DIR}/lib/cmake) - endif () - if (NOT EXISTS "${hipSYCL_DIR}") - message(FATAL_ERROR "Can't find the appropriate CMake definitions for hipSYCL") - endif () - - # register_definitions(_GLIBCXX_USE_CXX11_ABI=0) - find_package(hipSYCL CONFIG REQUIRED) - message(STATUS "ok") - elseif (${SYCL_COMPILER} STREQUAL "DPCPP") - set(CMAKE_CXX_COMPILER ${SYCL_COMPILER_DIR}/bin/clang++) - include_directories(${SYCL_COMPILER_DIR}/include/sycl) - register_append_cxx_flags(ANY -fsycl) - register_append_link_flags(-fsycl) - elseif (${SYCL_COMPILER} STREQUAL "ONEAPI-ICPX") - set(CMAKE_CXX_COMPILER icpx) - set(CMAKE_C_COMPILER icx) - register_append_cxx_flags(ANY -fsycl) - register_append_link_flags(-fsycl) - elseif (${SYCL_COMPILER} STREQUAL "ONEAPI-Clang") - set(CMAKE_CXX_COMPILER clang++) - set(CMAKE_C_COMPILER clang) - register_append_cxx_flags(ANY -fsycl) - register_append_link_flags(-fsycl) - else () - message(FATAL_ERROR "SYCL_COMPILER=${SYCL_COMPILER} is unsupported") - endif () - -endmacro() - - -macro(setup_target NAME) - if (${SYCL_COMPILER} STREQUAL "AdaptiveCpp") - # so AdaptiveCpp has this weird (and bad) CMake usage where they append their - # own custom integration header flags AFTER the target has been specified - # hence this macro here - add_sycl_to_target( - TARGET ${NAME} - SOURCES ${IMPL_SOURCES}) - elseif (${SYCL_COMPILER} STREQUAL "HIPSYCL") - # so hipSYCL has this weird (and bad) CMake usage where they append their - # own custom integration header flags AFTER the target has been specified - # hence this macro here - add_sycl_to_target( - TARGET ${NAME} - SOURCES ${IMPL_SOURCES}) - endif () -endmacro()