// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
// RUN:   -emit-llvm -o - %s | FileCheck --check-prefix=HOST %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN:   -emit-llvm -o - -fcuda-is-device %s | FileCheck --check-prefix=DEV %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN:   -fatomic-fine-grained-memory -fatomic-ignore-denormal-mode \
// RUN:   -emit-llvm -o - -fcuda-is-device %s | FileCheck --check-prefix=OPT %s

#include "Inputs/cuda.h"

// HOST-LABEL: define dso_local void @_Z12test_defaultPf(
// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {
// HOST-NEXT:  [[ENTRY:.*:]]
// HOST-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8
// HOST-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4
// HOST-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4
// HOST-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8
// HOST-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    ret void
//
// DEV-LABEL: define dso_local void @_Z12test_defaultPf(
// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {
// DEV-NEXT:  [[ENTRY:.*:]]
// DEV-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// DEV-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// DEV-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// DEV-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// DEV-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4:![0-9]+]], !amdgpu.no.remote.memory [[META4]]
// DEV-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    ret void
//
// OPT-LABEL: define dso_local void @_Z12test_defaultPf(
// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {
// OPT-NEXT:  [[ENTRY:.*:]]
// OPT-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// OPT-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// OPT-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// OPT-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// OPT-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4:![0-9]+]], !amdgpu.ignore.denormal.mode [[META4]]
// OPT-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    ret void
//
__device__ __host__ void test_default(float *a) {
  __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
}

// HOST-LABEL: define dso_local void @_Z8test_onePf(
// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// HOST-NEXT:  [[ENTRY:.*:]]
// HOST-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8
// HOST-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4
// HOST-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4
// HOST-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8
// HOST-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    ret void
//
// DEV-LABEL: define dso_local void @_Z8test_onePf(
// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// DEV-NEXT:  [[ENTRY:.*:]]
// DEV-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// DEV-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// DEV-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// DEV-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// DEV-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]]
// DEV-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    ret void
//
// OPT-LABEL: define dso_local void @_Z8test_onePf(
// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// OPT-NEXT:  [[ENTRY:.*:]]
// OPT-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// OPT-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// OPT-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// OPT-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// OPT-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]
// OPT-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    ret void
//
__device__ __host__ void test_one(float *a) {
  [[clang::atomic(no_remote_memory)]] {
    __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
  }
}

// HOST-LABEL: define dso_local void @_Z8test_twoPf(
// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// HOST-NEXT:  [[ENTRY:.*:]]
// HOST-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8
// HOST-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4
// HOST-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4
// HOST-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8
// HOST-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    ret void
//
// DEV-LABEL: define dso_local void @_Z8test_twoPf(
// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// DEV-NEXT:  [[ENTRY:.*:]]
// DEV-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// DEV-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// DEV-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// DEV-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// DEV-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]
// DEV-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    ret void
//
// OPT-LABEL: define dso_local void @_Z8test_twoPf(
// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// OPT-NEXT:  [[ENTRY:.*:]]
// OPT-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// OPT-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// OPT-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// OPT-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// OPT-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META4]]
// OPT-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    ret void
//
__device__ __host__ void test_two(float *a) {
  [[clang::atomic(remote_memory, ignore_denormal_mode)]] {
    __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
  }
}

// HOST-LABEL: define dso_local void @_Z10test_threePf(
// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// HOST-NEXT:  [[ENTRY:.*:]]
// HOST-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8
// HOST-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4
// HOST-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4
// HOST-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8
// HOST-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    ret void
//
// DEV-LABEL: define dso_local void @_Z10test_threePf(
// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// DEV-NEXT:  [[ENTRY:.*:]]
// DEV-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// DEV-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// DEV-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// DEV-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// DEV-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]]
// DEV-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    ret void
//
// OPT-LABEL: define dso_local void @_Z10test_threePf(
// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// OPT-NEXT:  [[ENTRY:.*:]]
// OPT-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// OPT-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// OPT-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// OPT-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// OPT-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]]
// OPT-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    ret void
//
__device__ __host__ void test_three(float *a) {
  [[clang::atomic(no_remote_memory, fine_grained_memory, no_ignore_denormal_mode)]] {
    __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
  }
}

// HOST-LABEL: define dso_local void @_Z19test_multiple_attrsPf(
// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// HOST-NEXT:  [[ENTRY:.*:]]
// HOST-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8
// HOST-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4
// HOST-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4
// HOST-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8
// HOST-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    ret void
//
// DEV-LABEL: define dso_local void @_Z19test_multiple_attrsPf(
// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// DEV-NEXT:  [[ENTRY:.*:]]
// DEV-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// DEV-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// DEV-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// DEV-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// DEV-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]]
// DEV-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    ret void
//
// OPT-LABEL: define dso_local void @_Z19test_multiple_attrsPf(
// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// OPT-NEXT:  [[ENTRY:.*:]]
// OPT-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// OPT-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// OPT-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// OPT-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// OPT-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META4]]
// OPT-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    ret void
//
__device__ __host__ void test_multiple_attrs(float *a) {
  [[clang::atomic(no_remote_memory)]] [[clang::atomic(remote_memory)]] {
    __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
  }
}

