diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index fda16e42d2c6b..4271bd2d55618 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -268,7 +268,7 @@ TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts") TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f64, "dd*1d", "t", "gfx90a-insts") TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f32, "ff*1f", "t", "atomic-fadd-rtn-insts") -TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2f16, "V2hV2h*1V2h", "t", "atomic-buffer-global-pk-add-f16-insts") +TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2f16, "V2xV2x*1V2x", "n", "atomic-buffer-global-pk-add-f16-insts") TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmin_f64, "dd*1d", "t", "gfx90a-insts") TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmax_f64, "dd*1d", "t", "gfx90a-insts") @@ -280,11 +280,11 @@ TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f64, "dd*3d", "t", "gfx90a-insts" TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f32, "ff*3f", "t", "gfx8-insts") TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f32, "ff*0f", "t", "gfx940-insts") -TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2f16, "V2hV2h*0V2h", "t", "atomic-flat-pk-add-16-insts") +TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2f16, "V2xV2x*0V2x", "n", "atomic-flat-pk-add-16-insts") TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2bf16, "V2sV2s*0V2s", "t", "atomic-flat-pk-add-16-insts") TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s", "t", "atomic-global-pk-add-bf16-inst") TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t", "atomic-ds-pk-add-16-insts") -TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2hV2h*3V2h", "t", "atomic-ds-pk-add-16-insts") +TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2xV2x*3V2x", "n", "atomic-ds-pk-add-16-insts") TARGET_BUILTIN(__builtin_amdgcn_load_to_lds, "vv*v*3IUiIiIUi", "", "vmem-to-lds-load-insts") TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t", "vmem-to-lds-load-insts") diff --git a/clang/test/Sema/builtin-amdgcn-atomic-fadd-v2f16-type-err.c b/clang/test/Sema/builtin-amdgcn-atomic-fadd-v2f16-type-err.c new file mode 100644 index 0000000000000..c6d87642d99b3 --- /dev/null +++ b/clang/test/Sema/builtin-amdgcn-atomic-fadd-v2f16-type-err.c @@ -0,0 +1,46 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -verify %s + +// Test semantic analysis for AMDGCN atomic fadd v2f16 builtins +// These tests ensure proper type checking for the builtin arguments + +typedef _Float16 v2f16 __attribute__((ext_vector_type(2))); +typedef float v2f32 __attribute__((ext_vector_type(2))); +typedef _Float16 v4f16 __attribute__((ext_vector_type(4))); + +void test_global_atomic_fadd_v2f16_negative() { + v2f16 val; + v2f32 val_f32; + v4f16 val_v4f16; + _Float16 __attribute__((ext_vector_type(2))) __attribute__((address_space(1))) *ptr_v2f16; + + __builtin_amdgcn_global_atomic_fadd_v2f16(ptr_v2f16, val_f32); // expected-error{{passing 'v2f32'}} + __builtin_amdgcn_global_atomic_fadd_v2f16(ptr_v2f16, val_v4f16); // expected-error{{passing 'v4f16'}} + __builtin_amdgcn_global_atomic_fadd_v2f16(ptr_v2f16); // expected-error{{too few arguments to function call}} + __builtin_amdgcn_global_atomic_fadd_v2f16(ptr_v2f16, val, val); // expected-error{{too many arguments to function call}} +} + +void test_flat_atomic_fadd_v2f16_negative() { + v2f16 val; + v2f32 val_f32; + v4f16 val_v4f16; + _Float16 __attribute__((ext_vector_type(2))) __attribute__((address_space(0))) *ptr_v2f16; + + // Same error patterns for flat atomic + __builtin_amdgcn_flat_atomic_fadd_v2f16(ptr_v2f16, val_f32); // expected-error{{passing 'v2f32'}} + __builtin_amdgcn_flat_atomic_fadd_v2f16(ptr_v2f16, val_v4f16); // expected-error{{passing 'v4f16'}} + __builtin_amdgcn_flat_atomic_fadd_v2f16(ptr_v2f16); // expected-error{{too few arguments to function call}} + __builtin_amdgcn_flat_atomic_fadd_v2f16(ptr_v2f16, val, val); // expected-error{{too many arguments to function call}} +} + +void test_ds_atomic_fadd_v2f16_negative() { + v2f16 val; + v2f32 val_f32; + v4f16 val_v4f16; + _Float16 __attribute__((ext_vector_type(2))) __attribute__((address_space(3))) *ptr_v2f16; + + // Same error patterns for ds atomic + __builtin_amdgcn_ds_atomic_fadd_v2f16(ptr_v2f16, val_f32); // expected-error{{passing 'v2f32'}} + __builtin_amdgcn_ds_atomic_fadd_v2f16(ptr_v2f16, val_v4f16); // expected-error{{passing 'v4f16'}} + __builtin_amdgcn_ds_atomic_fadd_v2f16(ptr_v2f16); // expected-error{{too few arguments to function call}} + __builtin_amdgcn_ds_atomic_fadd_v2f16(ptr_v2f16, val, val); // expected-error{{too many arguments to function call}} +}