yggdrasil/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

[AMDGPU] Report error in clang if wave32 is requested where unsupported (#97633)

Browse Source
This commit is contained in:
Stanislav Mekhanoshin 2024-07-09 14:25:58 -07:00 committed by GitHub
parent af21bc1917
commit f363e30f15
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
5 changed files with 31 additions and 16 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
clang/lib/Basic/Targets/AMDGPU.cpp
View File

@ -187,9 +187,15 @@ bool AMDGPUTargetInfo::initFeatureMap(
return false;
// TODO: Should move this logic into TargetParser
std::string ErrorMsg;
if (!insertWaveSizeFeature(CPU, getTriple(), Features, ErrorMsg)) {
Diags.Report(diag::err_invalid_feature_combination) << ErrorMsg;
auto HasError = insertWaveSizeFeature(CPU, getTriple(), Features);
switch (HasError.first) {
default:
break;
case llvm::AMDGPU::INVALID_FEATURE_COMBINATION:
Diags.Report(diag::err_invalid_feature_combination) << HasError.second;
return false;
case llvm::AMDGPU::UNSUPPORTED_TARGET_FEATURE:
Diags.Report(diag::err_opt_not_valid_on_target) << HasError.second;
return false;
}

2
clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl
View File

@ -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() {}

3
clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl
View File

@ -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);
}

11
llvm/include/llvm/TargetParser/TargetParser.h
View File

@ -157,6 +157,12 @@ enum ArchFeatureKind : uint32_t {
FEATURE_WGP = 1 << 9,
};
enum FeatureError : uint32_t {
NO_ERROR = 0,
INVALID_FEATURE_COMBINATION,
UNSUPPORTED_TARGET_FEATURE
};
StringRef getArchFamilyNameAMDGCN(GPUKind AK);
StringRef getArchNameAMDGCN(GPUKind AK);
@ -177,8 +183,9 @@ void fillAMDGPUFeatureMap(StringRef GPU, const Triple &T,
StringMap<bool> &Features);
/// Inserts wave size feature for given GPU into features map
bool insertWaveSizeFeature(StringRef GPU, const Triple &T,
StringMap<bool> &Features, std::string &ErrorMsg);
std::pair<FeatureError, StringRef>
insertWaveSizeFeature(StringRef GPU, const Triple &T,
StringMap<bool> &Features);
} // namespace AMDGPU
} // namespace llvm

19
llvm/lib/TargetParser/TargetParser.cpp
View File

@ -616,18 +616,19 @@ static bool isWave32Capable(StringRef GPU, const Triple &T) {
return IsWave32Capable;
}
bool AMDGPU::insertWaveSizeFeature(StringRef GPU, const Triple &T,
StringMap<bool> &Features,
std::string &ErrorMsg) {
std::pair<FeatureError, StringRef>
AMDGPU::insertWaveSizeFeature(StringRef GPU, const Triple &T,
StringMap<bool> &Features) {
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";
return false;
return {AMDGPU::INVALID_FEATURE_COMBINATION,
"'wavefrontsize32' and 'wavefrontsize64' are mutually exclusive"};
}
if (HaveWave32 && !IsNullGPU && !IsWave32Capable) {
return {AMDGPU::UNSUPPORTED_TARGET_FEATURE, "wavefrontsize32"};
}
// Don't assume any wavesize with an unknown subtarget.
if (!IsNullGPU) {
@ -638,5 +639,5 @@ bool AMDGPU::insertWaveSizeFeature(StringRef GPU, const Triple &T,
Features.insert(std::make_pair(DefaultWaveSizeFeature, true));
}
}
return true;
return {NO_ERROR, StringRef()};
}