416 lines · plain
1// REQUIRES: amdgpu-registered-target2// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,GCN %s3// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,GCN %s4// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,GCN %s5// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,GCN %s6// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,AMDGCNSPIRV %s7 8#define INVALID_MEMORY_SCOPE (__MEMORY_SCOPE_CLUSTR+1)9 10#pragma OPENCL EXTENSION cl_khr_fp16 : enable11 12typedef unsigned long ulong;13typedef unsigned int uint;14 15// CHECK-LABEL: @test_div_fixup_f1616// CHECK: {{.*}}call{{.*}} half @llvm.amdgcn.div.fixup.f1617void test_div_fixup_f16(global half* out, half a, half b, half c)18{19 *out = __builtin_amdgcn_div_fixuph(a, b, c);20}21 22// CHECK-LABEL: @test_rcp_f1623// CHECK: {{.*}}call{{.*}} half @llvm.amdgcn.rcp.f1624void test_rcp_f16(global half* out, half a)25{26 *out = __builtin_amdgcn_rcph(a);27}28 29// CHECK-LABEL: @test_sqrt_f1630// CHECK: {{.*}}call{{.*}} half @llvm.{{((amdgcn.){0,1})}}sqrt.f1631void test_sqrt_f16(global half* out, half a)32{33 *out = __builtin_amdgcn_sqrth(a);34}35 36// CHECK-LABEL: @test_rsq_f1637// CHECK: {{.*}}call{{.*}} half @llvm.amdgcn.rsq.f1638void test_rsq_f16(global half* out, half a)39{40 *out = __builtin_amdgcn_rsqh(a);41}42 43// CHECK-LABEL: @test_sin_f1644// CHECK: {{.*}}call{{.*}} half @llvm.amdgcn.sin.f1645void test_sin_f16(global half* out, half a)46{47 *out = __builtin_amdgcn_sinh(a);48}49 50// CHECK-LABEL: @test_cos_f1651// CHECK: {{.*}}call{{.*}} half @llvm.amdgcn.cos.f1652void test_cos_f16(global half* out, half a)53{54 *out = __builtin_amdgcn_cosh(a);55}56 57// CHECK-LABEL: @test_ldexp_f1658// CHECK: [[TRUNC:%[0-9a-z]+]] = trunc i3259// CHECK: {{.*}}call{{.*}} half @llvm.ldexp.f16.i16(half %a, i16 [[TRUNC]])60void test_ldexp_f16(global half* out, half a, int b)61{62 *out = __builtin_amdgcn_ldexph(a, b);63}64 65// CHECK-LABEL: @test_frexp_mant_f1666// CHECK: {{.*}}call{{.*}} half @llvm.amdgcn.frexp.mant.f1667void test_frexp_mant_f16(global half* out, half a)68{69 *out = __builtin_amdgcn_frexp_manth(a);70}71 72// CHECK-LABEL: @test_frexp_exp_f1673// CHECK: {{.*}}call{{.*}} i16 @llvm.amdgcn.frexp.exp.i16.f1674void test_frexp_exp_f16(global short* out, half a)75{76 *out = __builtin_amdgcn_frexp_exph(a);77}78 79// CHECK-LABEL: @test_fract_f1680// CHECK: {{.*}}call{{.*}} half @llvm.amdgcn.fract.f1681void test_fract_f16(global half* out, half a)82{83 *out = __builtin_amdgcn_fracth(a);84}85 86// CHECK-LABEL: @test_class_f1687// CHECK: {{.*}}call{{.*}} i1 @llvm.amdgcn.class.f1688void test_class_f16(global half* out, half a, int b)89{90 *out = __builtin_amdgcn_classh(a, b);91}92 93// CHECK-LABEL: @test_s_memrealtime94// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.s.memrealtime()95void test_s_memrealtime(global ulong* out)96{97 *out = __builtin_amdgcn_s_memrealtime();98}99 100// CHECK-LABEL: @test_s_dcache_wb()101// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.dcache.wb()102void test_s_dcache_wb()103{104 __builtin_amdgcn_s_dcache_wb();105}106 107// CHECK-LABEL: @test_mov_dpp_int108// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %src, i32 0, i32 0, i32 0, i1 false)109void test_mov_dpp_int(global int* out, int src)110{111 *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);112}113 114// CHECK-LABEL: @test_mov_dpp_long115// CHECK: %0 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 poison, i64 %x, i32 257, i32 15, i32 15, i1 false)116// CHECK-NEXT: store i64 %0,117void test_mov_dpp_long(long x, global long *p) {118 *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);119}120 121// CHECK-LABEL: @test_mov_dpp_float122// CHECK: %0 = bitcast float %x to i32123// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %0, i32 257, i32 15, i32 15, i1 false)124// CHECK-NEXT: store i32 %1,125void test_mov_dpp_float(float x, global float *p) {126 *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);127}128 129// CHECK-LABEL: @test_mov_dpp_double130// CHECK: %0 = bitcast double %x to i64131// CHECK-NEXT: %1 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 poison, i64 %0, i32 257, i32 15, i32 15, i1 false)132// CHECK-NEXT: store i64 %1,133void test_mov_dpp_double(double x, global double *p) {134 *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);135}136 137// CHECK-LABEL: @test_mov_dpp_short138// CHECK: %0 = zext i16 %x to i32139// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %0, i32 257, i32 15, i32 15, i1 false)140// CHECK-NEXT: %2 = trunc i32 %1 to i16141// CHECK-NEXT: store i16 %2,142void test_mov_dpp_short(short x, global short *p) {143 *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);144}145 146// CHECK-LABEL: @test_mov_dpp_char147// CHECK: %0 = zext i8 %x to i32148// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %0, i32 257, i32 15, i32 15, i1 false)149// CHECK-NEXT: %2 = trunc i32 %1 to i8150// CHECK-NEXT: store i8 %2,151void test_mov_dpp_char(char x, global char *p) {152 *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);153}154 155// CHECK-LABEL: @test_mov_dpp_half156// CHECK: %0 = load i16,157// CHECK: %1 = zext i16 %0 to i32158// CHECK-NEXT: %2 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %1, i32 257, i32 15, i32 15, i1 false)159// CHECK-NEXT: %3 = trunc i32 %2 to i16160// CHECK-NEXT: store i16 %3,161void test_mov_dpp_half(half *x, global half *p) {162 *p = __builtin_amdgcn_mov_dpp(*x, 0x101, 0xf, 0xf, 0);163}164 165// CHECK-LABEL: @test_update_dpp_int166// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false)167void test_update_dpp_int(global int* out, int arg1, int arg2)168{169 *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false);170}171 172// CHECK-LABEL: @test_update_dpp_long173// CHECK: %0 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 %x, i64 %x, i32 257, i32 15, i32 15, i1 false)174// CHECK-NEXT: store i64 %0,175void test_update_dpp_long(long x, global long *p) {176 *p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);177}178 179// CHECK-LABEL: @test_update_dpp_float180// CHECK: %0 = bitcast float %x to i32181// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %0, i32 257, i32 15, i32 15, i1 false)182// CHECK-NEXT: store i32 %1,183void test_update_dpp_float(float x, global float *p) {184 *p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);185}186 187// CHECK-LABEL: @test_update_dpp_double188// CHECK: %0 = bitcast double %x to i64189// CHECK-NEXT: %1 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 %0, i64 %0, i32 257, i32 15, i32 15, i1 false)190// CHECK-NEXT: store i64 %1,191void test_update_dpp_double(double x, global double *p) {192 *p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);193}194 195// CHECK-LABEL: @test_update_dpp_short196// CHECK: %0 = zext i16 %x to i32197// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %0, i32 257, i32 15, i32 15, i1 false)198// CHECK-NEXT: %2 = trunc i32 %1 to i16199// CHECK-NEXT: store i16 %2,200void test_update_dpp_short(short x, global short *p) {201 *p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);202}203 204// CHECK-LABEL: @test_update_dpp_char205// CHECK: %0 = zext i8 %x to i32206// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %0, i32 257, i32 15, i32 15, i1 false)207// CHECK-NEXT: %2 = trunc i32 %1 to i8208// CHECK-NEXT: store i8 %2,209void test_update_dpp_char(char x, global char *p) {210 *p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);211}212 213// CHECK-LABEL: @test_update_dpp_half214// CHECK: %0 = load i16,215// CHECK: %1 = zext i16 %0 to i32216// CHECK-NEXT: %2 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %1, i32 %1, i32 257, i32 15, i32 15, i1 false)217// CHECK-NEXT: %3 = trunc i32 %2 to i16218// CHECK-NEXT: store i16 %3,219void test_update_dpp_half(half *x, global half *p) {220 *p = __builtin_amdgcn_update_dpp(*x, *x, 0x101, 0xf, 0xf, 0);221}222 223// CHECK-LABEL: @test_update_dpp_int_uint224// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false)225void test_update_dpp_int_uint(global int* out, int arg1, unsigned int arg2)226{227 *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false);228}229 230// CHECK-LABEL: @test_update_dpp_lit_int231// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 5, i32 %arg1, i32 0, i32 0, i32 0, i1 false)232void test_update_dpp_lit_int(global int* out, int arg1)233{234 *out = __builtin_amdgcn_update_dpp(5, arg1, 0, 0, 0, false);235}236 237__constant int gi = 5;238 239// CHECK-LABEL: @test_update_dpp_const_int240// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 5, i32 %arg1, i32 0, i32 0, i32 0, i1 false)241void test_update_dpp_const_int(global int* out, int arg1)242{243 *out = __builtin_amdgcn_update_dpp(gi, arg1, 0, 0, 0, false);244}245 246// CHECK-LABEL: @test_ds_fadd247// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}}248// CHECK: atomicrmw volatile fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}}249 250// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src acquire, align 4{{$}}251// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src acquire, align 4{{$}}252// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src release, align 4{{$}}253// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src acq_rel, align 4{{$}}254// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src seq_cst, align 4{{$}}255// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src seq_cst, align 4{{$}}256 257// GCN: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("agent") monotonic, align 4{{$}}258// AMDGCNSPIRV: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("device") monotonic, align 4{{$}}259 260// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("workgroup") monotonic, align 4{{$}}261 262// GCN: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("cluster") monotonic, align 4{{$}}263// AMDGCNSPIRV: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("workgroup") monotonic, align 4{{$}}264 265// GCN: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("wavefront") monotonic, align 4{{$}}266// AMDGCNSPIRV: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("subgroup") monotonic, align 4{{$}}267 268// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src syncscope("singlethread") monotonic, align 4{{$}}269// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}}270#if !defined(__SPIRV__)271void test_ds_faddf(local float *out, float src) {272#else273 void test_ds_faddf(__attribute__((address_space(3))) float *out, float src) {274#endif275 276 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM, false);277 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM, true);278 279 // Test all orders.280 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_CONSUME, __MEMORY_SCOPE_SYSTEM, false);281 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_SYSTEM, false);282 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_RELEASE, __MEMORY_SCOPE_SYSTEM, false);283 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_ACQ_REL, __MEMORY_SCOPE_SYSTEM, false);284 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_SYSTEM, false);285 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_SYSTEM, false); // invalid286 287 // Test all syncscopes.288 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE, false);289 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_WRKGRP, false);290 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_CLUSTR, false);291 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_WVFRNT, false);292 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_SINGLE, false);293 *out = __builtin_amdgcn_ds_faddf(out, src, __ATOMIC_RELAXED, INVALID_MEMORY_SCOPE, false); // invalid294}295 296// CHECK-LABEL: @test_ds_fmin297// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src monotonic, align 4{{$}}298// CHECK: atomicrmw volatile fmin ptr addrspace(3) %out, float %src monotonic, align 4{{$}}299 300// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src acquire, align 4{{$}}301// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src acquire, align 4{{$}}302// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src release, align 4{{$}}303// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src acq_rel, align 4{{$}}304// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src seq_cst, align 4{{$}}305// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src seq_cst, align 4{{$}}306 307// GCN: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("agent") monotonic, align 4{{$}}308// AMDGCNSPIRV: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("device") monotonic, align 4{{$}}309 310// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("workgroup") monotonic, align 4{{$}}311 312// GCN: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("cluster") monotonic, align 4{{$}}313// AMDGCNSPIRV: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("workgroup") monotonic, align 4{{$}}314 315// GCN: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("wavefront") monotonic, align 4{{$}}316// AMDGCNSPIRV: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("subgroup") monotonic, align 4{{$}}317 318// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src syncscope("singlethread") monotonic, align 4{{$}}319// CHECK: atomicrmw fmin ptr addrspace(3) %out, float %src monotonic, align 4{{$}}320 321#if !defined(__SPIRV__)322void test_ds_fminf(local float *out, float src) {323#else324void test_ds_fminf(__attribute__((address_space(3))) float *out, float src) {325#endif326 *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, false);327 *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, true);328 329 // Test all orders.330 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_CONSUME, __MEMORY_SCOPE_SYSTEM, false);331 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_SYSTEM, false);332 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELEASE, __MEMORY_SCOPE_SYSTEM, false);333 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_ACQ_REL, __MEMORY_SCOPE_SYSTEM, false);334 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_SYSTEM, false);335 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_SYSTEM, false); // invalid336 337 // Test all syncscopes.338 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE, false);339 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_WRKGRP, false);340 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_CLUSTR, false);341 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_WVFRNT, false);342 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_SINGLE, false);343 *out = __builtin_amdgcn_ds_fminf(out, src, __ATOMIC_RELAXED, INVALID_MEMORY_SCOPE, false); // invalid344}345 346// CHECK-LABEL: @test_ds_fmax347// CHECK: atomicrmw fmax ptr addrspace(3) %out, float %src monotonic, align 4{{$}}348// CHECK: atomicrmw volatile fmax ptr addrspace(3) %out, float %src monotonic, align 4{{$}}349 350// CHECK: atomicrmw fmax ptr addrspace(3) %out, float %src acquire, align 4{{$}}351// CHECK: atomicrmw fmax ptr addrspace(3) %out, float %src acquire, align 4{{$}}352// CHECK: atomicrmw fmax ptr addrspace(3) %out, float %src release, align 4{{$}}353// CHECK: atomicrmw fmax ptr addrspace(3) %out, float %src acq_rel, align 4{{$}}354// CHECK: atomicrmw fmax ptr addrspace(3) %out, float %src seq_cst, align 4{{$}}355// CHECK: atomicrmw fmax ptr addrspace(3) %out, float %src seq_cst, align 4{{$}}356 357// GCN: atomicrmw fmax ptr addrspace(3) %out, float %src syncscope("agent") monotonic, align 4{{$}}358// AMDGCNSPIRV: atomicrmw fmax ptr addrspace(3) %out, float %src syncscope("device") monotonic, align 4{{$}}359 360// CHECK: atomicrmw fmax ptr addrspace(3) %out, float %src syncscope("workgroup") monotonic, align 4{{$}}361 362// GCN: atomicrmw fmax ptr addrspace(3) %out, float %src syncscope("cluster") monotonic, align 4{{$}}363// AMDGCNSPIRV: atomicrmw fmax ptr addrspace(3) %out, float %src syncscope("workgroup") monotonic, align 4{{$}}364 365// GCN: atomicrmw fmax ptr addrspace(3) %out, float %src syncscope("wavefront") monotonic, align 4{{$}}366// AMDGCNSPIRV: atomicrmw fmax ptr addrspace(3) %out, float %src syncscope("subgroup") monotonic, align 4{{$}}367 368// CHECK: atomicrmw fmax ptr addrspace(3) %out, float %src syncscope("singlethread") monotonic, align 4{{$}}369// CHECK: atomicrmw fmax ptr addrspace(3) %out, float %src monotonic, align 4{{$}}370 371#if !defined(__SPIRV__)372void test_ds_fmaxf(local float *out, float src) {373#else374void test_ds_fmaxf(__attribute__((address_space(3))) float *out, float src) {375#endif376 *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false);377 *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, true);378 379 // Test all orders.380 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_CONSUME, __MEMORY_SCOPE_SYSTEM, false);381 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_SYSTEM, false);382 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELEASE, __MEMORY_SCOPE_SYSTEM, false);383 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_ACQ_REL, __MEMORY_SCOPE_SYSTEM, false);384 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_SYSTEM, false);385 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_SYSTEM, false); // invalid386 387 // Test all syncscopes.388 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE, false);389 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_WRKGRP, false);390 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_CLUSTR, false);391 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_WVFRNT, false);392 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __MEMORY_SCOPE_SINGLE, false);393 *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, INVALID_MEMORY_SCOPE, false); // invalid394}395 396// CHECK-LABEL: @test_s_memtime397// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.s.memtime()398void test_s_memtime(global ulong* out)399{400 *out = __builtin_amdgcn_s_memtime();401}402 403// CHECK-LABEL: @test_perm404// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.perm(i32 %a, i32 %b, i32 %s)405void test_perm(global uint* out, uint a, uint b, uint s)406{407 *out = __builtin_amdgcn_perm(a, b, s);408}409 410// CHECK-LABEL: @test_groupstaticsize411// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.groupstaticsize()412void test_groupstaticsize(global uint* out)413{414 *out = __builtin_amdgcn_groupstaticsize();415}416