Skip to content

Commit e8ac5a0

Browse files
authored
[SYCL][FPGA] Adding function wrapper for function-level loop fusion attributes (#4716)
The function attributes [[intel::loop_fuse(N)]] and [[intel::loop_fuse_independent(N)]] are being replaced by equivalent function wrappers by the inclusion of a header file. These attributes can be accessed now by sycl::ext::intel::fpga_loop_fuse<N>(F) and sycl::ext::intel::fpga_loop_fuse_independent<N>(F) for a function or lambda F. There is a [lit test](https://github.com/intel/llvm/blob/e9b707dd1653ed74c639b987481bbfe151be2cbd/llvm-spirv/test/transcoding/SPV_INTEL_kernel_attributes/intel_fpga_function_attributes.ll) for the functionality of the loop fusion attributes. The attributes are tested on functions elsewhere (e.g. [here ](https://github.com/intel/llvm/blob/29a1f167060e69d4d90d1e6a15d4998fd4556c4d/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp) and [here](https://github.com/intel/llvm/blob/29a1f167060e69d4d90d1e6a15d4998fd4556c4d/clang/test/CodeGenSYCL/loop_fuse_device.cpp)).
1 parent 97d33b7 commit e8ac5a0

File tree

2 files changed

+29
-0
lines changed

2 files changed

+29
-0
lines changed

sycl/include/sycl/ext/intel/fpga_extensions.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,7 @@
88

99
#pragma once
1010
#include <sycl/ext/intel/fpga_device_selector.hpp>
11+
#include <sycl/ext/intel/fpga_loop_fuse.hpp>
1112
#include <sycl/ext/intel/fpga_lsu.hpp>
1213
#include <sycl/ext/intel/fpga_reg.hpp>
1314
#include <sycl/ext/intel/pipes.hpp>
Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,28 @@
1+
//==--------- fpga_loop_fuse.hpp --- SYCL FPGA Loop Fuse Extension ---------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
#pragma once
9+
10+
__SYCL_INLINE_NAMESPACE(cl) {
11+
namespace sycl {
12+
namespace ext {
13+
namespace intel {
14+
15+
template <int _N = 1, typename _F>
16+
void fpga_loop_fuse [[intel::loop_fuse(_N)]] (_F f) {
17+
f();
18+
}
19+
20+
template <int _N = 1, typename _F>
21+
void fpga_loop_fuse_independent [[intel::loop_fuse_independent(_N)]] (_F f) {
22+
f();
23+
}
24+
25+
} // namespace intel
26+
} // namespace ext
27+
} // namespace sycl
28+
} // __SYCL_INLINE_NAMESPACE(cl)

0 commit comments

Comments
 (0)