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