Skip to content

Commit 3f06cad

Browse files
premanandraopremanandrao
authored
[SYCL] Increase max _Bitint size in fpgs to 4096 (#6376)
The FPGA backend is able to support (and needs) a max _Bitint size of 4096. This change raises it from the previous max of 2048 and adjusts tests accordingly.
1 parent 2b8c118 commit 3f06cad

File tree

6 files changed

+52
-51
lines changed

6 files changed

+52
-51
lines changed

‎clang/lib/Basic/Targets/SPIR.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -230,7 +230,7 @@ class LLVM_LIBRARY_VISIBILITY SPIR64FPGATargetInfo : public SPIR64TargetInfo {
230230
public:
231231
SPIR64FPGATargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
232232
: SPIR64TargetInfo(Triple, Opts) {}
233-
virtual size_t getMaxBitIntWidth() const override { return 2048; }
233+
virtual size_t getMaxBitIntWidth() const override { return 4096; }
234234
};
235235

236236
// x86-32 SPIR Windows target

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

Lines changed: 15 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -3,23 +3,23 @@
33
// This test checks that we generate appropriate code for division
44
// operations of _BitInts of size greater than 128 bits, since it
55
// is allowed when -fintelfpga is enabled. The test uses a value of
6-
// 2048 for the bitsize as that is the maximum that is currently
6+
// 4096 for the bitsize as that is the maximum that is currently
77
// supported.
88

9-
// CHECK: define{{.*}} void @_Z3fooDB2048_S_(i2048* {{.*}} sret(i2048) align 8 %agg.result, i2048* {{.*}} byval(i2048) align 8 %[[ARG1:[0-9]+]], i2048* {{.*}} byval(i2048) align 8 %[[ARG2:[0-9]+]])
10-
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
11-
// CHECK: %[[VAR_A:a]].addr = alloca i2048, align 8
12-
// CHECK: %[[VAR_B:b]].addr = alloca i2048, align 8
13-
// CHECK: %[[VAR_A]] = load i2048, i2048* %[[ARG1]], align 8
14-
// CHECK: %[[VAR_B]] = load i2048, i2048* %[[ARG2]], align 8
15-
// CHECK: store i2048 %[[VAR_A]], i2048* %[[VAR_A]].addr, align 8
16-
// CHECK: store i2048 %[[VAR_B]], i2048* %[[VAR_B]].addr, align 8
17-
// CHECK: %[[TEMP1:[0-9]+]] = load i2048, i2048* %[[VAR_A]].addr, align 8
18-
// CHECK: %[[TEMP2:[0-9]+]] = load i2048, i2048* %[[VAR_B]].addr, align 8
19-
// CHECK: %div = sdiv i2048 %[[TEMP1]], %[[TEMP2]]
20-
// CHECK: store i2048 %div, i2048* %agg.result, align 8
21-
// CHECK: %[[RES:[0-9+]]] = load i2048, i2048* %agg.result, align 8
22-
// CHECK: store i2048 %[[RES]], i2048* %agg.result, align 8
9+
// CHECK: define{{.*}} void @_Z3fooDB4096_S_(i4096* {{.*}} sret(i4096) align 8 %agg.result, i4096* {{.*}} byval(i4096) align 8 %[[ARG1:[0-9]+]], i4096* {{.*}} byval(i4096) align 8 %[[ARG2:[0-9]+]])
10+
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
11+
// CHECK: %[[VAR_A:a]].addr = alloca i4096, align 8
12+
// CHECK: %[[VAR_B:b]].addr = alloca i4096, align 8
13+
// CHECK: %[[VAR_A]] = load i4096, i4096* %[[ARG1]], align 8
14+
// CHECK: %[[VAR_B]] = load i4096, i4096* %[[ARG2]], align 8
15+
// CHECK: store i4096 %[[VAR_A]], i4096* %[[VAR_A]].addr, align 8
16+
// CHECK: store i4096 %[[VAR_B]], i4096* %[[VAR_B]].addr, align 8
17+
// CHECK: %[[TEMP1:[0-9]+]] = load i4096, i4096* %[[VAR_A]].addr, align 8
18+
// CHECK: %[[TEMP2:[0-9]+]] = load i4096, i4096* %[[VAR_B]].addr, align 8
19+
// CHECK: %div = sdiv i4096 %[[TEMP1]], %[[TEMP2]]
20+
// CHECK: store i4096 %div, i4096* %agg.result, align 8
21+
// CHECK: %[[RES:[0-9+]]] = load i4096, i4096* %agg.result, align 8
22+
// CHECK: store i4096 %[[RES]], i4096* %agg.result, align 8
2323
// CHECK: ret void
2424
return a / b;
2525
}

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

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -3,24 +3,24 @@
33
// This test checks that we generate appropriate code for division
44
// operations of _BitInts of size greater than 128 bits, since it
55
// is allowed when -fintelfpga is enabled. The test uses a value
6-
// of 2048 for the bitsize, the max that is currently supported.
6+
// of 4096 for the bitsize, the max that is currently supported.
77

