1 // REQUIRES
: amdgpu-registered-target
3 // RUN
: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -verify -emit-llvm -o - %s
4 // RUN
: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -verify -emit-llvm -o - %s
6 typedef unsigned int uint
;
7 typedef half __attribute__
((ext_vector_type(2))) half2
;
8 typedef short __attribute__
((ext_vector_type(2))) short2
;
9 typedef unsigned short __attribute__
((ext_vector_type(2))) ushort2
;
11 #pragma OPENCL EXTENSION cl_khr_fp16
: enable
12 kernel void builtins_amdgcn_dl_insts_err
(
13 global float
*fOut
, global int
*siOut
, global uint
*uiOut
,
14 global short
*sOut
, global int
*iOut
, global half
*hOut
,
15 half2 v2hA
, half2 v2hB
, float fC
, half hC
,
16 short2 v2ssA
, short2 v2ssB
, short sC
, int siA
, int siB
, int siC
,
17 ushort2 v2usA
, ushort2 v2usB
, uint uiA
, uint uiB
, uint uiC
,
18 int A
, int B
, int C
) {
19 fOut
[0] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC, false); // expected-error {{'__builtin_amdgcn_fdot2' needs target feature dot10-insts}}
20 fOut[1] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC, true); // expected-error {{'__builtin_amdgcn_fdot2' needs target feature dot10-insts}}
22 hOut[0] = __builtin_amdgcn_fdot2_f16_f16
(v2hA, v2hB
, hC
); // expected-error {{'__builtin_amdgcn_fdot2_f16_f16' needs target feature dot9-insts}}
24 sOut
[0] = __builtin_amdgcn_fdot2_bf16_bf16(v2ssA, v2ssB, sC); // expected-error {{'__builtin_amdgcn_fdot2_bf16_bf16' needs target feature dot9-insts}}
26 fOut[3] = __builtin_amdgcn_fdot2_f32_bf16(v2ssA, v2ssB, fC, false); // expected-error {{'__builtin_amdgcn_fdot2_f32_bf16' needs target feature dot9-insts}}
27 fOut[4] = __builtin_amdgcn_fdot2_f32_bf16(v2ssA, v2ssB, fC, true); // expected-error {{'__builtin_amdgcn_fdot2_f32_bf16' needs target feature dot9-insts}}
29 siOut[0] = __builtin_amdgcn_sdot2
(v2ssA, v2ssB
, siC
, false
); // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dot2-insts}}
30 siOut
[1] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC, true); // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dot2-insts}}
32 uiOut[0] = __builtin_amdgcn_udot2(v2usA, v2usB, uiC, false); // expected-error {{'__builtin_amdgcn_udot2' needs target feature dot2-insts}}
33 uiOut[1] = __builtin_amdgcn_udot2
(v2usA, v2usB
, uiC
, true
); // expected-error {{'__builtin_amdgcn_udot2' needs target feature dot2-insts}}
35 siOut
[2] = __builtin_amdgcn_sdot4(siA, siB, siC, false); // expected-error {{'__builtin_amdgcn_sdot4' needs target feature dot1-insts}}
36 siOut[3] = __builtin_amdgcn_sdot4(siA, siB, siC, true); // expected-error {{'__builtin_amdgcn_sdot4' needs target feature dot1-insts}}
38 uiOut[2] = __builtin_amdgcn_udot4
(uiA, uiB
, uiC
, false
); // expected-error {{'__builtin_amdgcn_udot4' needs target feature dot7-insts}}
39 uiOut
[3] = __builtin_amdgcn_udot4(uiA, uiB, uiC, true); // expected-error {{'__builtin_amdgcn_udot4' needs target feature dot7-insts}}
41 iOut[0] = __builtin_amdgcn_sudot4(true, A, false, B, C, false); // expected-error {{'__builtin_amdgcn_sudot4' needs target feature dot8-insts}}
42 iOut[1] = __builtin_amdgcn_sudot4(false, A, true, B, C, true); // expected-error {{'__builtin_amdgcn_sudot4' needs target feature dot8-insts}}
44 siOut[4] = __builtin_amdgcn_sdot8(siA, siB, siC, false); // expected-error {{'__builtin_amdgcn_sdot8' needs target feature dot1-insts}}
45 siOut[5] = __builtin_amdgcn_sdot8(siA, siB, siC, true); // expected-error {{'__builtin_amdgcn_sdot8' needs target feature dot1-insts}}
47 uiOut[4] = __builtin_amdgcn_udot8(uiA, uiB, uiC, false); // expected-error {{'__builtin_amdgcn_udot8' needs target feature dot7-insts}}
48 uiOut[5] = __builtin_amdgcn_udot8(uiA, uiB, uiC, true); // expected-error {{'__builtin_amdgcn_udot8' needs target feature dot7-insts}}
50 iOut[3] = __builtin_amdgcn_sudot8
(false, A
, true
, B
, C
, false
); // expected-error {{'__builtin_amdgcn_sudot8' needs target feature dot8-insts}}
51 iOut
[4] = __builtin_amdgcn_sudot8(true, A, false, B, C, true); // expected-error {{'__builtin_amdgcn_sudot8' needs target feature dot8-insts}}
53 fOut[5] = __builtin_amdgcn_dot4_f32_fp8_bf8(uiA, uiB, fC); // expected-error {{'__builtin_amdgcn_dot4_f32_fp8_bf8' needs target feature dot11-insts}}
54 fOut[6] = __builtin_amdgcn_dot4_f32_bf8_fp8(uiA, uiB, fC); // expected-error {{'__builtin_amdgcn_dot4_f32_bf8_fp8' needs target feature dot11-insts}}
55 fOut[7] = __builtin_amdgcn_dot4_f32_fp8_fp8(uiA, uiB, fC); // expected-error {{'__builtin_amdgcn_dot4_f32_fp8_fp8' needs target feature dot11-insts}}
56 fOut[8] = __builtin_amdgcn_dot4_f32_bf8_bf8(uiA, uiB, fC); // expected-error {{'__builtin_amdgcn_dot4_f32_bf8_bf8' needs target feature dot11-insts}}