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@-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)}}
+}
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
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@-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)}}
+}
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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
