https://github.com/ssahasra updated https://github.com/llvm/llvm-project/pull/188890
>From 601d50289bd5ba36f92f0c636310901852cad31f Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <[email protected]> Date: Fri, 27 Mar 2026 06:33:19 +0530 Subject: [PATCH 1/5] [Clang][NFC] Pre-commit tests for #185408 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. --- clang/test/CodeGenHIP/incorrect-atomic-scope.hip | 10 ++++++++++ clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl | 13 +++++++++++++ 2 files changed, 23 insertions(+) create mode 100644 clang/test/CodeGenHIP/incorrect-atomic-scope.hip create mode 100644 clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl diff --git a/clang/test/CodeGenHIP/incorrect-atomic-scope.hip b/clang/test/CodeGenHIP/incorrect-atomic-scope.hip new file mode 100644 index 0000000000000..6ededb84c6eac --- /dev/null +++ b/clang/test/CodeGenHIP/incorrect-atomic-scope.hip @@ -0,0 +1,10 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x hip -emit-llvm -fcuda-is-device %s -o - | FileCheck %s + +#define __device__ __attribute__((device)) + +// CHECK: atomicrmw {{.*}} syncscope("singlethread") + +__device__ void test(__attribute__((address_space(3))) float *out, float src) { + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT, false); // produces the wrong scope, and there is no check for it. +} diff --git a/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl b/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl new file mode 100644 index 0000000000000..7e884b911c014 --- /dev/null +++ b/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl @@ -0,0 +1,13 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK %s +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK %s + +// CHECK: atomicrmw {{.*}} syncscope("workgroup") + +#if !defined(__SPIRV__) +void test(local float *out, float src) { +#else +void test(__attribute__((address_space(3))) float *out, float src) { +#endif + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE, false); // produces the wrong scope, and there is no check for it. +} >From 2f671fa8bbd69cfbe657e72b0c3fa670ae8cc52e Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <[email protected]> Date: Tue, 31 Mar 2026 12:03:47 +0530 Subject: [PATCH 2/5] add an example of scoped atomics --- clang/test/CodeGenHIP/incorrect-atomic-scope.hip | 10 ++++++++-- clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl | 12 +++++++++--- 2 files changed, 17 insertions(+), 5 deletions(-) diff --git a/clang/test/CodeGenHIP/incorrect-atomic-scope.hip b/clang/test/CodeGenHIP/incorrect-atomic-scope.hip index 6ededb84c6eac..ff5c33f88e451 100644 --- a/clang/test/CodeGenHIP/incorrect-atomic-scope.hip +++ b/clang/test/CodeGenHIP/incorrect-atomic-scope.hip @@ -3,8 +3,14 @@ #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: atomicrmw fmax {{.*}} syncscope("singlethread") // CHECK: atomicrmw {{.*}} syncscope("singlethread") -__device__ void test(__attribute__((address_space(3))) float *out, float src) { - *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT, false); // produces the wrong scope, and there is no check for it. +__device__ void test(__attribute__((address_space(3))) float *out, int *ptr, float src) { + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT, false); + __scoped_atomic_fetch_add(ptr, 1, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_AGENT); } diff --git a/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl b/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl index 7e884b911c014..07f20792e6fb5 100644 --- a/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl +++ b/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl @@ -2,12 +2,18 @@ // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK %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: atomicrmw fmax {{.*}} syncscope("workgroup") // CHECK: atomicrmw {{.*}} syncscope("workgroup") #if !defined(__SPIRV__) -void test(local float *out, float src) { +void test(local float *out, int *ptr, float src) { #else -void test(__attribute__((address_space(3))) float *out, float src) { +void test(__attribute__((address_space(3))) float *out, int *ptr, float src) { #endif - *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE, false); // produces the wrong scope, and there is no check for it. + *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE, false); + __scoped_atomic_fetch_add(ptr, 1, __ATOMIC_SEQ_CST, __OPENCL_MEMORY_SCOPE_DEVICE); } >From 0430ca081751f6367b9ab3eee8b0a5d284509eb3 Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <[email protected]> Date: Tue, 31 Mar 2026 13:08:22 +0530 Subject: [PATCH 3/5] add a builtin that becomes an intrinsic with metadata scope --- .../CodeGenHIP/incorrect-atomic-scope.hip | 20 +++++++++++++-- .../CodeGenOpenCL/incorrect-atomic-scope.cl | 25 ++++++++++++++++--- 2 files changed, 40 insertions(+), 5 deletions(-) diff --git a/clang/test/CodeGenHIP/incorrect-atomic-scope.hip b/clang/test/CodeGenHIP/incorrect-atomic-scope.hip index ff5c33f88e451..b257fda009f91 100644 --- a/clang/test/CodeGenHIP/incorrect-atomic-scope.hip +++ b/clang/test/CodeGenHIP/incorrect-atomic-scope.hip @@ -1,5 +1,5 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x hip -emit-llvm -fcuda-is-device %s -o - | FileCheck %s +// 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)) @@ -7,10 +7,26 @@ // 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(__attribute__((address_space(3))) float *out, int *ptr, float src) { +__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/CodeGenOpenCL/incorrect-atomic-scope.cl b/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl index 07f20792e6fb5..aa281698d5f39 100644 --- a/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl +++ b/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl @@ -1,19 +1,38 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK %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") +// +// CHECK-LABEL: test_intrinsic_metadata +// CHECK: call {{.*}} @llvm.amdgcn.flat.load.monitor{{.*}} metadata [[SCOPE:![0-9]+]] +// CHECK: [[SCOPE]] = !{!"workgroup"} #if !defined(__SPIRV__) -void test(local float *out, int *ptr, float src) { +void test_builtin_rmw(local float *out, float src) { #else -void test(__attribute__((address_space(3))) float *out, int *ptr, float src) { +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) +#else +int test_intrinsic_metadata(__attribute__((address_space(0)))int* ptr) +#endif +{ + return __builtin_amdgcn_flat_load_monitor_b32(ptr, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} >From 44b37e59072abb1e95d11c8d122d6e7c9771deca Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <[email protected]> Date: Wed, 1 Apr 2026 11:55:50 +0530 Subject: [PATCH 4/5] remove dependency on amdgpu target --- clang/test/CodeGenHIP/incorrect-atomic-scope.hip | 1 - clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl | 1 - 2 files changed, 2 deletions(-) diff --git a/clang/test/CodeGenHIP/incorrect-atomic-scope.hip b/clang/test/CodeGenHIP/incorrect-atomic-scope.hip index b257fda009f91..07499b412aa30 100644 --- a/clang/test/CodeGenHIP/incorrect-atomic-scope.hip +++ b/clang/test/CodeGenHIP/incorrect-atomic-scope.hip @@ -1,4 +1,3 @@ -// REQUIRES: amdgpu-registered-target // 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)) diff --git a/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl b/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl index aa281698d5f39..ca8604234aef0 100644 --- a/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl +++ b/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl @@ -1,4 +1,3 @@ -// REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK %s // RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK %s >From 1d52ab4eca500e56d0a0a3acf06663d9db5432e5 Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <[email protected]> Date: Wed, 1 Apr 2026 12:18:20 +0530 Subject: [PATCH 5/5] don't use gfx1250 with spirv target --- clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl | 12 +++++------- 1 file changed, 5 insertions(+), 7 deletions(-) diff --git a/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl b/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl index ca8604234aef0..add9309f4be30 100644 --- a/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl +++ b/clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK %s +// 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 @@ -11,9 +11,9 @@ // CHECK-LABEL: test_scoped_atomic // CHECK: atomicrmw {{.*}} syncscope("workgroup") // -// CHECK-LABEL: test_intrinsic_metadata -// CHECK: call {{.*}} @llvm.amdgcn.flat.load.monitor{{.*}} metadata [[SCOPE:![0-9]+]] -// CHECK: [[SCOPE]] = !{!"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) { @@ -29,9 +29,7 @@ void test_scoped_atomic(int *ptr) { #if !defined(__SPIRV__) int test_intrinsic_metadata(int* ptr) -#else -int test_intrinsic_metadata(__attribute__((address_space(0)))int* ptr) -#endif { 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
