3f4b5893ef
Add an option -munsafe-fp-atomics for AMDGPU target. When enabled, clang adds function attribute "amdgpu-unsafe-fp-atomics" to any functions for amdgpu target. This allows amdgpu backend to use unsafe fp atomic instructions in these functions. Differential Revision: https://reviews.llvm.org/D91546
23 lines
864 B
Plaintext
23 lines
864 B
Plaintext
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
|
|
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
|
|
// RUN: | FileCheck -check-prefixes=NO-UNSAFE-FP-ATOMICS %s
|
|
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
|
|
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
|
|
// RUN: -munsafe-fp-atomics \
|
|
// RUN: | FileCheck -check-prefixes=UNSAFE-FP-ATOMICS %s
|
|
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
|
|
// RUN: -o - -x hip %s -munsafe-fp-atomics \
|
|
// RUN: | FileCheck -check-prefix=NO-UNSAFE-FP-ATOMICS %s
|
|
|
|
#include "Inputs/cuda.h"
|
|
|
|
__device__ void test() {
|
|
// UNSAFE-FP-ATOMICS: define void @_Z4testv() [[ATTR:#[0-9]+]]
|
|
}
|
|
|
|
|
|
// Make sure this is silently accepted on other targets.
|
|
// NO-UNSAFE-FP-ATOMICS-NOT: "amdgpu-unsafe-fp-atomics"
|
|
|
|
// UNSAFE-FP-ATOMICS-DAG: attributes [[ATTR]] = {{.*}}"amdgpu-unsafe-fp-atomics"="true"
|