466 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 52// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \3// RUN: -emit-llvm -o - %s | FileCheck --check-prefix=HOST %s4// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \5// RUN: -emit-llvm -o - -fcuda-is-device %s | FileCheck --check-prefix=DEV %s6// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \7// RUN: -fatomic-fine-grained-memory -fatomic-ignore-denormal-mode \8// RUN: -emit-llvm -o - -fcuda-is-device %s | FileCheck --check-prefix=OPT %s9 10#include "Inputs/cuda.h"11 12// HOST-LABEL: define dso_local void @_Z12test_defaultPf(13// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {14// HOST-NEXT: [[ENTRY:.*:]]15// HOST-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 816// HOST-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 417// HOST-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 418// HOST-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 819// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 820// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 421// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 422// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 423// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 424// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 425// HOST-NEXT: ret void26//27// DEV-LABEL: define dso_local void @_Z12test_defaultPf(28// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {29// DEV-NEXT: [[ENTRY:.*:]]30// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)31// DEV-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)32// DEV-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)33// DEV-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr34// DEV-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr35// DEV-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr36// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 837// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 838// DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 439// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 440// 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]]41// DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 442// DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 443// DEV-NEXT: ret void44//45// OPT-LABEL: define dso_local void @_Z12test_defaultPf(46// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {47// OPT-NEXT: [[ENTRY:.*:]]48// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)49// OPT-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)50// OPT-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)51// OPT-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr52// OPT-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr53// OPT-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr54// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 855// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 856// OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 457// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 458// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4:![0-9]+]], !amdgpu.ignore.denormal.mode [[META4]]59// OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 460// OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 461// OPT-NEXT: ret void62//63__device__ __host__ void test_default(float *a) {64 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);65}66 67// HOST-LABEL: define dso_local void @_Z8test_onePf(68// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {69// HOST-NEXT: [[ENTRY:.*:]]70// HOST-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 871// HOST-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 472// HOST-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 473// HOST-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 874// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 875// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 476// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 477// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 478// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 479// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 480// HOST-NEXT: ret void81//82// DEV-LABEL: define dso_local void @_Z8test_onePf(83// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {84// DEV-NEXT: [[ENTRY:.*:]]85// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)86// DEV-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)87// DEV-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)88// DEV-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr89// DEV-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr90// DEV-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr91// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 892// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 893// DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 494// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 495// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]]96// DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 497// DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 498// DEV-NEXT: ret void99//100// OPT-LABEL: define dso_local void @_Z8test_onePf(101// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {102// OPT-NEXT: [[ENTRY:.*:]]103// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)104// OPT-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)105// OPT-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)106// OPT-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr107// OPT-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr108// OPT-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr109// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8110// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8111// OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4112// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4113// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]114// OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4115// OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4116// OPT-NEXT: ret void117//118__device__ __host__ void test_one(float *a) {119 [[clang::atomic(no_remote_memory)]] {120 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);121 }122}123 124// HOST-LABEL: define dso_local void @_Z8test_twoPf(125// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {126// HOST-NEXT: [[ENTRY:.*:]]127// HOST-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8128// HOST-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4129// HOST-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4130// HOST-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8131// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8132// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4133// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4134// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4135// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4136// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4137// HOST-NEXT: ret void138//139// DEV-LABEL: define dso_local void @_Z8test_twoPf(140// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {141// DEV-NEXT: [[ENTRY:.*:]]142// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)143// DEV-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)144// DEV-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)145// DEV-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr146// DEV-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr147// DEV-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr148// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8149// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8150// DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4151// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4152// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]153// DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4154// DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4155// DEV-NEXT: ret void156//157// OPT-LABEL: define dso_local void @_Z8test_twoPf(158// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {159// OPT-NEXT: [[ENTRY:.*:]]160// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)161// OPT-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)162// OPT-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)163// OPT-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr164// OPT-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr165// OPT-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr166// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8167// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8168// OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4169// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4170// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META4]]171// OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4172// OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4173// OPT-NEXT: ret void174//175__device__ __host__ void test_two(float *a) {176 [[clang::atomic(remote_memory, ignore_denormal_mode)]] {177 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);178 }179}180 181// HOST-LABEL: define dso_local void @_Z10test_threePf(182// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {183// HOST-NEXT: [[ENTRY:.*:]]184// HOST-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8185// HOST-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4186// HOST-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4187// HOST-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8188// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8189// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4190// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4191// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4192// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4193// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4194// HOST-NEXT: ret void195//196// DEV-LABEL: define dso_local void @_Z10test_threePf(197// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {198// DEV-NEXT: [[ENTRY:.*:]]199// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)200// DEV-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)201// DEV-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)202// DEV-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr203// DEV-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr204// DEV-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr205// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8206// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8207// DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4208// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4209// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]]210// DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4211// DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4212// DEV-NEXT: ret void213//214// OPT-LABEL: define dso_local void @_Z10test_threePf(215// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {216// OPT-NEXT: [[ENTRY:.*:]]217// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)218// OPT-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)219// OPT-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)220// OPT-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr221// OPT-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr222// OPT-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr223// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8224// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8225// OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4226// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4227// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]]228// OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4229// OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4230// OPT-NEXT: ret void231//232__device__ __host__ void test_three(float *a) {233 [[clang::atomic(no_remote_memory, fine_grained_memory, no_ignore_denormal_mode)]] {234 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);235 }236}237 238// HOST-LABEL: define dso_local void @_Z19test_multiple_attrsPf(239// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {240// HOST-NEXT: [[ENTRY:.*:]]241// HOST-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8242// HOST-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4243// HOST-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4244// HOST-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8245// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8246// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4247// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4248// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4249// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4250// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4251// HOST-NEXT: ret void252//253// DEV-LABEL: define dso_local void @_Z19test_multiple_attrsPf(254// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {255// DEV-NEXT: [[ENTRY:.*:]]256// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)257// DEV-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)258// DEV-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)259// DEV-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr260// DEV-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr261// DEV-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr262// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8263// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8264// DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4265// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4266// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]]267// DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4268// DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4269// DEV-NEXT: ret void270//271// OPT-LABEL: define dso_local void @_Z19test_multiple_attrsPf(272// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {273// OPT-NEXT: [[ENTRY:.*:]]274// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)275// OPT-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)276// OPT-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)277// OPT-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr278// OPT-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr279// OPT-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr280// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8281// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8282// OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4283// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4284// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META4]]285// OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4286// OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4287// OPT-NEXT: ret void288//289__device__ __host__ void test_multiple_attrs(float *a) {290 [[clang::atomic(no_remote_memory)]] [[clang::atomic(remote_memory)]] {291 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);292 }293}294 295// HOST-LABEL: define dso_local void @_Z11test_nestedPf(296// HOST-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {297// HOST-NEXT: [[ENTRY:.*:]]298// HOST-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8299// HOST-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4300// HOST-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4301// HOST-NEXT: [[DOTATOMICTMP1:%.*]] = alloca float, align 4302// HOST-NEXT: [[ATOMIC_TEMP2:%.*]] = alloca float, align 4303// HOST-NEXT: [[DOTATOMICTMP3:%.*]] = alloca float, align 4304// HOST-NEXT: [[ATOMIC_TEMP4:%.*]] = alloca float, align 4305// HOST-NEXT: [[DOTATOMICTMP5:%.*]] = alloca float, align 4306// HOST-NEXT: [[ATOMIC_TEMP6:%.*]] = alloca float, align 4307// HOST-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8308// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8309// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4310// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4311// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4312// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4313// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4314// HOST-NEXT: [[TMP4:%.*]] = load ptr, ptr [[A_ADDR]], align 8315// HOST-NEXT: store float 2.000000e+00, ptr [[DOTATOMICTMP1]], align 4316// HOST-NEXT: [[TMP5:%.*]] = load float, ptr [[DOTATOMICTMP1]], align 4317// HOST-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] seq_cst, align 4318// HOST-NEXT: store float [[TMP6]], ptr [[ATOMIC_TEMP2]], align 4319// HOST-NEXT: [[TMP7:%.*]] = load float, ptr [[ATOMIC_TEMP2]], align 4320// HOST-NEXT: [[TMP8:%.*]] = load ptr, ptr [[A_ADDR]], align 8321// HOST-NEXT: store float 3.000000e+00, ptr [[DOTATOMICTMP3]], align 4322// HOST-NEXT: [[TMP9:%.*]] = load float, ptr [[DOTATOMICTMP3]], align 4323// HOST-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] acquire, align 4324// HOST-NEXT: store float [[TMP10]], ptr [[ATOMIC_TEMP4]], align 4325// HOST-NEXT: [[TMP11:%.*]] = load float, ptr [[ATOMIC_TEMP4]], align 4326// HOST-NEXT: [[TMP12:%.*]] = load ptr, ptr [[A_ADDR]], align 8327// HOST-NEXT: store float 4.000000e+00, ptr [[DOTATOMICTMP5]], align 4328// HOST-NEXT: [[TMP13:%.*]] = load float, ptr [[DOTATOMICTMP5]], align 4329// HOST-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] release, align 4330// HOST-NEXT: store float [[TMP14]], ptr [[ATOMIC_TEMP6]], align 4331// HOST-NEXT: [[TMP15:%.*]] = load float, ptr [[ATOMIC_TEMP6]], align 4332// HOST-NEXT: ret void333//334// DEV-LABEL: define dso_local void @_Z11test_nestedPf(335// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {336// DEV-NEXT: [[ENTRY:.*:]]337// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)338// DEV-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)339// DEV-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)340// DEV-NEXT: [[DOTATOMICTMP1:%.*]] = alloca float, align 4, addrspace(5)341// DEV-NEXT: [[ATOMIC_TEMP2:%.*]] = alloca float, align 4, addrspace(5)342// DEV-NEXT: [[DOTATOMICTMP3:%.*]] = alloca float, align 4, addrspace(5)343// DEV-NEXT: [[ATOMIC_TEMP4:%.*]] = alloca float, align 4, addrspace(5)344// DEV-NEXT: [[DOTATOMICTMP5:%.*]] = alloca float, align 4, addrspace(5)345// DEV-NEXT: [[ATOMIC_TEMP6:%.*]] = alloca float, align 4, addrspace(5)346// DEV-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr347// DEV-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr348// DEV-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr349// DEV-NEXT: [[DOTATOMICTMP1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP1]] to ptr350// DEV-NEXT: [[ATOMIC_TEMP2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP2]] to ptr351// DEV-NEXT: [[DOTATOMICTMP3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP3]] to ptr352// DEV-NEXT: [[ATOMIC_TEMP4_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP4]] to ptr353// DEV-NEXT: [[DOTATOMICTMP5_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP5]] to ptr354// DEV-NEXT: [[ATOMIC_TEMP6_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP6]] to ptr355// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8356// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8357// DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4358// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4359// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]]360// DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4361// DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4362// DEV-NEXT: [[TMP4:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8363// DEV-NEXT: store float 2.000000e+00, ptr [[DOTATOMICTMP1_ASCAST]], align 4364// DEV-NEXT: [[TMP5:%.*]] = load float, ptr [[DOTATOMICTMP1_ASCAST]], align 4365// DEV-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4366// DEV-NEXT: store float [[TMP6]], ptr [[ATOMIC_TEMP2_ASCAST]], align 4367// DEV-NEXT: [[TMP7:%.*]] = load float, ptr [[ATOMIC_TEMP2_ASCAST]], align 4368// DEV-NEXT: [[TMP8:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8369// DEV-NEXT: store float 3.000000e+00, ptr [[DOTATOMICTMP3_ASCAST]], align 4370// DEV-NEXT: [[TMP9:%.*]] = load float, ptr [[DOTATOMICTMP3_ASCAST]], align 4371// DEV-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META4]]372// DEV-NEXT: store float [[TMP10]], ptr [[ATOMIC_TEMP4_ASCAST]], align 4373// DEV-NEXT: [[TMP11:%.*]] = load float, ptr [[ATOMIC_TEMP4_ASCAST]], align 4374// DEV-NEXT: [[TMP12:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8375// DEV-NEXT: store float 4.000000e+00, ptr [[DOTATOMICTMP5_ASCAST]], align 4376// DEV-NEXT: [[TMP13:%.*]] = load float, ptr [[DOTATOMICTMP5_ASCAST]], align 4377// DEV-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !amdgpu.no.fine.grained.memory [[META4]]378// DEV-NEXT: store float [[TMP14]], ptr [[ATOMIC_TEMP6_ASCAST]], align 4379// DEV-NEXT: [[TMP15:%.*]] = load float, ptr [[ATOMIC_TEMP6_ASCAST]], align 4380// DEV-NEXT: ret void381//382// OPT-LABEL: define dso_local void @_Z11test_nestedPf(383// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {384// OPT-NEXT: [[ENTRY:.*:]]385// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)386// OPT-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4, addrspace(5)387// OPT-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4, addrspace(5)388// OPT-NEXT: [[DOTATOMICTMP1:%.*]] = alloca float, align 4, addrspace(5)389// OPT-NEXT: [[ATOMIC_TEMP2:%.*]] = alloca float, align 4, addrspace(5)390// OPT-NEXT: [[DOTATOMICTMP3:%.*]] = alloca float, align 4, addrspace(5)391// OPT-NEXT: [[ATOMIC_TEMP4:%.*]] = alloca float, align 4, addrspace(5)392// OPT-NEXT: [[DOTATOMICTMP5:%.*]] = alloca float, align 4, addrspace(5)393// OPT-NEXT: [[ATOMIC_TEMP6:%.*]] = alloca float, align 4, addrspace(5)394// OPT-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr395// OPT-NEXT: [[DOTATOMICTMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP]] to ptr396// OPT-NEXT: [[ATOMIC_TEMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP]] to ptr397// OPT-NEXT: [[DOTATOMICTMP1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP1]] to ptr398// OPT-NEXT: [[ATOMIC_TEMP2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP2]] to ptr399// OPT-NEXT: [[DOTATOMICTMP3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP3]] to ptr400// OPT-NEXT: [[ATOMIC_TEMP4_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP4]] to ptr401// OPT-NEXT: [[DOTATOMICTMP5_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTATOMICTMP5]] to ptr402// OPT-NEXT: [[ATOMIC_TEMP6_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ATOMIC_TEMP6]] to ptr403// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8404// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8405// OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP_ASCAST]], align 4406// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP_ASCAST]], align 4407// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]408// OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP_ASCAST]], align 4409// OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP_ASCAST]], align 4410// OPT-NEXT: [[TMP4:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8411// OPT-NEXT: store float 2.000000e+00, ptr [[DOTATOMICTMP1_ASCAST]], align 4412// OPT-NEXT: [[TMP5:%.*]] = load float, ptr [[DOTATOMICTMP1_ASCAST]], align 4413// OPT-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4414// OPT-NEXT: store float [[TMP6]], ptr [[ATOMIC_TEMP2_ASCAST]], align 4415// OPT-NEXT: [[TMP7:%.*]] = load float, ptr [[ATOMIC_TEMP2_ASCAST]], align 4416// OPT-NEXT: [[TMP8:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8417// OPT-NEXT: store float 3.000000e+00, ptr [[DOTATOMICTMP3_ASCAST]], align 4418// OPT-NEXT: [[TMP9:%.*]] = load float, ptr [[DOTATOMICTMP3_ASCAST]], align 4419// OPT-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META4]]420// OPT-NEXT: store float [[TMP10]], ptr [[ATOMIC_TEMP4_ASCAST]], align 4421// OPT-NEXT: [[TMP11:%.*]] = load float, ptr [[ATOMIC_TEMP4_ASCAST]], align 4422// OPT-NEXT: [[TMP12:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8423// OPT-NEXT: store float 4.000000e+00, ptr [[DOTATOMICTMP5_ASCAST]], align 4424// OPT-NEXT: [[TMP13:%.*]] = load float, ptr [[DOTATOMICTMP5_ASCAST]], align 4425// OPT-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !amdgpu.no.fine.grained.memory [[META4]]426// OPT-NEXT: store float [[TMP14]], ptr [[ATOMIC_TEMP6_ASCAST]], align 4427// OPT-NEXT: [[TMP15:%.*]] = load float, ptr [[ATOMIC_TEMP6_ASCAST]], align 4428// OPT-NEXT: ret void429//430__device__ __host__ void test_nested(float *a) {431 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);432 {433 [[clang::atomic(remote_memory, fine_grained_memory, no_ignore_denormal_mode)]] {434 __scoped_atomic_fetch_max(a, 2, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);435 {436 [[clang::atomic(no_remote_memory)]] {437 __scoped_atomic_fetch_min(a, 3, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP);438 }439 }440 {441 [[clang::atomic(no_fine_grained_memory)]] {442 __scoped_atomic_fetch_sub(a, 4, __ATOMIC_RELEASE, __MEMORY_SCOPE_WVFRNT);443 }444 }445 }446 }447}448 449//450//451//452//453template<typename T> __device__ __host__ void test_template(T *a) {454 [[clang::atomic(no_remote_memory, fine_grained_memory)]] {455 __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);456 }457}458 459template __device__ __host__ void test_template<float>(float *a);460 461//.462// DEV: [[META4]] = !{}463//.464// OPT: [[META4]] = !{}465//.466