Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2024-11-23 13:50:11 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

[AMDGPU] Remove gws feature from GFX12 (#78711)

Browse Source 
This was already done for LLVM. This patch just updates the Clang
builtin handling to match.
This commit is contained in:
Jay Foad 2024-01-19 15:45:53 +00:00 committed by GitHub
parent 97747467f1
commit e21b0b083e
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
3 changed files with 31 additions and 19 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

25
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl
View File

@ -2,23 +2,12 @@
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -verify -S -emit-llvm -o - %s
kernel void builtins_amdgcn_s_barrier_signal_err(global int* in, global int* out, int barrier) {
typedef unsigned int uint;
__builtin_amdgcn_s_barrier_signal(barrier); // expected-error {{'__builtin_amdgcn_s_barrier_signal' must be a constant integer}}
__builtin_amdgcn_s_barrier_wait(-1);
*out = *in;
}
kernel void builtins_amdgcn_s_barrier_wait_err(global int* in, global int* out, int barrier) {
__builtin_amdgcn_s_barrier_signal(-1);
__builtin_amdgcn_s_barrier_wait(barrier); // expected-error {{'__builtin_amdgcn_s_barrier_wait' must be a constant integer}}
*out = *in;
}
kernel void builtins_amdgcn_s_barrier_signal_isfirst_err(global int* in, global int* out, int barrier) {
__builtin_amdgcn_s_barrier_signal_isfirst(barrier); // expected-error {{'__builtin_amdgcn_s_barrier_signal_isfirst' must be a constant integer}}
__builtin_amdgcn_s_barrier_wait(-1);
*out = *in;
kernel void test_builtins_amdgcn_gws_insts(uint a, uint b) {
__builtin_amdgcn_ds_gws_init(a, b); // expected-error {{'__builtin_amdgcn_ds_gws_init' needs target feature gws}}
__builtin_amdgcn_ds_gws_barrier(a, b); // expected-error {{'__builtin_amdgcn_ds_gws_barrier' needs target feature gws}}
__builtin_amdgcn_ds_gws_sema_v(a); // expected-error {{'__builtin_amdgcn_ds_gws_sema_v' needs target feature gws}}
__builtin_amdgcn_ds_gws_sema_br(a, b); // expected-error {{'__builtin_amdgcn_ds_gws_sema_br' needs target feature gws}}
__builtin_amdgcn_ds_gws_sema_p(a); // expected-error {{'__builtin_amdgcn_ds_gws_sema_p' needs target feature gws}}
}

24
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl Normal file
View File

@ -0,0 +1,24 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -verify -S -emit-llvm -o - %s
kernel void builtins_amdgcn_s_barrier_signal_err(global int* in, global int* out, int barrier) {
__builtin_amdgcn_s_barrier_signal(barrier); // expected-error {{'__builtin_amdgcn_s_barrier_signal' must be a constant integer}}
__builtin_amdgcn_s_barrier_wait(-1);
*out = *in;
}
kernel void builtins_amdgcn_s_barrier_wait_err(global int* in, global int* out, int barrier) {
__builtin_amdgcn_s_barrier_signal(-1);
__builtin_amdgcn_s_barrier_wait(barrier); // expected-error {{'__builtin_amdgcn_s_barrier_wait' must be a constant integer}}
*out = *in;
}
kernel void builtins_amdgcn_s_barrier_signal_isfirst_err(global int* in, global int* out, int barrier) {
__builtin_amdgcn_s_barrier_signal_isfirst(barrier); // expected-error {{'__builtin_amdgcn_s_barrier_signal_isfirst' must be a constant integer}}
__builtin_amdgcn_s_barrier_wait(-1);
*out = *in;
}

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

@ -294,7 +294,6 @@ void AMDGPU::fillAMDGPUFeatureMap(StringRef GPU, const Triple &T,
Features["gfx12-insts"] = true;
Features["atomic-fadd-rtn-insts"] = true;
Features["image-insts"] = true;
Features["gws"] = true;
break;
case GK_GFX1151:
case GK_GFX1150: