diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index 4f1337ce4d657..33155e9e12fd6 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -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 diff --git a/clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-host-intelfpga-bitint.cpp b/clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-host-intelfpga-bitint.cpp index c5a3d0949adeb..f65d311cead68 100644 --- a/clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-host-intelfpga-bitint.cpp +++ b/clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-host-intelfpga-bitint.cpp @@ -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; } diff --git a/clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-intelfpga-bitint.cpp b/clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-intelfpga-bitint.cpp index 95b7988033ea9..e132278a542a3 100644 --- a/clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-intelfpga-bitint.cpp +++ b/clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-intelfpga-bitint.cpp @@ -3,16 +3,16 @@ // 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; } @@ -20,7 +20,7 @@ signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) 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); diff --git a/clang/test/CodeGenSYCL/sycl-host-intelfpga-bitint.cpp b/clang/test/CodeGenSYCL/sycl-host-intelfpga-bitint.cpp index 006514989eaee..725314be239d1 100644 --- a/clang/test/CodeGenSYCL/sycl-host-intelfpga-bitint.cpp +++ b/clang/test/CodeGenSYCL/sycl-host-intelfpga-bitint.cpp @@ -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; } diff --git a/clang/test/CodeGenSYCL/sycl-intelfpga-bitint.cpp b/clang/test/CodeGenSYCL/sycl-intelfpga-bitint.cpp index e69ffb2ca3742..ed428f01d600a 100644 --- a/clang/test/CodeGenSYCL/sycl-intelfpga-bitint.cpp +++ b/clang/test/CodeGenSYCL/sycl-intelfpga-bitint.cpp @@ -3,16 +3,16 @@ // 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; } @@ -20,7 +20,7 @@ signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) 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); diff --git a/clang/test/SemaSYCL/sycl-intelfpga.cpp b/clang/test/SemaSYCL/sycl-intelfpga.cpp index 63f32ce57a594..e43762530c160 100644 --- a/clang/test/SemaSYCL/sycl-intelfpga.cpp +++ b/clang/test/SemaSYCL/sycl-intelfpga.cpp @@ -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}}