[clang] [Clang][NFC] tests showcasing incorrect use of HIP and OpenCL memory scope macros (PR #188890)
Sameer Sahasrabuddhe via cfe-commits
cfe-commits at lists.llvm.org
Tue Mar 31 23:49:20 PDT 2026
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 <sameer.sahasrabuddhe at amd.com>
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 <sameer.sahasrabuddhe at amd.com>
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 <sameer.sahasrabuddhe at amd.com>
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 <sameer.sahasrabuddhe at amd.com>
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 <sameer.sahasrabuddhe at amd.com>
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
More information about the cfe-commits
mailing list