112 lines · plain
1// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1200 \2// RUN: %s -emit-llvm -o - | FileCheck %s3 4// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1200 \5// RUN: -S -o - %s | FileCheck -check-prefix=GFX12 %s6 7// REQUIRES: amdgpu-registered-target8 9typedef half __attribute__((ext_vector_type(2))) half2;10typedef short __attribute__((ext_vector_type(2))) short2;11 12// CHECK-LABEL: test_local_add_2bf1613// CHECK: [[BC0:%.+]] = bitcast <2 x i16> {{.+}} to <2 x bfloat>14// CHECK-NEXT: [[RMW:%.+]] = atomicrmw fadd ptr addrspace(3) %{{.+}}, <2 x bfloat> [[BC0]] syncscope("agent") monotonic, align 415// CHECK-NEXT: bitcast <2 x bfloat> [[RMW]] to <2 x i16>16 17// GFX12-LABEL: test_local_add_2bf1618// GFX12: ds_pk_add_rtn_bf1619short2 test_local_add_2bf16(__local short2 *addr, short2 x) {20 return __builtin_amdgcn_ds_atomic_fadd_v2bf16(addr, x);21}22 23// CHECK-LABEL: test_local_add_2bf16_noret24// CHECK: [[BC0:%.+]] = bitcast <2 x i16> {{.+}} to <2 x bfloat>25// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr addrspace(3) %{{.+}}, <2 x bfloat> [[BC0]] syncscope("agent") monotonic, align 426// CHECK-NEXT: bitcast <2 x bfloat> [[RMW]] to <2 x i16>27 28// GFX12-LABEL: test_local_add_2bf16_noret29// GFX12: ds_pk_add_bf1630void test_local_add_2bf16_noret(__local short2 *addr, short2 x) {31 __builtin_amdgcn_ds_atomic_fadd_v2bf16(addr, x);32}33 34// CHECK-LABEL: test_local_add_2f1635// CHECK: = atomicrmw fadd ptr addrspace(3) %{{.+}}, <2 x half> %{{.+}} monotonic, align 436// GFX12-LABEL: test_local_add_2f1637// GFX12: ds_pk_add_rtn_f1638half2 test_local_add_2f16(__local half2 *addr, half2 x) {39 return __builtin_amdgcn_ds_atomic_fadd_v2f16(addr, x);40}41 42// CHECK-LABEL: test_local_add_2f16_noret43// CHECK: = atomicrmw fadd ptr addrspace(3) %{{.+}}, <2 x half> %{{.+}} monotonic, align 444// GFX12-LABEL: test_local_add_2f16_noret45// GFX12: ds_pk_add_f1646void test_local_add_2f16_noret(__local half2 *addr, half2 x) {47 __builtin_amdgcn_ds_atomic_fadd_v2f16(addr, x);48}49 50// CHECK-LABEL: test_flat_add_2f1651// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr %{{.+}}, <2 x half> %{{.+}} syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}52 53// GFX12-LABEL: test_flat_add_2f1654// GFX12: flat_atomic_pk_add_f1655half2 test_flat_add_2f16(__generic half2 *addr, half2 x) {56 return __builtin_amdgcn_flat_atomic_fadd_v2f16(addr, x);57}58 59// CHECK-LABEL: test_flat_add_2bf1660// CHECK: [[BC:%.+]] = bitcast <2 x i16> %{{.+}} to <2 x bfloat>61// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr %{{.+}}, <2 x bfloat> [[BC]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}62// CHECK: bitcast <2 x bfloat> [[RMW]] to <2 x i16>63 64// GFX12-LABEL: test_flat_add_2bf1665// GFX12: flat_atomic_pk_add_bf1666short2 test_flat_add_2bf16(__generic short2 *addr, short2 x) {67 return __builtin_amdgcn_flat_atomic_fadd_v2bf16(addr, x);68}69 70// CHECK-LABEL: test_global_add_half271// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr addrspace(1) %{{.+}}, <2 x half> %{{.+}} syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}72 73// GFX12-LABEL: test_global_add_half274// GFX12: global_atomic_pk_add_f16 v2, v[{{[0-9]+}}:{{[0-9]+}}], v{{[0-9]+}}, off th:TH_ATOMIC_RETURN75void test_global_add_half2(__global half2 *addr, half2 x) {76 half2 *rtn;77 *rtn = __builtin_amdgcn_global_atomic_fadd_v2f16(addr, x);78}79 80// CHECK-LABEL: test_global_add_half2_noret81// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr addrspace(1) %{{.+}}, <2 x half> %{{.+}} syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}82 83// GFX12-LABEL: test_global_add_half2_noret84// GFX12: global_atomic_pk_add_f16 v[0:1], v2, off85void test_global_add_half2_noret(__global half2 *addr, half2 x) {86 __builtin_amdgcn_global_atomic_fadd_v2f16(addr, x);87}88 89// CHECK-LABEL: test_global_add_2bf1690// CHECK: [[BC:%.+]] = bitcast <2 x i16> %{{.+}} to <2 x bfloat>91// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr addrspace(1) %{{.+}}, <2 x bfloat> [[BC]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}92// CHECK: bitcast <2 x bfloat> [[RMW]] to <2 x i16>93 94 95// GFX12-LABEL: test_global_add_2bf1696// GFX12: global_atomic_pk_add_bf16 v2, v[{{[0-9]+}}:{{[0-9]+}}], v{{[0-9]+}}, off th:TH_ATOMIC_RETURN97void test_global_add_2bf16(__global short2 *addr, short2 x) {98 short2 *rtn;99 *rtn = __builtin_amdgcn_global_atomic_fadd_v2bf16(addr, x);100}101 102// CHECK-LABEL: test_global_add_2bf16_noret103// CHECK: [[BC:%.+]] = bitcast <2 x i16> %{{.+}} to <2 x bfloat>104// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr addrspace(1) %{{.+}}, <2 x bfloat> [[BC]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}105// CHECK: bitcast <2 x bfloat> [[RMW]] to <2 x i16>106 107// GFX12-LABEL: test_global_add_2bf16_noret108// GFX12: global_atomic_pk_add_bf16 v[0:1], v2, off109void test_global_add_2bf16_noret(__global short2 *addr, short2 x) {110 __builtin_amdgcn_global_atomic_fadd_v2bf16(addr, x);111}112