From 068788cad0ef1d82207d7d0c57168129970c0db8 Mon Sep 17 00:00:00 2001 From: "Kannan, Ajaykumar" Date: Tue, 13 Apr 2021 16:07:41 -0400 Subject: [PATCH 1/6] Change fpga_reg interface to only allow trivially copyable structs, and imply that a copy is created so that the x86 and the FPGA model are equivalent. We want to only allow trivially copyable structs, because we're creating an exact copy on device. --- sycl/include/CL/sycl/INTEL/fpga_reg.hpp | 17 +++++++---------- 1 file changed, 7 insertions(+), 10 deletions(-) diff --git a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp index 805c29501f84f..c9e27c99ddb4f 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp @@ -14,7 +14,12 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace INTEL { -template _T fpga_reg(const _T &t) { +// Returns a registered copy of the input +// This function is intended for FPGA users to instruct the compiler to insert +// at least one register stage between the input and the return value. +template +typename std::enable_if::value, T>::type +fpga_reg(T t) { #if __has_builtin(__builtin_intel_fpga_reg) return __builtin_intel_fpga_reg(t); #else @@ -24,12 +29,4 @@ template _T fpga_reg(const _T &t) { } // namespace INTEL } // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) - -// Keep it consistent with FPGA attributes like intelfpga::memory() -// Currently clang does not support nested namespace for attributes -namespace intelfpga { -template _T fpga_reg(const _T &t) { - return cl::sycl::INTEL::fpga_reg(t); -} -} // namespace intelfpga +} // __SYCL_INLINE_NAMESPACE(cl) \ No newline at end of file From 18696ab6f85b9a50d84dd2f0b6cf5ebf21b53efa Mon Sep 17 00:00:00 2001 From: "Kannan, Ajaykumar" Date: Wed, 21 Apr 2021 15:33:13 -0400 Subject: [PATCH 2/6] Add newline at EOF --- sycl/include/CL/sycl/INTEL/fpga_reg.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp index c9e27c99ddb4f..c617ae1f98c76 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp @@ -29,4 +29,5 @@ fpga_reg(T t) { } // namespace INTEL } // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) \ No newline at end of file +} // __SYCL_INLINE_NAMESPACE(cl) + From 1fa673b1d91458063cbf266070b0e638c351b76c Mon Sep 17 00:00:00 2001 From: "Kannan, Ajaykumar" Date: Wed, 21 Apr 2021 15:49:29 -0400 Subject: [PATCH 3/6] Fix a typo in fpga_reg.hpp --- sycl/include/CL/sycl/INTEL/fpga_reg.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp index c617ae1f98c76..38e588d5bfe4a 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp @@ -18,8 +18,8 @@ namespace INTEL { // This function is intended for FPGA users to instruct the compiler to insert // at least one register stage between the input and the return value. template -typename std::enable_if::value, T>::type -fpga_reg(T t) { +typename std::enable_if::value, _T>::type +fpga_reg(_T t) { #if __has_builtin(__builtin_intel_fpga_reg) return __builtin_intel_fpga_reg(t); #else From 37739acc98c3b08f3e9ae9d048f452fef1eb784d Mon Sep 17 00:00:00 2001 From: "Kannan, Ajaykumar" Date: Thu, 13 May 2021 00:35:01 -0400 Subject: [PATCH 4/6] Add type_traits as include to fpga_reg.hpp --- sycl/include/CL/sycl/INTEL/fpga_reg.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp index 38e588d5bfe4a..12ea5dd0b3ea3 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { From ccca3ba01e5adc48cb14569bfdea216dbe09bfe4 Mon Sep 17 00:00:00 2001 From: "Kannan, Ajaykumar" Date: Mon, 17 May 2021 19:21:29 -0400 Subject: [PATCH 5/6] Using a static assert in fpga_reg instead of a type_trait in the function signature --- sycl/include/CL/sycl/INTEL/fpga_reg.hpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp index 12ea5dd0b3ea3..4ccfc218ac6af 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp @@ -18,9 +18,9 @@ namespace INTEL { // Returns a registered copy of the input // This function is intended for FPGA users to instruct the compiler to insert // at least one register stage between the input and the return value. -template -typename std::enable_if::value, _T>::type -fpga_reg(_T t) { +template _T fpga_reg(_T t) { + static_assert(std::is_trivially_copyable<_T>::value, + "Type is not trivially_copyable."); #if __has_builtin(__builtin_intel_fpga_reg) return __builtin_intel_fpga_reg(t); #else @@ -31,4 +31,3 @@ fpga_reg(_T t) { } // namespace INTEL } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) - From 079f31ab803d81cd3a4113ab454bbb8bb5ffcf26 Mon Sep 17 00:00:00 2001 From: "Kannan, Ajaykumar" Date: Tue, 25 May 2021 10:55:38 -0400 Subject: [PATCH 6/6] Change fpga_reg to use a deprecation warning for old APIs --- sycl/include/CL/sycl/INTEL/fpga_reg.hpp | 29 ++++++++++++++++++++++--- 1 file changed, 26 insertions(+), 3 deletions(-) diff --git a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp index 4ccfc218ac6af..6142d0cd1996d 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp @@ -18,9 +18,22 @@ namespace INTEL { // Returns a registered copy of the input // This function is intended for FPGA users to instruct the compiler to insert // at least one register stage between the input and the return value. -template _T fpga_reg(_T t) { - static_assert(std::is_trivially_copyable<_T>::value, - "Type is not trivially_copyable."); +template +typename std::enable_if::value, _T>::type +fpga_reg(_T t) { +#if __has_builtin(__builtin_intel_fpga_reg) + return __builtin_intel_fpga_reg(t); +#else + return t; +#endif +} + +template +[[deprecated("INTEL::fpga_reg will only support trivially_copyable types in a " + "future release. The type used here will be disallowed.")]] +typename std::enable_if::value == false, + _T>::type +fpga_reg(_T t) { #if __has_builtin(__builtin_intel_fpga_reg) return __builtin_intel_fpga_reg(t); #else @@ -31,3 +44,13 @@ template _T fpga_reg(_T t) { } // namespace INTEL } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) + +// Keep it consistent with FPGA attributes like intelfpga::memory() +// Currently clang does not support nested namespace for attributes +namespace intelfpga { +template +[[deprecated("intelfpga::fpga_reg will be removed in a future release.")]] _T +fpga_reg(const _T &t) { + return cl::sycl::INTEL::fpga_reg(t); +} +} // namespace intelfpga