[AMDGPU] Diagnose unaligned atomic for amdgpu

[yaxunl]

rebase and simplify code

[yaxunl]

ping

[tra]

clang/lib/Driver/ToolChains/Clang.cpp

6451

Should it be done from HIPAMDToolChain::addClangWarningOptions ?

That's where Darwin does similar propotion from a waring to an error.

[MaskRay]

clang/test/Driver/hip-options.hip

56

Prefer --target= to deprecated (since 3.x) -target

[yaxunl]

clang/lib/Driver/ToolChains/Clang.cpp

6451

will do. Since this option is shared between all AMDGPU toolchains, I will add it to AMDGPUToolChain::addClangWarningOptions and let the derived toolchains call it.

clang/test/Driver/hip-options.hip

56

will fix

[yaxunl]

revised by Artem's and Fangrui's comments

clang/lib/Driver/ToolChains/Clang.cpp

   }
 
   if (Triple.isAMDGPU()) {
     handleAMDGPUCodeObjectVersionOptions(D, Args, CmdArgs);
 
     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
   // information using -fopenmp-targets= option.
   if (JA.isHostOffloading(Action::OFK_OpenMP)) {
     SmallString<128> TargetInfo("-fopenmp-targets=");
 
     Arg *Tgts = Args.getLastArg(options::OPT_fopenmp_targets_EQ);

clang/test/Driver/hip-options.hip

 // FIX-OVERLOAD: clang{{.*}} "-triple" "x86_64-unknown-linux-gnu" {{.*}} "-fgpu-exclude-wrong-side-overloads" "-fgpu-defer-diag"
 
 // Check -mconstructor-aliases is not passed to device compilation.
 
 // RUN: %clang -### -target x86_64-unknown-linux-gnu -nogpuinc -nogpulib \
 // 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"

clang/test/SemaCUDA/amdgpu-atomic-ops.cu

+// 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@-1 {{misaligned atomic operation may incur significant performance penalty; the expected alignment (512 bytes) exceeds the actual alignment (4 bytes)}}
+  // expected-error@-2 {{large atomic operation may incur significant performance penalty; the access size (512 bytes) exceeds the max lock-free size (8  bytes)}}
+}