[AMDGPU] Report error in clang if wave32 is requested where unsupported #97633
Conversation
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
llvmbot
added
clang
|
@llvm/pr-subscribers-clang @llvm/pr-subscribers-backend-amdgpu Author: Stanislav Mekhanoshin (rampitec) ChangesFull diff: https://github.com/llvm/llvm-project/pull/97633.diff 5 Files Affected:
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index cc7be64656e5b..ea20acdb930fa 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -188,8 +188,12 @@ bool AMDGPUTargetInfo::initFeatureMap(
// TODO: Should move this logic into TargetParser
std::string ErrorMsg;
- if (!insertWaveSizeFeature(CPU, getTriple(), Features, ErrorMsg)) {
- Diags.Report(diag::err_invalid_feature_combination) << ErrorMsg;
+ bool IsCombinationError;
+ if (!insertWaveSizeFeature(CPU, getTriple(), Features, ErrorMsg,
+ IsCombinationError)) {
+ Diags.Report(IsCombinationError ? diag::err_invalid_feature_combination
+ : diag::err_opt_not_valid_on_target)
+ << ErrorMsg;
return false;
}
diff --git a/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl b/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl
index 7dbf5c3c6cd59..4e2f7f86e8402 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl
@@ -1,6 +1,8 @@
// RUN: not %clang_cc1 -triple amdgcn -target-feature +wavefrontsize32 -target-feature +wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck %s
// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx1103 -target-feature +wavefrontsize32 -target-feature +wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck %s
+// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx900 -target-feature +wavefrontsize32 -o /dev/null %s 2>&1 | FileCheck %s --check-prefix=GFX9
// CHECK: error: invalid feature combination: 'wavefrontsize32' and 'wavefrontsize64' are mutually exclusive
+// GFX9: error: option 'wavefrontsize32' cannot be specified on this target
kernel void test() {}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl
index 52f31c1ff0575..e0e3872b566d9 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl
@@ -12,8 +12,7 @@ void test_ballot_wave32(global uint* out, int a, int b) {
*out = __builtin_amdgcn_ballot_w32(a == b); // expected-error {{'__builtin_amdgcn_ballot_w32' needs target feature wavefrontsize32}}
}
-// FIXME: Should error for subtargets that don't support wave32
-__attribute__((target("wavefrontsize32")))
+__attribute__((target("wavefrontsize32"))) // gfx9-error@*:* {{option 'wavefrontsize32' cannot be specified on this target}}
void test_ballot_wave32_target_attr(global uint* out, int a, int b) {
*out = __builtin_amdgcn_ballot_w32(a == b);
}
diff --git a/llvm/include/llvm/TargetParser/TargetParser.h b/llvm/include/llvm/TargetParser/TargetParser.h
index e03d8f6eebfca..858a1fdc01b37 100644
--- a/llvm/include/llvm/TargetParser/TargetParser.h
+++ b/llvm/include/llvm/TargetParser/TargetParser.h
@@ -178,7 +178,8 @@ void fillAMDGPUFeatureMap(StringRef GPU, const Triple &T,
/// Inserts wave size feature for given GPU into features map
bool insertWaveSizeFeature(StringRef GPU, const Triple &T,
- StringMap<bool> &Features, std::string &ErrorMsg);
+ StringMap<bool> &Features, std::string &ErrorMsg,
+ bool &IsCombinationError);
} // namespace AMDGPU
} // namespace llvm
diff --git a/llvm/lib/TargetParser/TargetParser.cpp b/llvm/lib/TargetParser/TargetParser.cpp
index 00df92e0aaded..4bcd966183c67 100644
--- a/llvm/lib/TargetParser/TargetParser.cpp
+++ b/llvm/lib/TargetParser/TargetParser.cpp
@@ -618,15 +618,20 @@ static bool isWave32Capable(StringRef GPU, const Triple &T) {
bool AMDGPU::insertWaveSizeFeature(StringRef GPU, const Triple &T,
StringMap<bool> &Features,
- std::string &ErrorMsg) {
+ std::string &ErrorMsg,
+ bool &IsCombinationError) {
bool IsWave32Capable = isWave32Capable(GPU, T);
const bool IsNullGPU = GPU.empty();
- // FIXME: Not diagnosing wavefrontsize32 on wave64 only targets.
- const bool HaveWave32 =
- (IsWave32Capable || IsNullGPU) && Features.count("wavefrontsize32");
+ const bool HaveWave32 = Features.count("wavefrontsize32");
const bool HaveWave64 = Features.count("wavefrontsize64");
if (HaveWave32 && HaveWave64) {
ErrorMsg = "'wavefrontsize32' and 'wavefrontsize64' are mutually exclusive";
+ IsCombinationError = true;
+ return false;
+ }
+ if (HaveWave32 && !IsNullGPU && !IsWave32Capable) {
+ ErrorMsg = "wavefrontsize32";
+ IsCombinationError = false;
return false;
}
// Don't assume any wavesize with an unknown subtarget.
|
shiltian
reviewed
Jul 4, 2024
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/157/builds/2114 Here is the relevant piece of the build log for the reference: |
rampitec
added a commit
to rampitec/llvm-project
that referenced
this pull request
Jul 9, 2024
rampitec
added a commit
that referenced
this pull request
Jul 9, 2024
|
Fixed #98231
Sorry.
/home/buildbots/llvm-external-buildbots/workers/ppc64le-flang-rhel-test/ppc64le-flang-rhel-clang-build/llvm-project/flang/lib/Frontend/CompilerInstance.cpp:226:44: error: too many arguments to function call, expected 3, have 4
errorMsg)) {
^~~~~~~~
/home/buildbots/llvm-external-buildbots/workers/ppc64le-flang-rhel-test/ppc64le-flang-rhel-clang-build/llvm-project/llvm/include/llvm/TargetParser/TargetParser.h:187:1: note: 'insertWaveSizeFeature' declared here
insertWaveSizeFeature(StringRef GPU, const Triple &T,
^
1 error generated.
|
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/153/builds/2433 Here is the relevant piece of the build log for the reference: |
Fixed. |
aaryanshukla
pushed a commit
to aaryanshukla/llvm-project
that referenced
this pull request
Jul 14, 2024
aaryanshukla
pushed a commit
to aaryanshukla/llvm-project
that referenced
this pull request
Jul 14, 2024
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
No description provided.