Skip to content

[SYCL] Increase max _BitInt size in FPGA to 4096 #6376

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Jun 30, 2022
Merged

[SYCL] Increase max _BitInt size in FPGA to 4096 #6376

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion clang/lib/Basic/Targets/SPIR.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -230,7 +230,7 @@ class LLVM_LIBRARY_VISIBILITY SPIR64FPGATargetInfo : public SPIR64TargetInfo {
public:
SPIR64FPGATargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
: SPIR64TargetInfo(Triple, Opts) {}
virtual size_t getMaxBitIntWidth() const override { return 2048; }
virtual size_t getMaxBitIntWidth() const override { return 4096; }
};

// x86-32 SPIR Windows target
Expand Down
30 changes: 15 additions & 15 deletions clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-host-intelfpga-bitint.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,23 +3,23 @@
// This test checks that we generate appropriate code for division
// operations of _BitInts of size greater than 128 bits, since it
// is allowed when -fintelfpga is enabled. The test uses a value of
// 2048 for the bitsize as that is the maximum that is currently
// 4096 for the bitsize as that is the maximum that is currently
// supported.

// 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]+]])
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
// CHECK: %[[VAR_A:a]].addr = alloca i2048, align 8
// CHECK: %[[VAR_B:b]].addr = alloca i2048, align 8
// CHECK: %[[VAR_A]] = load i2048, i2048* %[[ARG1]], align 8
// CHECK: %[[VAR_B]] = load i2048, i2048* %[[ARG2]], align 8
// CHECK: store i2048 %[[VAR_A]], i2048* %[[VAR_A]].addr, align 8
// CHECK: store i2048 %[[VAR_B]], i2048* %[[VAR_B]].addr, align 8
// CHECK: %[[TEMP1:[0-9]+]] = load i2048, i2048* %[[VAR_A]].addr, align 8
// CHECK: %[[TEMP2:[0-9]+]] = load i2048, i2048* %[[VAR_B]].addr, align 8
// CHECK: %div = sdiv i2048 %[[TEMP1]], %[[TEMP2]]
// CHECK: store i2048 %div, i2048* %agg.result, align 8
// CHECK: %[[RES:[0-9+]]] = load i2048, i2048* %agg.result, align 8
// CHECK: store i2048 %[[RES]], i2048* %agg.result, align 8
// 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]+]])
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
// CHECK: %[[VAR_A:a]].addr = alloca i4096, align 8
// CHECK: %[[VAR_B:b]].addr = alloca i4096, align 8
// CHECK: %[[VAR_A]] = load i4096, i4096* %[[ARG1]], align 8
// CHECK: %[[VAR_B]] = load i4096, i4096* %[[ARG2]], align 8
// CHECK: store i4096 %[[VAR_A]], i4096* %[[VAR_A]].addr, align 8
// CHECK: store i4096 %[[VAR_B]], i4096* %[[VAR_B]].addr, align 8
// CHECK: %[[TEMP1:[0-9]+]] = load i4096, i4096* %[[VAR_A]].addr, align 8
// CHECK: %[[TEMP2:[0-9]+]] = load i4096, i4096* %[[VAR_B]].addr, align 8
// CHECK: %div = sdiv i4096 %[[TEMP1]], %[[TEMP2]]
// CHECK: store i4096 %div, i4096* %agg.result, align 8
// CHECK: %[[RES:[0-9+]]] = load i4096, i4096* %agg.result, align 8
// CHECK: store i4096 %[[RES]], i4096* %agg.result, align 8
// CHECK: ret void
return a / b;
}
16 changes: 8 additions & 8 deletions clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-intelfpga-bitint.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,24 +3,24 @@
// This test checks that we generate appropriate code for division
// operations of _BitInts of size greater than 128 bits, since it
// is allowed when -fintelfpga is enabled. The test uses a value
// of 2048 for the bitsize, the max that is currently supported.
// of 4096 for the bitsize, the max that is currently supported.

#include "Inputs/sycl.hpp"

// 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]+]])
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
// CHECK: %[[VAR_A:a]] = load i2048, i2048* %[[ARG1]], align 8
// CHECK: %[[VAR_B:b]] = load i2048, i2048* %[[ARG2]], align 8
// CHECK: %[[RES:div]] = sdiv i2048 %[[VAR_A]], %[[VAR_B]]
// CHECK: store i2048 %[[RES]], i2048 addrspace(4)* %agg.result, align 8
// 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]+]])
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
// CHECK: %[[VAR_A:a]] = load i4096, i4096* %[[ARG1]], align 8
// CHECK: %[[VAR_B:b]] = load i4096, i4096* %[[ARG2]], align 8
// CHECK: %[[RES:div]] = sdiv i4096 %[[VAR_A]], %[[VAR_B]]
// CHECK: store i4096 %[[RES]], i4096 addrspace(4)* %agg.result, align 8
// CHECK: ret void
return a / b;
}

