Author: Sameer Sahasrabuddhe Date: 2026-04-09T07:33:08+05:30 New Revision: f7365375d9ecda7b715066e25fcb8730957a3c6a
URL: https://github.com/llvm/llvm-project/commit/f7365375d9ecda7b715066e25fcb8730957a3c6a DIFF: https://github.com/llvm/llvm-project/commit/f7365375d9ecda7b715066e25fcb8730957a3c6a.diff LOG: [Clang][NFC] tests showcasing incorrect use of HIP and OpenCL memory scope macros (#188890) The tests demonstrate how incorrect LLVM IR is generated without diagnostics, when an OpenCL or HIP scope number is passed to an AMDGPU intrinsic. #185408 lays the groundwork for properly diagnosing this situation by internally using a separate enum type to represent each set of scope numbers. Added: clang/test/SemaHIP/incorrect-atomic-scope.hip clang/test/SemaOpenCL/incorrect-atomic-scope.cl Modified: Removed: ################################################################################ diff --git a/clang/test/SemaHIP/incorrect-atomic-scope.hip b/clang/test/SemaHIP/incorrect-atomic-scope.hip new file mode 100644 index 0000000000000..07499b412aa30 --- /dev/null +++ b/clang/test/SemaHIP/incorrect-atomic-scope.hip @@ -0,0 +1,31 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1250 -x hip -emit-llvm -fcuda-is-device %s -o - | FileCheck %s + +#define __device__ __attribute__((device)) + +// Both atomics produce the wrong scope in LLVM IR because a HIP scope was +// incorrectly passed where a Clang scope was expected. But no error or warning +// is generated. +// +// CHECK-LABEL: test_builtin_rmw +// CHECK: atomicrmw fmax {{.*}} syncscope("singlethread") +// +// CHECK-LABEL: test_scoped_atomic +// CHECK: atomicrmw {{.*}} syncscope("singlethread") +// +// CHECK-LABEL: test_intrinsic_metadata +// CHECK: call i32 @llvm.amdgcn.flat.load.monitor{{.*}} metadata [[SCOPE:![0-9]+]] +// CHECK: [[SCOPE]] = !{!"wavefront"} + +__device__ void test_builtin_rmw(__attribute__((address_space(3))) float *out, float src) { + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT, false); +} + +__device__ void test_scoped_atomic(int *ptr) { + + __scoped_atomic_fetch_add(ptr, 1, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_AGENT); +} + +__device__ int test_intrinsic_metadata(int* ptr) +{ + return __builtin_amdgcn_flat_load_monitor_b32(ptr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); +} diff --git a/clang/test/SemaOpenCL/incorrect-atomic-scope.cl b/clang/test/SemaOpenCL/incorrect-atomic-scope.cl new file mode 100644 index 0000000000000..add9309f4be30 --- /dev/null +++ b/clang/test/SemaOpenCL/incorrect-atomic-scope.cl @@ -0,0 +1,35 @@ +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,AMDGCN %s +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK %s + +// Both atomics produce the wrong scope in LLVM IR because a HIP scope was +// incorrectly passed where a Clang scope was expected. But no error or warning +// is generated. +// +// CHECK-LABEL: test_builtin_rmw +// CHECK: atomicrmw fmax {{.*}} syncscope("workgroup") +// +// CHECK-LABEL: test_scoped_atomic +// CHECK: atomicrmw {{.*}} syncscope("workgroup") +// +// AMDGCN-LABEL: test_intrinsic_metadata +// AMDGCN: call {{.*}} @llvm.amdgcn.flat.load.monitor{{.*}} metadata [[SCOPE:![0-9]+]] +// AMDGCN: [[SCOPE]] = !{!"workgroup"} + +#if !defined(__SPIRV__) +void test_builtin_rmw(local float *out, float src) { +#else +void test_builtin_rmw(__attribute__((address_space(3))) float *out, float src) { +#endif + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE, false); +} + +void test_scoped_atomic(int *ptr) { + __scoped_atomic_fetch_add(ptr, 1, __ATOMIC_SEQ_CST, __OPENCL_MEMORY_SCOPE_DEVICE); +} + +#if !defined(__SPIRV__) +int test_intrinsic_metadata(int* ptr) +{ + return __builtin_amdgcn_flat_load_monitor_b32(ptr, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} +#endif _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
