[SYCL][FPGA] Add compute units library. #2395

jessicadavies-intel · 2020-08-31T21:17:13Z

Add a header library that allows the user to duplicate a single-task kernel k times, for any k>=1.

For spatial architectures such as FPGA, users often wish to duplicate their kernels several times in order to fill the available area on the chip. Code duplication is easy to achieve using C++ template metaprogramming. We designed this API to improve the usability.

Usage:

// The name of the kernel is templated on the compute unit ID,
// so each compute unit receives a unique name, e.g., foo_kernel<0>
template <std::size_t ID> class foo_kernel;

// Create 5 compute units whose functionality is determined by the generic lambda
intelfpga::submit_compute_units<5, foo_kernel>(q, [=](auto ID) {
   // kernel function body
  // compute unit functionality can be specialized on ID
});

Add a header library that allows the user to duplicate a single-task kernel k times, for any k>=1.

MrSidims

How is it expected to work with pipes? Will a programmer read/writer from different pipes within these kernels (aka by using templates in pipes naming) or we expect some sort of serialization in the runtime?

jessicadavies-intel · 2020-09-01T13:54:53Z

How is it expected to work with pipes? Will a programmer read/writer from different pipes within these kernels (aka by using templates in pipes naming) or we expect some sort of serialization in the runtime?

We have a working example that uses pipes. Yes, the intended use is for each compute unit to read/write to particular pipes, depending on the compute unit's ID.

s-kanaev · 2020-09-01T14:27:12Z

sycl/include/CL/sycl/INTEL/compute_units.hpp

