131 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py2// REQUIRES: amdgpu-registered-target3// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX12504 5typedef int v2i __attribute__((ext_vector_type(2)));6typedef int v3i __attribute__((ext_vector_type(3)));7typedef int v4i __attribute__((ext_vector_type(4)));8typedef short v8s __attribute__((ext_vector_type(8)));9typedef half v8h __attribute__((ext_vector_type(8)));10typedef __bf16 v8y __attribute__((ext_vector_type(8)));11 12// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_tr4_b64_v2i32(13// CHECK-GFX1250-NEXT: entry:14// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.global.load.tr4.b64.v2i32(ptr addrspace(1) [[INPTR:%.*]])15// CHECK-GFX1250-NEXT: ret <2 x i32> [[TMP0]]16//17v2i test_amdgcn_global_load_tr4_b64_v2i32(global v2i* inptr)18{19 return __builtin_amdgcn_global_load_tr4_b64_v2i32(inptr);20}21 22// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_tr8_b64_v2i32(23// CHECK-GFX1250-NEXT: entry:24// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.global.load.tr.b64.v2i32(ptr addrspace(1) [[INPTR:%.*]])25// CHECK-GFX1250-NEXT: ret <2 x i32> [[TMP0]]26//27v2i test_amdgcn_global_load_tr8_b64_v2i32(global v2i* inptr)28{29 return __builtin_amdgcn_global_load_tr8_b64_v2i32(inptr);30}31 32// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_tr6_b96_v3i32(33// CHECK-GFX1250-NEXT: entry:34// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <3 x i32> @llvm.amdgcn.global.load.tr6.b96.v3i32(ptr addrspace(1) [[INPTR:%.*]])35// CHECK-GFX1250-NEXT: ret <3 x i32> [[TMP0]]36//37v3i test_amdgcn_global_load_tr6_b96_v3i32(global v3i* inptr)38{39 return __builtin_amdgcn_global_load_tr6_b96_v3i32(inptr);40}41 42// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_tr16_b128_v8i16(43// CHECK-GFX1250-NEXT: entry:44// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <8 x i16> @llvm.amdgcn.global.load.tr.b128.v8i16(ptr addrspace(1) [[INPTR:%.*]])45// CHECK-GFX1250-NEXT: ret <8 x i16> [[TMP0]]46//47v8s test_amdgcn_global_load_tr16_b128_v8i16(global v8s* inptr)48{49 return __builtin_amdgcn_global_load_tr16_b128_v8i16(inptr);50}51 52// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_tr16_b128_v8f16(53// CHECK-GFX1250-NEXT: entry:54// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <8 x half> @llvm.amdgcn.global.load.tr.b128.v8f16(ptr addrspace(1) [[INPTR:%.*]])55// CHECK-GFX1250-NEXT: ret <8 x half> [[TMP0]]56//57v8h test_amdgcn_global_load_tr16_b128_v8f16(global v8h* inptr)58{59 return __builtin_amdgcn_global_load_tr16_b128_v8f16(inptr);60}61 62// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_tr16_b128_v8bf16(63// CHECK-GFX1250-NEXT: entry:64// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <8 x bfloat> @llvm.amdgcn.global.load.tr.b128.v8bf16(ptr addrspace(1) [[INPTR:%.*]])65// CHECK-GFX1250-NEXT: ret <8 x bfloat> [[TMP0]]66//67v8y test_amdgcn_global_load_tr16_b128_v8bf16(global v8y* inptr)68{69 return __builtin_amdgcn_global_load_tr16_b128_v8bf16(inptr);70}71 72// CHECK-GFX1250-LABEL: @test_amdgcn_ds_load_tr4_b64_v2i32(73// CHECK-GFX1250-NEXT: entry:74// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.ds.load.tr4.b64.v2i32(ptr addrspace(3) [[INPTR:%.*]])75// CHECK-GFX1250-NEXT: ret <2 x i32> [[TMP0]]76//77v2i test_amdgcn_ds_load_tr4_b64_v2i32(local v2i* inptr)78{79 return __builtin_amdgcn_ds_load_tr4_b64_v2i32(inptr);80}81 82// CHECK-GFX1250-LABEL: @test_amdgcn_ds_load_tr8_b64_v2i32(83// CHECK-GFX1250-NEXT: entry:84// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.ds.load.tr8.b64.v2i32(ptr addrspace(3) [[INPTR:%.*]])85// CHECK-GFX1250-NEXT: ret <2 x i32> [[TMP0]]86//87v2i test_amdgcn_ds_load_tr8_b64_v2i32(local v2i* inptr)88{89 return __builtin_amdgcn_ds_load_tr8_b64_v2i32(inptr);90}91 92// CHECK-GFX1250-LABEL: @test_amdgcn_ds_load_tr6_b96_v3i32(93// CHECK-GFX1250-NEXT: entry:94// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <3 x i32> @llvm.amdgcn.ds.load.tr6.b96.v3i32(ptr addrspace(3) [[INPTR:%.*]])95// CHECK-GFX1250-NEXT: ret <3 x i32> [[TMP0]]96//97v3i test_amdgcn_ds_load_tr6_b96_v3i32(local v3i* inptr)98{99 return __builtin_amdgcn_ds_load_tr6_b96_v3i32(inptr);100}101 102// CHECK-GFX1250-LABEL: @test_amdgcn_ds_load_tr16_b128_v8i16(103// CHECK-GFX1250-NEXT: entry:104// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <8 x i16> @llvm.amdgcn.ds.load.tr16.b128.v8i16(ptr addrspace(3) [[INPTR:%.*]])105// CHECK-GFX1250-NEXT: ret <8 x i16> [[TMP0]]106//107v8s test_amdgcn_ds_load_tr16_b128_v8i16(local v8s* inptr)108{109 return __builtin_amdgcn_ds_load_tr16_b128_v8i16(inptr);110}111 112// CHECK-GFX1250-LABEL: @test_amdgcn_ds_load_tr16_b128_v8f16(113// CHECK-GFX1250-NEXT: entry:114// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) [[INPTR:%.*]])115// CHECK-GFX1250-NEXT: ret <8 x half> [[TMP0]]116//117v8h test_amdgcn_ds_load_tr16_b128_v8f16(local v8h* inptr)118{119 return __builtin_amdgcn_ds_load_tr16_b128_v8f16(inptr);120}121 122// CHECK-GFX1250-LABEL: @test_amdgcn_ds_load_tr16_b128_v8bf16(123// CHECK-GFX1250-NEXT: entry:124// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <8 x bfloat> @llvm.amdgcn.ds.load.tr16.b128.v8bf16(ptr addrspace(3) [[INPTR:%.*]])125// CHECK-GFX1250-NEXT: ret <8 x bfloat> [[TMP0]]126//127v8y test_amdgcn_ds_load_tr16_b128_v8bf16(local v8y* inptr)128{129 return __builtin_amdgcn_ds_load_tr16_b128_v8bf16(inptr);130}131