88
#include "Inputs/sycl.hpp"
99

10-
// CHECK: define{{.*}} void @_Z3fooDB2048_S_(i2048 addrspace(4)* {{.*}} sret(i2048) align 8 %agg.result, i2048* {{.*}} byval(i2048) align 8 %[[ARG1:[0-9]+]], i2048* {{.*}} byval(i2048) align 8 %[[ARG2:[0-9]+]])
11-
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
12-
// CHECK: %[[VAR_A:a]] = load i2048, i2048* %[[ARG1]], align 8
13-
// CHECK: %[[VAR_B:b]] = load i2048, i2048* %[[ARG2]], align 8
14-
// CHECK: %[[RES:div]] = sdiv i2048 %[[VAR_A]], %[[VAR_B]]
15-
// CHECK: store i2048 %[[RES]], i2048 addrspace(4)* %agg.result, align 8
10+
// CHECK: define{{.*}} void @_Z3fooDB4096_S_(i4096 addrspace(4)* {{.*}} sret(i4096) align 8 %agg.result, i4096* {{.*}} byval(i4096) align 8 %[[ARG1:[0-9]+]], i4096* {{.*}} byval(i4096) align 8 %[[ARG2:[0-9]+]])
11+
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
12+
// CHECK: %[[VAR_A:a]] = load i4096, i4096* %[[ARG1]], align 8
13+
// CHECK: %[[VAR_B:b]] = load i4096, i4096* %[[ARG2]], align 8
14+
// CHECK: %[[RES:div]] = sdiv i4096 %[[VAR_A]], %[[VAR_B]]
15+
// CHECK: store i4096 %[[RES]], i4096 addrspace(4)* %agg.result, align 8
1616
// CHECK: ret void
1717
return a / b;
1818
}
1919

2020
int main() {
2121
sycl::handler h;
2222
auto lambda = []() {
23-
_BitInt(2048) a, b = 3, c = 4;
23+
_BitInt(4096) a, b = 3, c = 4;
2424
a = foo(b, c);
2525
};
2626
h.single_task(lambda);

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

Lines changed: 15 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -3,22 +3,22 @@
33
// This test checks that we generate appropriate code for division
44
// operations of _BitInts of size greater than 128 bits, since it
55
// is allowed when -fintelfpga is enabled. The test uses a value of
6-
// 2048, the maximum bitsize that is currently supported.
6+
// 4096, the maximum bitsize that is currently supported.
77

8-
// CHECK: define{{.*}} void @_Z3fooDB2048_S_(ptr {{.*}} sret(i2048) align 8 %agg.result, ptr {{.*}} byval(i2048) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i2048) align 8 %[[ARG2:[0-9]+]])
9-
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
10-
// CHECK: %[[VAR_A:a]].addr = alloca i2048, align 8
11-
// CHECK: %[[VAR_B:b]].addr = alloca i2048, align 8
12-
// CHECK: %[[VAR_A]] = load i2048, ptr %[[ARG1]], align 8
13-
// CHECK: %[[VAR_B]] = load i2048, ptr %[[ARG2]], align 8
14-
// CHECK: store i2048 %[[VAR_A]], ptr %[[VAR_A]].addr, align 8
15-
// CHECK: store i2048 %[[VAR_B]], ptr %[[VAR_B]].addr, align 8
16-
// CHECK: %[[TEMP1:[0-9]+]] = load i2048, ptr %[[VAR_A]].addr, align 8
17-
// CHECK: %[[TEMP2:[0-9]+]] = load i2048, ptr %[[VAR_B]].addr, align 8
18-
// CHECK: %div = sdiv i2048 %[[TEMP1]], %[[TEMP2]]
19-
// CHECK: store i2048 %div, ptr %agg.result, align 8
20-
// CHECK: %[[RES:[0-9+]]] = load i2048, ptr %agg.result, align 8
21-
// CHECK: store i2048 %[[RES]], ptr %agg.result, align 8
8+
// CHECK: define{{.*}} void @_Z3fooDB4096_S_(ptr {{.*}} sret(i4096) align 8 %agg.result, ptr {{.*}} byval(i4096) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i4096) align 8 %[[ARG2:[0-9]+]])
9+
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
10+
// CHECK: %[[VAR_A:a]].addr = alloca i4096, align 8
11+
// CHECK: %[[VAR_B:b]].addr = alloca i4096, align 8
12+
// CHECK: %[[VAR_A]] = load i4096, ptr %[[ARG1]], align 8
13+
// CHECK: %[[VAR_B]] = load i4096, ptr %[[ARG2]], align 8
14+
// CHECK: store i4096 %[[VAR_A]], ptr %[[VAR_A]].addr, align 8
15+
// CHECK: store i4096 %[[VAR_B]], ptr %[[VAR_B]].addr, align 8
16+
// CHECK: %[[TEMP1:[0-9]+]] = load i4096, ptr %[[VAR_A]].addr, align 8
17+
// CHECK: %[[TEMP2:[0-9]+]] = load i4096, ptr %[[VAR_B]].addr, align 8
18+
// CHECK: %div = sdiv i4096 %[[TEMP1]], %[[TEMP2]]
19+
// CHECK: store i4096 %div, ptr %agg.result, align 8
20+
// CHECK: %[[RES:[0-9+]]] = load i4096, ptr %agg.result, align 8
21+
// CHECK: store i4096 %[[RES]], ptr %agg.result, align 8
2222
// CHECK: ret void
2323
return a / b;
2424
}

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

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -3,24 +3,24 @@
33
// This test checks that we generate appropriate code for division
44
// operations of _BitInts of size greater than 128 bits, since it
55
// is allowed when -fintelfpga is enabled. The test uses a value of
6-
// 2048 for bitint size, the maximum that is currently supported.
6+
// 4096 for bitint size, the maximum that is currently supported.
77

