[SYCL] Allow _Bitint of size greater than 128 bits when -fintelfpga is used #6295
There are no files selected for viewing
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,25 @@ | ||
| // RUN: %clang_cc1 -no-opaque-pointers -fsycl-is-host -fintelfpga -triple x86_64 -aux-triple spir64_fpga -emit-llvm %s -o - | FileCheck %s | ||
|
|
||
| // 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 | ||
| // 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: ret void | ||
| return a / b; | ||
| } |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,27 @@ | ||
| // RUN: %clang_cc1 -no-opaque-pointers -fsycl-is-device -fintelfpga -triple spir64_fpga -aux-triple x86_64-unknown-linux-gnu -IInputs -emit-llvm %s -o - | FileCheck %s | ||
|
|
||
| // 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. | ||
|
|
||
| #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: ret void | ||
| return a / b; | ||
| } | ||
|
|
||
| int main() { | ||
| sycl::handler h; | ||
| auto lambda = []() { | ||
| _BitInt(2048) a, b = 3, c = 4; | ||
| a = foo(b, c); | ||
| }; | ||
| h.single_task(lambda); | ||
| } |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,24 @@ | ||
| // RUN: %clang_cc1 -opaque-pointers -fsycl-is-host -fintelfpga -triple x86_64 -aux-triple spir64_fpga -emit-llvm %s -o - | FileCheck %s | ||
|
|
||
| // 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. | ||
|
|
||
| // 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: ret void | ||
| return a / b; | ||
| } |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,27 @@ | ||
| // RUN: %clang_cc1 -opaque-pointers -fsycl-is-device -fintelfpga -triple spir64_fpga -aux-triple x86_64-unknown-linux-gnu -IInputs -emit-llvm %s -o - | FileCheck %s | ||
|
|
||
| // 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. | ||
|
|
||
| #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: ret void | ||
| return a / b; | ||
| } | ||
|
|
||
| int main() { | ||
| sycl::handler h; | ||
| auto lambda = []() { | ||
| _BitInt(2048) a, b = 3, c = 4; | ||
| a = foo(b, c); | ||
| }; | ||
| h.single_task(lambda); | ||
| } |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,23 @@ | ||
| // RUN: %clang_cc1 -fsycl-is-device -fintelfpga -verify=device-intelfpga -fsyntax-only %s -triple spir64_fpga -aux-triple x86_64-unknown-linux-gnu | ||
| // RUN: %clang_cc1 -fsycl-is-host -fintelfpga -verify=host-intelfpga -fsyntax-only %s -triple x86_64 -aux-triple spir64_fpga | ||
| // RUN: %clang_cc1 -fsycl-is-device -verify=device -fsyntax-only %s | ||
| // RUN: %clang_cc1 -fsycl-is-host -verify=host -fsyntax-only %s | ||
|
|
||
| // Tests that we do not issue errors for _Bitints of size greater than 128 | ||
|
There was a problem hiding this comment. Should we add a test where we do not diagnose since that is what this comment implies? I don't have strong preference since CodeGen tests test this anyway. So I leave it up to you. There was a problem hiding this comment. Yeah, I added the "passing" cases in CodeGenSYCL. I don't mind adding a case that would compile cleanly here. Is a value of 215 (greater than x86_64 limit of 128 but lower than fpga limit of 2048) okay with you? There was a problem hiding this comment. Sounds good to me |
||
| // when -fintelfpga is enabled. The backend is expected to be able to handle | ||
| // this. When -fintelfpga is not passed, we continue to diagnose. | ||
|
|
||
| // 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-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) { | ||
| return a / b; | ||
| } | ||
| // device-error@+4 3{{signed _BitInt of bit sizes greater than 128 not supported}} | ||
| // host-error@+3 3{{signed _BitInt of bit sizes greater than 128 not supported}} | ||
| // device-intelfpga-no-diagnostic@+2 | ||
| // host-intelfpga-no-diagnostic@+1 | ||
| signed _BitInt(215) foo(signed _BitInt(215) a, signed _BitInt(215) b) { | ||
| return a + b; | ||
| } | ||
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Hmm... I guess this would cause issues if the code is executed on the other platform, right? So we DO want to enforce the 'host' limit, but only if the 'host' bit-int gets emitted/used/etc? And vise-versa for the device? Or are we just hoping the failure is not quite silent later?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Right, I think that is what I was hoping. In this case we are in a configuration where one platform allows it but the other does not. Rather than preventing compilation here, I was relying on later phases catching it "nicely".
(In the specific case of intelfpga, which is where this can happen currently, there wouldn't be an issue later.)
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
That is incredibly unfortunate... We don't really have a good way of diagnosing with 'deferred diagnostics' on host unfortunately, so I suspect we're going to have to live with this.