int main() {
sycl::handler h;
auto lambda = []() {
_BitInt(2048) a, b = 3, c = 4;
_BitInt(4096) a, b = 3, c = 4;
a = foo(b, c);
};
h.single_task(lambda);
Expand Down
30 changes: 15 additions & 15 deletions clang/test/CodeGenSYCL/sycl-host-intelfpga-bitint.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,22 +3,22 @@
// This test checks that we generate appropriate code for division
// operations of _BitInts of size greater than 128 bits, since it
// is allowed when -fintelfpga is enabled. The test uses a value of
// 2048, the maximum bitsize that is currently supported.
// 4096, the maximum bitsize that is currently supported.

// 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]+]])
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
// CHECK: %[[VAR_A:a]].addr = alloca i2048, align 8
// CHECK: %[[VAR_B:b]].addr = alloca i2048, align 8
// CHECK: %[[VAR_A]] = load i2048, ptr %[[ARG1]], align 8
// CHECK: %[[VAR_B]] = load i2048, ptr %[[ARG2]], align 8
// CHECK: store i2048 %[[VAR_A]], ptr %[[VAR_A]].addr, align 8
// CHECK: store i2048 %[[VAR_B]], ptr %[[VAR_B]].addr, align 8
// CHECK: %[[TEMP1:[0-9]+]] = load i2048, ptr %[[VAR_A]].addr, align 8
// CHECK: %[[TEMP2:[0-9]+]] = load i2048, ptr %[[VAR_B]].addr, align 8
// CHECK: %div = sdiv i2048 %[[TEMP1]], %[[TEMP2]]
// CHECK: store i2048 %div, ptr %agg.result, align 8
// CHECK: %[[RES:[0-9+]]] = load i2048, ptr %agg.result, align 8
// CHECK: store i2048 %[[RES]], ptr %agg.result, align 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]+]])
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
// CHECK: %[[VAR_A:a]].addr = alloca i4096, align 8
// CHECK: %[[VAR_B:b]].addr = alloca i4096, align 8
// CHECK: %[[VAR_A]] = load i4096, ptr %[[ARG1]], align 8
// CHECK: %[[VAR_B]] = load i4096, ptr %[[ARG2]], align 8
// CHECK: store i4096 %[[VAR_A]], ptr %[[VAR_A]].addr, align 8
// CHECK: store i4096 %[[VAR_B]], ptr %[[VAR_B]].addr, align 8
// CHECK: %[[TEMP1:[0-9]+]] = load i4096, ptr %[[VAR_A]].addr, align 8
// CHECK: %[[TEMP2:[0-9]+]] = load i4096, ptr %[[VAR_B]].addr, align 8
// CHECK: %div = sdiv i4096 %[[TEMP1]], %[[TEMP2]]
// CHECK: store i4096 %div, ptr %agg.result, align 8
// CHECK: %[[RES:[0-9+]]] = load i4096, ptr %agg.result, align 8
// CHECK: store i4096 %[[RES]], ptr %agg.result, align 8
// CHECK: ret void
return a / b;
}
16 changes: 8 additions & 8 deletions clang/test/CodeGenSYCL/sycl-intelfpga-bitint.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,24 +3,24 @@
// This test checks that we generate appropriate code for division
// operations of _BitInts of size greater than 128 bits, since it
// is allowed when -fintelfpga is enabled. The test uses a value of
// 2048 for bitint size, the maximum that is currently supported.
// 4096 for bitint size, the maximum that is currently supported.

#include "Inputs/sycl.hpp"

// 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]+]])
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
// CHECK: %[[VAR_A:a]] = load i2048, ptr %[[ARG1]], align 8
// CHECK: %[[VAR_B:b]] = load i2048, ptr %[[ARG2]], align 8
// CHECK: %[[RES:div]] = sdiv i2048 %[[VAR_A]], %[[VAR_B]]
// CHECK: store i2048 %[[RES]], ptr addrspace(4) %agg.result, align 8
// 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]+]])
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
// CHECK: %[[VAR_A:a]] = load i4096, ptr %[[ARG1]], align 8
// CHECK: %[[VAR_B:b]] = load i4096, ptr %[[ARG2]], align 8
// CHECK: %[[RES:div]] = sdiv i4096 %[[VAR_A]], %[[VAR_B]]
// CHECK: store i4096 %[[RES]], ptr addrspace(4) %agg.result, align 8
// CHECK: ret void
return a / b;
}

int main() {
sycl::handler h;
auto lambda = []() {
_BitInt(2048) a, b = 3, c = 4;
_BitInt(4096) a, b = 3, c = 4;
a = foo(b, c);
};
h.single_task(lambda);
Expand Down
9 changes: 5 additions & 4 deletions clang/test/SemaSYCL/sycl-intelfpga.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,13 +5,14 @@

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

// device-intelfpga-error@+4 3{{signed _BitInt of bit sizes greater than 2048 not supported}}
// host-intelfpga-error@+3 3{{signed _BitInt of bit sizes greater than 2048 not supported}}
// device-intelfpga-error@+4 3{{signed _BitInt of bit sizes greater than 4096 not supported}}
// host-intelfpga-error@+3 3{{signed _BitInt of bit sizes greater than 4096 not supported}}
// device-error@+2 3{{signed _BitInt of bit sizes greater than 128 not supported}}
// host-error@+1 3{{signed _BitInt of bit sizes greater than 128 not supported}}
signed _BitInt(2049) foo(signed _BitInt(2049) a, signed _BitInt(2049) b) {
signed _BitInt(4097) foo(signed _BitInt(4097) a, signed _BitInt(4097) b) {
return a / b;
}
// device-error@+4 3{{signed _BitInt of bit sizes greater than 128 not supported}}
Expand Down