88
#include "Inputs/sycl.hpp"
99

10-
// CHECK: define{{.*}} void @_Z3fooDB2048_S_(ptr addrspace(4) {{.*}} sret(i2048) align 8 %agg.result, ptr {{.*}} byval(i2048) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i2048) align 8 %[[ARG2:[0-9]+]])
11-
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
12-
// CHECK: %[[VAR_A:a]] = load i2048, ptr %[[ARG1]], align 8
13-
// CHECK: %[[VAR_B:b]] = load i2048, ptr %[[ARG2]], align 8
14-
// CHECK: %[[RES:div]] = sdiv i2048 %[[VAR_A]], %[[VAR_B]]
15-
// CHECK: store i2048 %[[RES]], ptr addrspace(4) %agg.result, align 8
10+
// CHECK: define{{.*}} void @_Z3fooDB4096_S_(ptr addrspace(4) {{.*}} sret(i4096) align 8 %agg.result, ptr {{.*}} byval(i4096) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i4096) align 8 %[[ARG2:[0-9]+]])
11+
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
12+
// CHECK: %[[VAR_A:a]] = load i4096, ptr %[[ARG1]], align 8
13+
// CHECK: %[[VAR_B:b]] = load i4096, ptr %[[ARG2]], align 8
14+
// CHECK: %[[RES:div]] = sdiv i4096 %[[VAR_A]], %[[VAR_B]]
15+
// CHECK: store i4096 %[[RES]], ptr addrspace(4) %agg.result, align 8
1616
// CHECK: ret void
1717
return a / b;
1818
}
1919

2020
int main() {
2121
sycl::handler h;
2222
auto lambda = []() {
23-
_BitInt(2048) a, b = 3, c = 4;
23+
_BitInt(4096) a, b = 3, c = 4;
2424
a = foo(b, c);
2525
};
2626
h.single_task(lambda);

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

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5,13 +5,14 @@
55

66
// Tests that we do not issue errors for _Bitints of size greater than 128
77
// when -fintelfpga is enabled. The backend is expected to be able to handle
8-
// this. When -fintelfpga is not passed, we continue to diagnose.
8+
// this, upto a maximum size of 4096. When -fintelfpga is not passed,
9+
// we continue to diagnose size greater than 128.
910

10-
// device-intelfpga-error@+4 3{{signed _BitInt of bit sizes greater than 2048 not supported}}
11-
// host-intelfpga-error@+3 3{{signed _BitInt of bit sizes greater than 2048 not supported}}
11+
// device-intelfpga-error@+4 3{{signed _BitInt of bit sizes greater than 4096 not supported}}
12+
// host-intelfpga-error@+3 3{{signed _BitInt of bit sizes greater than 4096 not supported}}
1213
// device-error@+2 3{{signed _BitInt of bit sizes greater than 128 not supported}}
1314
// host-error@+1 3{{signed _BitInt of bit sizes greater than 128 not supported}}
14-
signed _BitInt(2049) foo(signed _BitInt(2049) a, signed _BitInt(2049) b) {
15+
signed _BitInt(4097) foo(signed _BitInt(4097) a, signed _BitInt(4097) b) {
1516
return a / b;
1617
}
1718
// device-error@+4 3{{signed _BitInt of bit sizes greater than 128 not supported}}

0 commit comments

Comments
 (0)