Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions clang/include/clang/Basic/BuiltinsAMDGPU.def
Original file line number Diff line number Diff line change
Expand Up @@ -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")

Expand All @@ -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")
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Also make sure to test this with all of these similar intrinsics, the "t"s are all suspicious (I'm guessing if it was for anything, it was hacking around the address space of the pointer across the languages)

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")

Expand Down
46 changes: 46 additions & 0 deletions clang/test/Sema/builtin-amdgcn-atomic-fadd-v2f16-type-err.c
Original file line number Diff line number Diff line change
@@ -0,0 +1,46 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -verify %s
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can you verify we have test coverage using OpenCL plus one of the C/C++/HIP variants


// 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)));

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can you also test some of the cases with the type embedded in a structure like was crashing

Copy link
Author

@tcgu-amd tcgu-amd Oct 14, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think the type was from hip though. How do I add a type from hip_fp16 to a clang test? Do I need to put the test in one of the hip/opencl folders? Thanks!

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}}
}
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

bf16 vector cases and f32 cases are also marked with t, and probably broken in the same way