Skip to content

Commit 51f2ea3

Browse files
premanandraopremanandrao
authored
[SYCL] Allow _Bitint of size greater than 128 bits when -fintelfpga is used (#6152)
1 parent f799cf7 commit 51f2ea3

File tree

10 files changed

+128
-3
lines changed

10 files changed

+128
-3
lines changed

‎clang/include/clang/Basic/LangOptions.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -269,6 +269,7 @@ LANGOPT(GPUExcludeWrongSideOverloads, 1, 0, "always exclude wrong side overloads
269269

270270
LANGOPT(SYCLIsDevice , 1, 0, "Generate code for SYCL device")
271271
LANGOPT(SYCLIsHost , 1, 0, "SYCL host compilation")
272+
LANGOPT(IntelFPGA , 1, 0, "Perform ahead-of-time compilation for FPGA")
272273
LANGOPT(SYCLAllowFuncPtr , 1, 0, "Allow function pointers in SYCL device code")
273274
LANGOPT(SYCLStdLayoutKernelParams, 1, 0, "Enable standard layout requirement for SYCL kernel parameters")
274275
LANGOPT(SYCLUnnamedLambda , 1, 0, "Allow unnamed lambda SYCL kernels")

‎clang/include/clang/Driver/Options.td

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2754,7 +2754,8 @@ defm strict_vtable_pointers : BoolFOption<"strict-vtable-pointers",
27542754
NegFlag<SetFalse>>;
27552755
def fstrict_overflow : Flag<["-"], "fstrict-overflow">, Group<f_Group>;
27562756
def fintelfpga : Flag<["-"], "fintelfpga">, Group<f_Group>,
2757-
Flags<[CC1Option, CoreOption]>, HelpText<"Perform ahead of time compilation for FPGA">;
2757+
Flags<[CC1Option, CoreOption]>, MarshallingInfoFlag<LangOpts<"IntelFPGA">>,
2758+
HelpText<"Perform ahead-of-time compilation for FPGA">;
27582759
def fsycl_device_only : Flag<["-"], "fsycl-device-only">, Flags<[CoreOption]>,
27592760
HelpText<"Compile SYCL kernels for device">;
27602761
def fsycl_targets_EQ : CommaJoined<["-"], "fsycl-targets=">, Flags<[NoXarchOption, CC1Option, CoreOption]>,

‎clang/lib/Driver/ToolChains/Clang.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4965,8 +4965,11 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
49654965
HasFPGA = true;
49664966
break;
49674967
}
4968-
if (HasFPGA)
4968+
if (HasFPGA) {
49694969
CmdArgs.push_back("-fsycl-disable-range-rounding");
4970+
// Pass -fintelfpga to both the host and device SYCL compilations if set.
4971+
CmdArgs.push_back("-fintelfpga");
4972+
}
49704973

49714974
// Add any options that are needed specific to SYCL offload while
49724975
// performing the host side compilation.

‎clang/lib/Sema/SemaType.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2292,7 +2292,7 @@ QualType Sema::BuildBitIntType(bool IsUnsigned, Expr *BitWidth,
22922292
}
22932293

22942294
const TargetInfo &TI = getASTContext().getTargetInfo();
2295-
if (NumBits > TI.getMaxBitIntWidth()) {
2295+
if (NumBits > TI.getMaxBitIntWidth() && !Context.getLangOpts().IntelFPGA) {
22962296
Diag(Loc, diag::err_bit_int_max_size)
22972297
<< IsUnsigned << static_cast<uint64_t>(TI.getMaxBitIntWidth());
22982298
return QualType();

‎clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-host-intelfpga-bitint.cpp

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
// RUN: %clang_cc1 -no-opaque-pointers -fsycl-is-host -fintelfpga -triple x86_64 -emit-llvm %s -o - | FileCheck %s
2+
3+
// This test checks that we generate appropriate code for division
4+
// operations of _BitInts of size greater than 128 bits, since it
5+
// is allowed when -fintelfpga is enabled.
6+
7+
// CHECK: define{{.*}} void @_Z3fooDB211_S_(i211* {{.*}} sret(i211) align 8 %agg.result, i211* {{.*}} byval(i211) align 8 %[[ARG1:[0-9]+]], i211* {{.*}} byval(i211) align 8 %[[ARG2:[0-9]+]])
8+
signed _BitInt(211) foo(signed _BitInt(211) a, signed _BitInt(211) b) {
9+
// CHECK: %[[VAR_A:a]].addr = alloca i211, align 8
10+
// CHECK: %[[VAR_B:b]].addr = alloca i211, align 8
11+
// CHECK: %[[VAR_A]] = load i211, i211* %[[ARG1]], align 8
12+
// CHECK: %[[VAR_B]] = load i211, i211* %[[ARG2]], align 8
13+
// CHECK: store i211 %[[VAR_A]], i211* %[[VAR_A]].addr, align 8
14+
// CHECK: store i211 %[[VAR_B]], i211* %[[VAR_B]].addr, align 8
15+
// CHECK: %[[TEMP1:[0-9]+]] = load i211, i211* %[[VAR_A]].addr, align 8
16+
// CHECK: %[[TEMP2:[0-9]+]] = load i211, i211* %[[VAR_B]].addr, align 8
17+
// CHECK: %div = sdiv i211 %[[TEMP1]], %[[TEMP2]]
18+
// CHECK: store i211 %div, i211* %agg.result, align 8
19+
// CHECK: %[[RES:[0-9+]]] = load i211, i211* %agg.result, align 8
20+
// CHECK: store i211 %[[RES]], i211* %agg.result, align 8
21+
// CHECK: ret void
22+
return a / b;
23+
}

‎clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-intelfpga-bitint.cpp

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
// RUN: %clang_cc1 -no-opaque-pointers -fsycl-is-device -fintelfpga -triple spir64_fpga -IInputs -emit-llvm %s -o - | FileCheck %s
2+
3+
// This test checks that we generate appropriate code for division
4+
// operations of _BitInts of size greater than 128 bits, since it
5+
// is allowed when -fintelfpga is enabled.
6+
7+
#include "Inputs/sycl.hpp"
8+
9+
// CHECK: define{{.*}} void @_Z3fooDB211_S_(i211 addrspace(4)* {{.*}} sret(i211) align 8 %agg.result, i211* {{.*}} byval(i211) align 8 %[[ARG1:[0-9]+]], i211* {{.*}} byval(i211) align 8 %[[ARG2:[0-9]+]])
10+
signed _BitInt(211) foo(signed _BitInt(211) a, signed _BitInt(211) b) {
11+
// CHECK: %[[VAR_A:a]] = load i211, i211* %[[ARG1]], align 8
12+
// CHECK: %[[VAR_B:b]] = load i211, i211* %[[ARG2]], align 8
13+
// CHECK: %[[RES:div]] = sdiv i211 %[[VAR_A]], %[[VAR_B]]
14+
// CHECK: store i211 %[[RES]], i211 addrspace(4)* %agg.result, align 8
15+
// CHECK: ret void
16+
return a / b;
17+
}
18+
19+
int main() {
20+
sycl::handler h;
21+
auto lambda = []() {
22+
_BitInt(211) a, b = 3, c = 4;
23+
a = foo(b, c);
24+
};
25+
h.single_task(lambda);
26+
}

‎clang/test/CodeGenSYCL/sycl-host-intelfpga-bitint.cpp

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
// RUN: %clang_cc1 -opaque-pointers -fsycl-is-host -fintelfpga -triple x86_64 -emit-llvm %s -o - | FileCheck %s
2+
3+
// This test checks that we generate appropriate code for division
4+
// operations of _BitInts of size greater than 128 bits, since it
5+
// is allowed when -fintelfpga is enabled.
6+
7+
// CHECK: define{{.*}} void @_Z3fooDB211_S_(ptr {{.*}} sret(i211) align 8 %agg.result, ptr {{.*}} byval(i211) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i211) align 8 %[[ARG2:[0-9]+]])
8+
signed _BitInt(211) foo(signed _BitInt(211) a, signed _BitInt(211) b) {
9+
// CHECK: %[[VAR_A:a]].addr = alloca i211, align 8
10+
// CHECK: %[[VAR_B:b]].addr = alloca i211, align 8
11+
// CHECK: %[[VAR_A]] = load i211, ptr %[[ARG1]], align 8
12+
// CHECK: %[[VAR_B]] = load i211, ptr %[[ARG2]], align 8
13+
// CHECK: store i211 %[[VAR_A]], ptr %[[VAR_A]].addr, align 8
14+
// CHECK: store i211 %[[VAR_B]], ptr %[[VAR_B]].addr, align 8
15+
// CHECK: %[[TEMP1:[0-9]+]] = load i211, ptr %[[VAR_A]].addr, align 8
16+
// CHECK: %[[TEMP2:[0-9]+]] = load i211, ptr %[[VAR_B]].addr, align 8
17+
// CHECK: %div = sdiv i211 %[[TEMP1]], %[[TEMP2]]
18+
// CHECK: store i211 %div, ptr %agg.result, align 8
19+
// CHECK: %[[RES:[0-9+]]] = load i211, ptr %agg.result, align 8
20+
// CHECK: store i211 %[[RES]], ptr %agg.result, align 8
21+
// CHECK: ret void
22+
return a / b;
23+
}

‎clang/test/CodeGenSYCL/sycl-intelfpga-bitint.cpp

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
// RUN: %clang_cc1 -opaque-pointers -fsycl-is-device -fintelfpga -triple spir64_fpga -IInputs -emit-llvm %s -o - | FileCheck %s
2+
3+
// This test checks that we generate appropriate code for division
4+
// operations of _BitInts of size greater than 128 bits, since it
5+
// is allowed when -fintelfpga is enabled.
6+
7+
#include "Inputs/sycl.hpp"
8+
9+
// CHECK: define{{.*}} void @_Z3fooDB211_S_(ptr addrspace(4) {{.*}} sret(i211) align 8 %agg.result, ptr {{.*}} byval(i211) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i211) align 8 %[[ARG2:[0-9]+]])
10+
signed _BitInt(211) foo(signed _BitInt(211) a, signed _BitInt(211) b) {
11+
// CHECK: %[[VAR_A:a]] = load i211, ptr %[[ARG1]], align 8
12+
// CHECK: %[[VAR_B:b]] = load i211, ptr %[[ARG2]], align 8
13+
// CHECK: %[[RES:div]] = sdiv i211 %[[VAR_A]], %[[VAR_B]]
14+
// CHECK: store i211 %[[RES]], ptr addrspace(4) %agg.result, align 8
15+
// CHECK: ret void
16+
return a / b;
17+
}
18+
19+
int main() {
20+
sycl::handler h;
21+
auto lambda = []() {
22+
_BitInt(211) a, b = 3, c = 4;
23+
a = foo(b, c);
24+
};
25+
h.single_task(lambda);
26+
}

‎clang/test/Driver/sycl-offload-intelfpga.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,12 @@
2020
// RUN: | FileCheck -check-prefix=CHK-TOOLS-INTELFPGA-G0 %s
2121
// CHK-TOOLS-INTELFPGA-G0-NOT: clang{{.*}} "-debug-info-kind=constructor"
2222

23+
/// -fintelfpga passes it to host and device cc1 compilations
24+
// RUN: %clangxx -### -fsycl -fintelfpga %s 2>&1 \
25+
// RUN: | FileCheck -check-prefix=CHK-HOST-DEVICE %s
26+
// CHK-HOST-DEVICE: clang{{.*}} "-cc1"{{.*}} "-fsycl-is-device"{{.*}} "-fintelfpga"
27+
// CHK-HOST-DEVICE: clang{{.*}} "-cc1"{{.*}} "-fintelfpga"{{.*}} "-fsycl-is-host"
28+
2329
/// FPGA target implies -fsycl-disable-range-rounding
2430
// RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -fintelfpga %s 2>&1 \
2531
// RUN: | FileCheck -check-prefix=CHK-RANGE-ROUNDING %s

‎clang/test/SemaSYCL/sycl-intelfpga.cpp

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,16 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -fintelfpga -verify=device-intelfpga -fsyntax-only %s
2+
// RUN: %clang_cc1 -fsycl-is-host -fintelfpga -verify=host-intelfpga -fsyntax-only %s
3+
// RUN: %clang_cc1 -fsycl-is-device -verify=device -fsyntax-only %s
4+
// RUN: %clang_cc1 -fsycl-is-host -verify=host -fsyntax-only %s
5+
6+
// Tests that we do not issue errors for _Bitints of size greater than 128
7+
// when -fintelfpga is enabled. The backend is expected to be able to handle
8+
// this. When -fintelfpga is not passed, we continue to diagnose.
9+
10+
// device-intelfpga-no-diagnostics
11+
// host-intelfpga-no-diagnostics
12+
// device-error@+2 3{{signed _BitInt of bit sizes greater than 128 not supported}}
13+
// host-error@+1 3{{signed _BitInt of bit sizes greater than 128 not supported}}
14+
signed _BitInt(211) foo(signed _BitInt(211) a, signed _BitInt(211) b) {
15+
return a / b;
16+
}

0 commit comments

Comments
 (0)