106 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 52 3// REQUIRES: amdgpu-registered-target4// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck %s5// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1251 -emit-llvm -o - %s | FileCheck %s6 7typedef int v2i __attribute__((ext_vector_type(2)));8typedef int v4i __attribute__((ext_vector_type(4)));9 10// CHECK-LABEL: define dso_local void @test_amdgcn_cooperative_atomic_store_32x4B(11// CHECK-SAME: ptr addrspace(1) noundef writeonly captures(none) [[GADDR:%.*]], i32 noundef [[VAL:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {12// CHECK-NEXT: [[ENTRY:.*:]]13// CHECK-NEXT: tail call void @llvm.amdgcn.cooperative.atomic.store.32x4B.p1(ptr addrspace(1) [[GADDR]], i32 [[VAL]], i32 0, metadata [[META4:![0-9]+]])14// CHECK-NEXT: ret void15//16void test_amdgcn_cooperative_atomic_store_32x4B(global int* gaddr, int val)17{18 __builtin_amdgcn_cooperative_atomic_store_32x4B(gaddr, val, __ATOMIC_RELAXED, "agent");19}20 21// CHECK-LABEL: define dso_local i32 @test_amdgcn_cooperative_atomic_load_32x4B(22// CHECK-SAME: ptr noundef readonly captures(none) [[ADDR:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] {23// CHECK-NEXT: [[ENTRY:.*:]]24// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.cooperative.atomic.load.32x4B.p0(ptr [[ADDR]], i32 0, metadata [[META5:![0-9]+]])25// CHECK-NEXT: ret i32 [[TMP0]]26//27int test_amdgcn_cooperative_atomic_load_32x4B(int* addr)28{29 return __builtin_amdgcn_cooperative_atomic_load_32x4B(addr, __ATOMIC_RELAXED, "");30}31 32// CHECK-LABEL: define dso_local void @test_amdgcn_cooperative_atomic_store_16x8B(33// CHECK-SAME: ptr addrspace(1) noundef writeonly captures(none) [[GADDR:%.*]], <2 x i32> noundef [[VAL:%.*]]) local_unnamed_addr #[[ATTR0]] {34// CHECK-NEXT: [[ENTRY:.*:]]35// CHECK-NEXT: tail call void @llvm.amdgcn.cooperative.atomic.store.16x8B.p1(ptr addrspace(1) [[GADDR]], <2 x i32> [[VAL]], i32 0, metadata [[META5]])36// CHECK-NEXT: ret void37//38void test_amdgcn_cooperative_atomic_store_16x8B(global v2i* gaddr, v2i val)39{40 __builtin_amdgcn_cooperative_atomic_store_16x8B(gaddr, val, __ATOMIC_RELAXED, "");41}42 43// CHECK-LABEL: define dso_local <2 x i32> @test_amdgcn_cooperative_atomic_load_16x8B(44// CHECK-SAME: ptr addrspace(1) noundef readonly captures(none) [[GADDR:%.*]]) local_unnamed_addr #[[ATTR2]] {45// CHECK-NEXT: [[ENTRY:.*:]]46// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.cooperative.atomic.load.16x8B.p1(ptr addrspace(1) [[GADDR]], i32 0, metadata [[META6:![0-9]+]])47// CHECK-NEXT: ret <2 x i32> [[TMP0]]48//49v2i test_amdgcn_cooperative_atomic_load_16x8B(global v2i* gaddr)50{51 return __builtin_amdgcn_cooperative_atomic_load_16x8B(gaddr, __ATOMIC_RELAXED, "workgroup");52}53 54// CHECK-LABEL: define dso_local void @test_amdgcn_cooperative_atomic_store_8x16B(55// CHECK-SAME: ptr addrspace(1) noundef writeonly captures(none) [[GADDR:%.*]], <4 x i32> noundef [[VAL:%.*]]) local_unnamed_addr #[[ATTR0]] {56// CHECK-NEXT: [[ENTRY:.*:]]57// CHECK-NEXT: tail call void @llvm.amdgcn.cooperative.atomic.store.8x16B.p1(ptr addrspace(1) [[GADDR]], <4 x i32> [[VAL]], i32 0, metadata [[META7:![0-9]+]])58// CHECK-NEXT: ret void59//60void test_amdgcn_cooperative_atomic_store_8x16B(global v4i* gaddr, v4i val)61{62 __builtin_amdgcn_cooperative_atomic_store_8x16B(gaddr, val, __ATOMIC_RELAXED, "singlethread");63}64 65// CHECK-LABEL: define dso_local <4 x i32> @test_amdgcn_cooperative_atomic_load_8x16B(66// CHECK-SAME: ptr addrspace(1) noundef readonly captures(none) [[GADDR:%.*]]) local_unnamed_addr #[[ATTR2]] {67// CHECK-NEXT: [[ENTRY:.*:]]68// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.cooperative.atomic.load.8x16B.p1(ptr addrspace(1) [[GADDR]], i32 0, metadata [[META4]])69// CHECK-NEXT: ret <4 x i32> [[TMP0]]70//71v4i test_amdgcn_cooperative_atomic_load_8x16B(global v4i* gaddr)72{73 return __builtin_amdgcn_cooperative_atomic_load_8x16B(gaddr, __ATOMIC_RELAXED, "agent");74}75 76// CHECK-LABEL: define dso_local void @test_amdgcn_cooperative_atomic_store_32x4B_truncated(77// CHECK-SAME: ptr addrspace(1) noundef writeonly captures(none) [[GADDR:%.*]], i64 noundef [[VAL:%.*]]) local_unnamed_addr #[[ATTR0]] {78// CHECK-NEXT: [[ENTRY:.*:]]79// CHECK-NEXT: [[CONV:%.*]] = trunc i64 [[VAL]] to i3280// CHECK-NEXT: tail call void @llvm.amdgcn.cooperative.atomic.store.32x4B.p1(ptr addrspace(1) [[GADDR]], i32 [[CONV]], i32 0, metadata [[META4]])81// CHECK-NEXT: ret void82//83void test_amdgcn_cooperative_atomic_store_32x4B_truncated(global int* gaddr, long val)84{85 __builtin_amdgcn_cooperative_atomic_store_32x4B(gaddr, val, __ATOMIC_RELAXED, "agent");86}87 88// CHECK-LABEL: define dso_local void @test_amdgcn_cooperative_atomic_store_32x4B_extended(89// CHECK-SAME: ptr addrspace(1) noundef writeonly captures(none) [[GADDR:%.*]], i8 noundef signext [[VAL:%.*]]) local_unnamed_addr #[[ATTR0]] {90// CHECK-NEXT: [[ENTRY:.*:]]91// CHECK-NEXT: [[CONV:%.*]] = sext i8 [[VAL]] to i3292// CHECK-NEXT: tail call void @llvm.amdgcn.cooperative.atomic.store.32x4B.p1(ptr addrspace(1) [[GADDR]], i32 [[CONV]], i32 0, metadata [[META4]])93// CHECK-NEXT: ret void94//95void test_amdgcn_cooperative_atomic_store_32x4B_extended(global int* gaddr, char val)96{97 __builtin_amdgcn_cooperative_atomic_store_32x4B(gaddr, val, __ATOMIC_RELAXED, "agent");98}99 100//.101// CHECK: [[META4]] = !{!"agent"}102// CHECK: [[META5]] = !{!""}103// CHECK: [[META6]] = !{!"workgroup"}104// CHECK: [[META7]] = !{!"singlethread"}105//.106