// HOST-LABEL: define dso_local void @_Z11test_nestedPf(
// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// HOST-NEXT:  [[ENTRY:.*:]]
// HOST-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8
// HOST-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4
// HOST-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4
// HOST-NEXT:    [[DOTATOMICTMP1:%.*]] = alloca float, align 4
// HOST-NEXT:    [[ATOMIC_TEMP2:%.*]] = alloca float, align 4
// HOST-NEXT:    [[DOTATOMICTMP3:%.*]] = alloca float, align 4
// HOST-NEXT:    [[ATOMIC_TEMP4:%.*]] = alloca float, align 4
// HOST-NEXT:    [[DOTATOMICTMP5:%.*]] = alloca float, align 4
// HOST-NEXT:    [[ATOMIC_TEMP6:%.*]] = alloca float, align 4
// HOST-NEXT:    store ptr [[A]], ptr [[A_ADDR]], align 8
// HOST-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT:    store float 2.000000e+00, ptr [[DOTATOMICTMP1]], align 4
// HOST-NEXT:    [[TMP5:%.*]] = load float, ptr [[DOTATOMICTMP1]], align 4
// HOST-NEXT:    [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] seq_cst, align 4
// HOST-NEXT:    store float [[TMP6]], ptr [[ATOMIC_TEMP2]], align 4
// HOST-NEXT:    [[TMP7:%.*]] = load float, ptr [[ATOMIC_TEMP2]], align 4
// HOST-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT:    store float 3.000000e+00, ptr [[DOTATOMICTMP3]], align 4
// HOST-NEXT:    [[TMP9:%.*]] = load float, ptr [[DOTATOMICTMP3]], align 4
// HOST-NEXT:    [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] acquire, align 4
// HOST-NEXT:    store float [[TMP10]], ptr [[ATOMIC_TEMP4]], align 4
// HOST-NEXT:    [[TMP11:%.*]] = load float, ptr [[ATOMIC_TEMP4]], align 4
// HOST-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT:    store float 4.000000e+00, ptr [[DOTATOMICTMP5]], align 4
// HOST-NEXT:    [[TMP13:%.*]] = load float, ptr [[DOTATOMICTMP5]], align 4
// HOST-NEXT:    [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] release, align 4
// HOST-NEXT:    store float [[TMP14]], ptr [[ATOMIC_TEMP6]], align 4
// HOST-NEXT:    [[TMP15:%.*]] = load float, ptr [[ATOMIC_TEMP6]], align 4
// HOST-NEXT:    ret void
//
// DEV-LABEL: define dso_local void @_Z11test_nestedPf(
// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// DEV-NEXT:  [[ENTRY:.*:]]
// DEV-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// DEV-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[DOTATOMICTMP1:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[ATOMIC_TEMP2:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[DOTATOMICTMP3:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[ATOMIC_TEMP4:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[DOTATOMICTMP5:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[ATOMIC_TEMP6:%.*]] = alloca float, align 4, addrspace(5)
// DEV-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// DEV-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// DEV-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// DEV-NEXT:    [[DOTATOMICTMP1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP1]] to ptr
// DEV-NEXT:    [[ATOMIC_TEMP2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP2]] to ptr
// DEV-NEXT:    [[DOTATOMICTMP3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP3]] to ptr
// DEV-NEXT:    [[ATOMIC_TEMP4_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP4]] to ptr
// DEV-NEXT:    [[DOTATOMICTMP5_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP5]] to ptr
// DEV-NEXT:    [[ATOMIC_TEMP6_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP6]] to ptr
// DEV-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]]
// DEV-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// DEV-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    store float 2.000000e+00, ptr [[DOTATOMICTMP1_ASCAST]], align 4
// DEV-NEXT:    [[TMP5:%.*]] = load float, ptr [[DOTATOMICTMP1_ASCAST]], align 4
// DEV-NEXT:    [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4
// DEV-NEXT:    store float [[TMP6]], ptr [[ATOMIC_TEMP2_ASCAST]], align 4
// DEV-NEXT:    [[TMP7:%.*]] = load float, ptr [[ATOMIC_TEMP2_ASCAST]], align 4
// DEV-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    store float 3.000000e+00, ptr [[DOTATOMICTMP3_ASCAST]], align 4
// DEV-NEXT:    [[TMP9:%.*]] = load float, ptr [[DOTATOMICTMP3_ASCAST]], align 4
// DEV-NEXT:    [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META4]]
// DEV-NEXT:    store float [[TMP10]], ptr [[ATOMIC_TEMP4_ASCAST]], align 4
// DEV-NEXT:    [[TMP11:%.*]] = load float, ptr [[ATOMIC_TEMP4_ASCAST]], align 4
// DEV-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT:    store float 4.000000e+00, ptr [[DOTATOMICTMP5_ASCAST]], align 4
// DEV-NEXT:    [[TMP13:%.*]] = load float, ptr [[DOTATOMICTMP5_ASCAST]], align 4
// DEV-NEXT:    [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !amdgpu.no.fine.grained.memory [[META4]]
// DEV-NEXT:    store float [[TMP14]], ptr [[ATOMIC_TEMP6_ASCAST]], align 4
// DEV-NEXT:    [[TMP15:%.*]] = load float, ptr [[ATOMIC_TEMP6_ASCAST]], align 4
// DEV-NEXT:    ret void
//
// OPT-LABEL: define dso_local void @_Z11test_nestedPf(
// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
// OPT-NEXT:  [[ENTRY:.*:]]
// OPT-NEXT:    [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// OPT-NEXT:    [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[DOTATOMICTMP1:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[ATOMIC_TEMP2:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[DOTATOMICTMP3:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[ATOMIC_TEMP4:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[DOTATOMICTMP5:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[ATOMIC_TEMP6:%.*]] = alloca float, align 4, addrspace(5)
// OPT-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
// OPT-NEXT:    [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr
// OPT-NEXT:    [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr
// OPT-NEXT:    [[DOTATOMICTMP1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP1]] to ptr
// OPT-NEXT:    [[ATOMIC_TEMP2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP2]] to ptr
// OPT-NEXT:    [[DOTATOMICTMP3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP3]] to ptr
// OPT-NEXT:    [[ATOMIC_TEMP4_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP4]] to ptr
// OPT-NEXT:    [[DOTATOMICTMP5_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP5]] to ptr
// OPT-NEXT:    [[ATOMIC_TEMP6_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP6]] to ptr
// OPT-NEXT:    store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]
// OPT-NEXT:    store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4
// OPT-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    store float 2.000000e+00, ptr [[DOTATOMICTMP1_ASCAST]], align 4
// OPT-NEXT:    [[TMP5:%.*]] = load float, ptr [[DOTATOMICTMP1_ASCAST]], align 4
// OPT-NEXT:    [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4
// OPT-NEXT:    store float [[TMP6]], ptr [[ATOMIC_TEMP2_ASCAST]], align 4
// OPT-NEXT:    [[TMP7:%.*]] = load float, ptr [[ATOMIC_TEMP2_ASCAST]], align 4
// OPT-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    store float 3.000000e+00, ptr [[DOTATOMICTMP3_ASCAST]], align 4
// OPT-NEXT:    [[TMP9:%.*]] = load float, ptr [[DOTATOMICTMP3_ASCAST]], align 4
// OPT-NEXT:    [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META4]]
// OPT-NEXT:    store float [[TMP10]], ptr [[ATOMIC_TEMP4_ASCAST]], align 4
// OPT-NEXT:    [[TMP11:%.*]] = load float, ptr [[ATOMIC_TEMP4_ASCAST]], align 4
// OPT-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT:    store float 4.000000e+00, ptr [[DOTATOMICTMP5_ASCAST]], align 4
// OPT-NEXT:    [[TMP13:%.*]] = load float, ptr [[DOTATOMICTMP5_ASCAST]], align 4
// OPT-NEXT:    [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !amdgpu.no.fine.grained.memory [[META4]]
// OPT-NEXT:    store float [[TMP14]], ptr [[ATOMIC_TEMP6_ASCAST]], align 4
// OPT-NEXT:    [[TMP15:%.*]] = load float, ptr [[ATOMIC_TEMP6_ASCAST]], align 4
// OPT-NEXT:    ret void
//
__device__ __host__ void test_nested(float *a) {
  __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
  {
    [[clang::atomic(remote_memory, fine_grained_memory, no_ignore_denormal_mode)]] {
      __scoped_atomic_fetch_max(a, 2, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
      {
        [[clang::atomic(no_remote_memory)]] {
          __scoped_atomic_fetch_min(a, 3, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP);
        }
      }
      {
        [[clang::atomic(no_fine_grained_memory)]] {
          __scoped_atomic_fetch_sub(a, 4, __ATOMIC_RELEASE, __MEMORY_SCOPE_WVFRNT);
        }
      }
    }
  }
}

//
//
//
//
template<typename T> __device__ __host__ void test_template(T *a) {
  [[clang::atomic(no_remote_memory, fine_grained_memory)]] {
    __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
  }
}

template __device__ __host__ void test_template<float>(float *a);

//.
// DEV: [[META4]] = !{}
//.
// OPT: [[META4]] = !{}
//.