[clang] [Clang][NFC] Pre-commit tests for #185408 (PR #188890)

[via cfe-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang

Author: Sameer Sahasrabuddhe (ssahasra)

<details>
<summary>Changes</summary>

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.

---
Full diff: https://github.com/llvm/llvm-project/pull/188890.diff


2 Files Affected:

- (added) clang/test/CodeGenHIP/incorrect-atomic-scope.hip (+10) 
- (added) clang/test/CodeGenOpenCL/incorrect-atomic-scope.cl (+13) 


``````````diff
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.
+}

``````````

</details>


https://github.com/llvm/llvm-project/pull/188890