[PATCH] D99201: [HIP] Diagnose unaligned atomic for amdgpu

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added reviewers: tra, rjmccall, jfb.
Herald added subscribers: kerbowa, t-tye, tpr, dstuttard, nhaehnle, jvesely, kzhuravl.
yaxunl requested review of this revision.
Herald added a subscriber: wdng.

https://reviews.llvm.org/D99201

Files:
  clang/lib/Driver/ToolChains/Clang.cpp
  clang/test/Driver/hip-options.hip
  clang/test/SemaCUDA/amdgpu-atomic-ops.cu


Index: clang/test/SemaCUDA/amdgpu-atomic-ops.cu
===================================================================
--- /dev/null
+++ clang/test/SemaCUDA/amdgpu-atomic-ops.cu
@@ -0,0 +1,26 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 %s -verify -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
+// RUN:   -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
+// RUN:   -fnative-half-arguments-and-returns -Werror=atomic-alignment
+
+#include "Inputs/cuda.h"
+#include <stdatomic.h>
+
+__device__ _Float16 test_Flot16(_Float16 *p) {
+  return __atomic_fetch_sub(p, 1.0f16, memory_order_relaxed);
+}
+
+__device__ __fp16 test_fp16(__fp16 *p) {
+  return __atomic_fetch_sub(p, 1.0f16, memory_order_relaxed);
+}
+
+struct BigStruct {
+  int data[128];
+};
+
+__device__ void test_big(BigStruct *p1, BigStruct *p2) {
+  __atomic_load(p1, p2, memory_order_relaxed);
+  // expected-error at -1 {{misaligned atomic operation may incur significant performance penalty; the expected alignment (512 bytes) exceeds the actual alignment (4 bytes)}}
+  // expected-error at -2 {{large atomic operation may incur significant performance penalty; the access size (512 bytes) exceeds the max lock-free size (8  bytes)}}
+}
Index: clang/test/Driver/hip-options.hip
===================================================================
--- clang/test/Driver/hip-options.hip
+++ clang/test/Driver/hip-options.hip
@@ -51,3 +51,15 @@
 // RUN:   --cuda-gpu-arch=gfx906  %s 2>&1 | FileCheck -check-prefix=CTA %s
 // CTA: clang{{.*}} "-triple" "x86_64-unknown-linux-gnu" {{.*}} "-mconstructor-aliases"
 // CTA-NOT: clang{{.*}} "-triple" "amdgcn-amd-amdhsa" {{.*}} "-mconstructor-aliases"
+
+// Check -Werror=atomic-alignment is passed for amdpu by default.
+
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -nogpuinc -nogpulib \
+// RUN:   --cuda-gpu-arch=gfx906 %s 2>&1 | FileCheck -check-prefix=WARN-ATOMIC %s
+// WARN-ATOMIC: clang{{.*}} "-triple" "amdgcn-amd-amdhsa" {{.*}} "-Werror=atomic-alignment"
+// WARN-ATOMIC-NOT: clang{{.*}} "-triple" "x86_64-unknown-linux-gnu" {{.*}} "-Werror=atomic-alignment"
+
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -nogpuinc -nogpulib \
+// RUN:   --cuda-gpu-arch=gfx906 -Wno-error=atomic-alignment %s 2>&1 | FileCheck -check-prefix=NO-WARN-ATOMIC %s
+// NO-WARN-ATOMIC-NOT: clang{{.*}} "-triple" "amdgcn-amd-amdhsa" {{.*}} "-Werror=atomic-alignment"
+// NO-WARN-ATOMIC-NOT: clang{{.*}} "-triple" "x86_64-unknown-linux-gnu" {{.*}} "-Werror=atomic-alignment"
Index: clang/lib/Driver/ToolChains/Clang.cpp
===================================================================
--- clang/lib/Driver/ToolChains/Clang.cpp
+++ clang/lib/Driver/ToolChains/Clang.cpp
@@ -6447,6 +6447,18 @@
     if (Args.hasFlag(options::OPT_munsafe_fp_atomics,
                      options::OPT_mno_unsafe_fp_atomics, /*Default=*/false))
       CmdArgs.push_back("-munsafe-fp-atomics");
+
+    // AMDGPU does not support atomic lib call. Treat atomic alignment
+    // warnings as errors by default unless it is disabled explicitly.
+    bool DiagAtomicLibCall = true;
+    for (auto *A : Args.filtered(options::OPT_W_Joined)) {
+      if (StringRef(A->getValue()) == "no-error=atomic-alignment")
+        DiagAtomicLibCall = false;
+      if (StringRef(A->getValue()) == "error=atomic-alignment")
+        DiagAtomicLibCall = true;
+    }
+    if (DiagAtomicLibCall)
+      CmdArgs.push_back("-Werror=atomic-alignment");
   }
 
   // For all the host OpenMP offloading compile jobs we need to pass the targets


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D99201.332731.patch
Type: text/x-patch
Size: 3525 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210323/2d0a9c52/attachment.bin>