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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