+#include <CL/sycl.hpp>
+#include <utility>
+
+namespace intelfpga {


Is it a proper namespace? I mean, it's only ::intelfpga, not smth like cl::sycl::intelfpga. Is it OK?

No, it's not OK. It shall be at leastcl::sycl::INTEL. And I'm not sure about intelfpga.

Thanks. I think it should be cl::sycl::INTEL.

jessicadavies-intel · 2020-09-01T18:10:15Z

@MrSidims @s-kanaev do you think this should go in ONEAPI instead of INTEL? It isn't really specific to Intel hardware, so maybe it belongs in ONEAPI instead?

MrSidims · 2020-09-01T18:35:56Z

@MrSidims @s-kanaev do you think this should go in ONEAPI instead of INTEL? It isn't really specific to Intel hardware, so maybe it belongs in ONEAPI instead?

ONEAPI namespace serves for DPCPP language extensions (in theory anyone outside of Intel can add an extension here), while INTEL namespace is for Intel hardware extensions, each extension is for the specific hardware (f. e. pipes are for FPGA and not used anywhere else and ESIMD is for GEN and not used anywhere else). So I'd say this extension belongs to INTEL namespace as it is FPGA H/W extension.

Some comments for the patch. Functional changes are LGTM, but please add a test. Please also consider to add the extension specification. I think, it can be done separately and after the initial commit (yet it's just INTEL extension, not ONEAPI, where the rules are more strict).

Add a test for the compute units header. Also add the pipe array header, since both headers can be easily tested at the same time.

MrSidims · 2020-09-01T20:06:58Z

sycl/test/fpga_tests/compute_units.cpp

@@ -0,0 +1,78 @@
+//==------------ compute_units.cpp - SYCL FPGA compute units test ----------==//


You need to add compile and run string for the test. You can take one from pipes.cpp

MrSidims · 2020-09-01T20:07:12Z

sycl/include/CL/sycl/INTEL/pipe_array.hpp

+
+#pragma once
+
+#include <CL/sycl/intel/pipes.hpp>


Suggested change

#include <CL/sycl/intel/pipes.hpp>

#include <CL/sycl/INTEL/pipes.hpp>

keryell

Thank you for not doing this with #pragma, devconstexpr or [[attribute]]! :-)

keryell · 2020-09-01T23:12:10Z

sycl/include/CL/sycl/INTEL/compute_units.hpp

+  SubmitOneComputeUnit(Func &&f, sycl::queue &q) {
+
+    q.submit([&](sycl::handler &h) {
+      h.single_task<Name<index>>([=]() {


Suggested change

h.single_task<Name<index>>([=]() {

h.single_task<Name<index>>([=] {

keryell · 2020-09-01T23:16:00Z

sycl/test/fpga_tests/compute_units.cpp

+void SourceKernel(queue &q, float data) {
+
+  q.submit([&](handler &h) {
+    h.single_task<source_kernel>([=]() { Pipes::PipeAt<0>::write(data); });


Suggested change

h.single_task<source_kernel>([=]() { Pipes::PipeAt<0>::write(data); });

h.single_task<source_kernel>([=] { Pipes::PipeAt<0>::write(data); });

Otherwise, I was thinking to this alternative API:

h.single_task<source_kernel>([=] { Pipes::write<0>(data); });

keryell · 2020-09-01T23:20:14Z

sycl/include/CL/sycl/INTEL/pipe_array.hpp

+} // namespace
+
+template <class Id, typename BaseTy, size_t depth, size_t... dims>
+struct PipeArray {


This is not really an array... But I cannot find a better name.
pipe_array looks more similar to SYCL coding style, otherwise.

keryell · 2020-09-01T23:22:36Z

sycl/include/CL/sycl/INTEL/pipe_array.hpp

+  };
+
+  template <size_t... idxs>
+  using PipeAt = typename VerifyIndices<idxs...>::VerifiedPipe;


Perhaps get to be closer to std::tuple and even std::array through the tuple interface?
Or just at?
I am not convinced you need to repeat pipe in the name.

keryell · 2020-09-01T23:23:04Z

sycl/test/fpga_tests/compute_units.cpp

+
+#include <CL/sycl/INTEL/fpga_extensions.hpp>
+
+using namespace sycl;


Bad coding style for a tutorial but probably OK in a unit test...

keryell · 2020-09-01T23:29:02Z

sycl/test/fpga_tests/compute_units.cpp

+  q.submit([&](handler &h) {
+    auto out_accessor = out_buf.get_access<access::mode::write>(h);
+    h.single_task<sink_kernel>(
+        [=]() { out_accessor[0] = Pipes::PipeAt<kComputeUnits>::read(); });


Suggested change

[=]() { out_accessor[0] = Pipes::PipeAt<kComputeUnits>::read(); });

[=] { out_accessor[0] = Pipes::PipeAt<kComputeUnits>::read(); });

Otherwise, I was thinking to this alternative API:

[=] { out_accessor[0] = Pipes::read<kComputeUnits>(); });

jessicadavies-intel · 2020-09-02T14:30:35Z

@keryell thank you for the review. I will use your suggestions to simplify the read/write API.

After talking to @jbrodman we are planning to move these headers to a separate repo. Therefore, I'm closing this PR.

In translation from __spirv_AtomicCompareExchange to OpenCL builtin atomic_compare_exchange_strong_explicit, a new alloca `expected` is created and read/written in the OpenCL builtin. The OpenCL builtin call can't have tail marker since the marker requires that callee doesn't access alloca from the caller. Otherwise llvm alias analysis deduces that the alloca isn't accessed by the call, and instcombine pass replaces the load from the alloca after the call with the value stored to the alloca before the call. Original commit: KhronosGroup/SPIRV-LLVM-Translator@1ff4a764cd0f97c

fix event caching

[SYCL][FPGA] Add compute units library.

4dc7cc2

Add a header library that allows the user to duplicate a single-task kernel k times, for any k>=1.

jessicadavies-intel requested a review from a team as a code owner August 31, 2020 21:17

jessicadavies-intel requested a review from s-kanaev August 31, 2020 21:17

MrSidims reviewed Aug 31, 2020

View reviewed changes

s-kanaev reviewed Sep 1, 2020

View reviewed changes

jessicadavies-intel added 2 commits September 1, 2020 13:38

Change namespace to cl::sycl::INTEL from intelfpga

e09f711

Run clang-format

1fe9887

jessicadavies-intel added 2 commits September 1, 2020 15:25

Add test for compute units header

04a9515

Add a test for the compute units header. Also add the pipe array header, since both headers can be easily tested at the same time.

Run clang-format to correct include order

e133934

MrSidims requested changes Sep 1, 2020

View reviewed changes

jessicadavies-intel added 2 commits September 1, 2020 16:09

Fix compute units test

3926899

More fixes caught by CI

96e7682

keryell reviewed Sep 1, 2020

View reviewed changes

jessicadavies-intel closed this Sep 2, 2020

keryell mentioned this pull request Sep 22, 2020

Add compute_units DPC++ FPGA tutorial oneapi-src/oneAPI-samples#167

Merged

2 tasks

martygrant added a commit to martygrant/llvm that referenced this pull request Dec 2, 2024

[UR] Update tag to 7a38cc3 for UR intel#2395.

50b7e3e

Chenyang-L pushed a commit that referenced this pull request Feb 18, 2025

Merge pull request #2395 from pbalcer/fix-event-pooling

db011cd

fix event caching

		@@ -0,0 +1,78 @@
		//==------------ compute_units.cpp - SYCL FPGA compute units test ----------==//

	#include <CL/sycl/intel/pipes.hpp>
	#include <CL/sycl/INTEL/pipes.hpp>

	h.single_task<Name<index>>([=]() {
	h.single_task<Name<index>>([=] {

	h.single_task<source_kernel>([=]() { Pipes::PipeAt<0>::write(data); });
	h.single_task<source_kernel>([=] { Pipes::PipeAt<0>::write(data); });


		#include <CL/sycl/INTEL/fpga_extensions.hpp>

		using namespace sycl;

	[=]() { out_accessor[0] = Pipes::PipeAt<kComputeUnits>::read(); });
	[=] { out_accessor[0] = Pipes::PipeAt<kComputeUnits>::read(); });

[SYCL][FPGA] Add compute units library. #2395

[SYCL][FPGA] Add compute units library. #2395

Uh oh!

Conversation

jessicadavies-intel commented Aug 31, 2020

Uh oh!

MrSidims left a comment

Choose a reason for hiding this comment

Uh oh!

jessicadavies-intel commented Sep 1, 2020

Uh oh!

Choose a reason for hiding this comment

Uh oh!

MrSidims Sep 1, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jessicadavies-intel commented Sep 1, 2020

Uh oh!

MrSidims commented Sep 1, 2020

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

keryell left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jessicadavies-intel commented Sep 2, 2020

Uh oh!

Uh oh!

MrSidims Sep 1, 2020 •

edited

Loading