Skip to content

Commit a102342

Browse files
shiltianrampitec
shiltian
and
rampitec
authored
[AMDGPU] Add support for v_sin_bf16 on gfx1250 (#149241)
Co-authored-by: Mekhanoshin, Stanislav <[email protected]>
1 parent a6b5ece commit a102342

23 files changed

+819
-0
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -674,6 +674,7 @@ TARGET_BUILTIN(__builtin_amdgcn_rcp_bf16, "yy", "nc", "bf16-trans-insts")
674674
TARGET_BUILTIN(__builtin_amdgcn_rsq_bf16, "yy", "nc", "bf16-trans-insts")
675675
TARGET_BUILTIN(__builtin_amdgcn_log_bf16, "yy", "nc", "bf16-trans-insts")
676676
TARGET_BUILTIN(__builtin_amdgcn_exp2_bf16, "yy", "nc", "bf16-trans-insts")
677+
TARGET_BUILTIN(__builtin_amdgcn_sin_bf16, "yy", "nc", "bf16-trans-insts")
677678

678679
TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_fp8, "hiIi", "nc", "gfx1250-insts")
679680
TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_bf8, "hiIi", "nc", "gfx1250-insts")

‎clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -429,6 +429,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
429429
Intrinsic::amdgcn_rsq_clamp);
430430
case AMDGPU::BI__builtin_amdgcn_sinf:
431431
case AMDGPU::BI__builtin_amdgcn_sinh:
432+
case AMDGPU::BI__builtin_amdgcn_sin_bf16:
432433
return emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::amdgcn_sin);
433434
case AMDGPU::BI__builtin_amdgcn_cosf:
434435
case AMDGPU::BI__builtin_amdgcn_cosh:

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -137,6 +137,25 @@ void test_exp2_bf16(global __bf16* out, __bf16 a)
137137
*out = __builtin_amdgcn_exp2_bf16(a);
138138
}
139139

140+
// CHECK-LABEL: @test_sin_bf16(
141+
// CHECK-NEXT: entry:
142+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
143+
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca bfloat, align 2, addrspace(5)
144+
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
145+
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
146+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
147+
// CHECK-NEXT: store bfloat [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 2
148+
// CHECK-NEXT: [[TMP0:%.*]] = load bfloat, ptr [[A_ADDR_ASCAST]], align 2
149+
// CHECK-NEXT: [[TMP1:%.*]] = call bfloat @llvm.amdgcn.sin.bf16(bfloat [[TMP0]])
150+
// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
151+
// CHECK-NEXT: store bfloat [[TMP1]], ptr addrspace(1) [[TMP2]], align 2
152+
// CHECK-NEXT: ret void
153+
//
154+
void test_sin_bf16(global __bf16* out, __bf16 a)
155+
{
156+
*out = __builtin_amdgcn_sin_bf16(a);
157+
}
158+
140159
// CHECK-LABEL: @test_cvt_f16_fp8(
141160
// CHECK-NEXT: entry:
142161
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)

‎llvm/lib/Target/AMDGPU/VOP1Instructions.td

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -534,6 +534,7 @@ defm V_SQRT_BF16 : VOP1Inst_t16 <"v_sqrt_bf16", VOP_BF16_BF16, any_amdgcn_sqrt>;
534534
defm V_RSQ_BF16 : VOP1Inst_t16 <"v_rsq_bf16", VOP_BF16_BF16, AMDGPUrsq>;
535535
defm V_LOG_BF16 : VOP1Inst_t16 <"v_log_bf16", VOP_BF16_BF16, AMDGPUlogf16>;
536536
defm V_EXP_BF16 : VOP1Inst_t16 <"v_exp_bf16", VOP_BF16_BF16, AMDGPUexpf16>;
537+
defm V_SIN_BF16 : VOP1Inst_t16 <"v_sin_bf16", VOP_BF16_BF16, AMDGPUsin>;
537538
}
538539
} // End TRANS = 1, SchedRW = [WriteTrans32]
539540
defm V_FREXP_MANT_F16 : VOP1Inst_t16 <"v_frexp_mant_f16", VOP_F16_F16, int_amdgcn_frexp_mant>;
@@ -1147,6 +1148,7 @@ defm V_SQRT_BF16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x07a>;
11471148
defm V_RSQ_BF16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x07b>;
11481149
defm V_LOG_BF16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x07c>;
11491150
defm V_EXP_BF16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x07d>;
1151+
defm V_SIN_BF16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x07e>;
11501152

11511153
//===----------------------------------------------------------------------===//
11521154
// GFX10.

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sin.bf16.ll

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck -check-prefixes=GCN %s
2+
; xUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck -check-prefix=GCN %s
3+
4+
; FIXME: GlobalISel does not work with bf16
5+
6+
declare bfloat @llvm.amdgcn.sin.bf16(bfloat) #0
7+
8+
; GCN-LABEL: {{^}}sin_bf16:
9+
; GCN: v_sin_bf16_e32 {{v[0-9]+}}, {{s[0-9]+}}
10+
define amdgpu_kernel void @sin_bf16(ptr addrspace(1) %out, bfloat %src) #1 {
11+
%sin = call bfloat @llvm.amdgcn.sin.bf16(bfloat %src) #0
12+
store bfloat %sin, ptr addrspace(1) %out, align 2
13+
ret void
14+
}
15+
16+
; GCN-LABEL: {{^}}sin_bf16_constant_4
17+
; GCN: v_sin_bf16_e32 v0, 4.0
18+
define amdgpu_kernel void @sin_bf16_constant_4(ptr addrspace(1) %out) #1 {
19+
%sin = call bfloat @llvm.amdgcn.sin.bf16(bfloat 4.0) #0
20+
store bfloat %sin, ptr addrspace(1) %out, align 2
21+
ret void
22+
}
23+
24+
; GCN-LABEL: {{^}}sin_bf16_constant_100
25+
; GCN: v_sin_bf16_e32 {{v[0-9]+}}, 0x42c8
26+
define amdgpu_kernel void @sin_bf16_constant_100(ptr addrspace(1) %out) #1 {
27+
%sin = call bfloat @llvm.amdgcn.sin.bf16(bfloat 100.0) #0
28+
store bfloat %sin, ptr addrspace(1) %out, align 2
29+
ret void
30+
}
31+
32+
attributes #0 = { nounwind readnone }
33+
attributes #1 = { nounwind }

‎llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -298,6 +298,51 @@ v_exp_bf16 v5, src_scc
298298
v_exp_bf16 v127, 0x8000
299299
// GFX1250: v_exp_bf16_e32 v127, 0x8000 ; encoding: [0xff,0xfa,0xfe,0x7e,0x00,0x80,0x00,0x00]
300300

301+
v_sin_bf16 v5, v1
302+
// GFX1250: v_sin_bf16_e32 v5, v1 ; encoding: [0x01,0xfd,0x0a,0x7e]
303+
304+
v_sin_bf16 v5, v127
305+
// GFX1250: v_sin_bf16_e32 v5, v127 ; encoding: [0x7f,0xfd,0x0a,0x7e]
306+
307+
v_sin_bf16 v5, s1
308+
// GFX1250: v_sin_bf16_e32 v5, s1 ; encoding: [0x01,0xfc,0x0a,0x7e]
309+
310+
v_sin_bf16 v5, s105
311+
// GFX1250: v_sin_bf16_e32 v5, s105 ; encoding: [0x69,0xfc,0x0a,0x7e]
312+
313+
v_sin_bf16 v5, vcc_lo
314+
// GFX1250: v_sin_bf16_e32 v5, vcc_lo ; encoding: [0x6a,0xfc,0x0a,0x7e]
315+
316+
v_sin_bf16 v5, vcc_hi
317+
// GFX1250: v_sin_bf16_e32 v5, vcc_hi ; encoding: [0x6b,0xfc,0x0a,0x7e]
318+
319+
v_sin_bf16 v5, ttmp15
320+
// GFX1250: v_sin_bf16_e32 v5, ttmp15 ; encoding: [0x7b,0xfc,0x0a,0x7e]
321+
322+
v_sin_bf16 v5, m0
323+
// GFX1250: v_sin_bf16_e32 v5, m0 ; encoding: [0x7d,0xfc,0x0a,0x7e]
324+
325+
v_sin_bf16 v5, exec_lo
326+
// GFX1250: v_sin_bf16_e32 v5, exec_lo ; encoding: [0x7e,0xfc,0x0a,0x7e]
327+
328+
v_sin_bf16 v5, exec_hi
329+
// GFX1250: v_sin_bf16_e32 v5, exec_hi ; encoding: [0x7f,0xfc,0x0a,0x7e]
330+
331+
v_sin_bf16 v5, null
332+
// GFX1250: v_sin_bf16_e32 v5, null ; encoding: [0x7c,0xfc,0x0a,0x7e]
333+
334+
v_sin_bf16 v5, -1
335+
// GFX1250: v_sin_bf16_e32 v5, -1 ; encoding: [0xc1,0xfc,0x0a,0x7e]
336+
337+
v_sin_bf16 v5, 0.5
338+
// GFX1250: v_sin_bf16_e32 v5, 0.5 ; encoding: [0xf0,0xfc,0x0a,0x7e]
339+
340+
v_sin_bf16 v5, src_scc
341+
// GFX1250: v_sin_bf16_e32 v5, src_scc ; encoding: [0xfd,0xfc,0x0a,0x7e]
342+
343+
v_sin_bf16 v127, 0x8000
344+
// GFX1250: v_sin_bf16_e32 v127, 0x8000 ; encoding: [0xff,0xfc,0xfe,0x7e,0x00,0x80,0x00,0x00]
345+
301346
v_cvt_f32_bf16 v5, v1
302347
// GFX1250: v_cvt_f32_bf16_e32 v5, v1 ; encoding: [0x01,0xe5,0x0a,0x7e]
303348

‎llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s

Lines changed: 48 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -316,6 +316,54 @@ v_exp_bf16 v127, 0x8000
316316
v_exp_bf16 v5.h, v1.h
317317
// GFX1250: v_exp_bf16_e32 v5.h, v1.h ; encoding: [0x81,0xfb,0x0a,0x7f]
318318

319+
v_sin_bf16 v5, v1
320+
// GFX1250: v_sin_bf16_e32 v5, v1 ; encoding: [0x01,0xfd,0x0a,0x7e]
321+
322+
v_sin_bf16 v5, v127
323+
// GFX1250: v_sin_bf16_e32 v5, v127 ; encoding: [0x7f,0xfd,0x0a,0x7e]
324+
325+
v_sin_bf16 v5, s1
326+
// GFX1250: v_sin_bf16_e32 v5, s1 ; encoding: [0x01,0xfc,0x0a,0x7e]
327+
328+
v_sin_bf16 v5, s105
329+
// GFX1250: v_sin_bf16_e32 v5, s105 ; encoding: [0x69,0xfc,0x0a,0x7e]
330+
331+
v_sin_bf16 v5, vcc_lo
332+
// GFX1250: v_sin_bf16_e32 v5, vcc_lo ; encoding: [0x6a,0xfc,0x0a,0x7e]
333+
334+
v_sin_bf16 v5, vcc_hi
335+
// GFX1250: v_sin_bf16_e32 v5, vcc_hi ; encoding: [0x6b,0xfc,0x0a,0x7e]
336+
337+
v_sin_bf16 v5, ttmp15
338+
// GFX1250: v_sin_bf16_e32 v5, ttmp15 ; encoding: [0x7b,0xfc,0x0a,0x7e]
339+
340+
v_sin_bf16 v5, m0
341+
// GFX1250: v_sin_bf16_e32 v5, m0 ; encoding: [0x7d,0xfc,0x0a,0x7e]
342+
343+
v_sin_bf16 v5, exec_lo
344+
// GFX1250: v_sin_bf16_e32 v5, exec_lo ; encoding: [0x7e,0xfc,0x0a,0x7e]
345+
346+
v_sin_bf16 v5, exec_hi
347+
// GFX1250: v_sin_bf16_e32 v5, exec_hi ; encoding: [0x7f,0xfc,0x0a,0x7e]
348+
349+
v_sin_bf16 v5, null
350+
// GFX1250: v_sin_bf16_e32 v5, null ; encoding: [0x7c,0xfc,0x0a,0x7e]
351+
352+
v_sin_bf16 v5, -1
353+
// GFX1250: v_sin_bf16_e32 v5, -1 ; encoding: [0xc1,0xfc,0x0a,0x7e]
354+
355+
v_sin_bf16 v5, 0.5
356+
// GFX1250: v_sin_bf16_e32 v5, 0.5 ; encoding: [0xf0,0xfc,0x0a,0x7e]
357+
358+
v_sin_bf16 v5, src_scc
359+
// GFX1250: v_sin_bf16_e32 v5, src_scc ; encoding: [0xfd,0xfc,0x0a,0x7e]
360+
361+
v_sin_bf16 v127, 0x8000
362+
// GFX1250: v_sin_bf16_e32 v127, 0x8000 ; encoding: [0xff,0xfc,0xfe,0x7e,0x00,0x80,0x00,0x00]
363+
364+
v_sin_bf16 v5.h, v1.h
365+
// GFX1250: v_sin_bf16_e32 v5.h, v1.h ; encoding: [0x81,0xfd,0x0a,0x7f]
366+
319367
v_cvt_f32_bf16 v5, v1
320368
// GFX1250: v_cvt_f32_bf16_e32 v5, v1 ; encoding: [0x01,0xe5,0x0a,0x7e]
321369

‎llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -338,6 +338,62 @@ v_exp_bf16 v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 bound_ctrl:0 fi
338338
// GFX1250: v_exp_bf16_dpp v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 fi:1 ; encoding: [0xfa,0xfa,0xfe,0x7e,0x7f,0x6f,0x35,0x30]
339339
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
340340

341+
v_sin_bf16 v5, v1 quad_perm:[3,2,1,0]
342+
// GFX1250: v_sin_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x1b,0x00,0xff]
343+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
344+
345+
v_sin_bf16 v5, v1 quad_perm:[0,1,2,3]
346+
// GFX1250: v_sin_bf16_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0xe4,0x00,0xff]
347+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
348+
349+
v_sin_bf16 v5, v1 row_mirror
350+
// GFX1250: v_sin_bf16_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x40,0x01,0xff]
351+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
352+
353+
v_sin_bf16 v5, v1 row_half_mirror
354+
// GFX1250: v_sin_bf16_dpp v5, v1 row_half_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x41,0x01,0xff]
355+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
356+
357+
v_sin_bf16 v5, v1 row_shl:1
358+
// GFX1250: v_sin_bf16_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x01,0x01,0xff]
359+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
360+
361+
v_sin_bf16 v5, v1 row_shl:15
362+
// GFX1250: v_sin_bf16_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x0f,0x01,0xff]
363+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
364+
365+
v_sin_bf16 v5, v1 row_shr:1
366+
// GFX1250: v_sin_bf16_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x11,0x01,0xff]
367+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
368+
369+
v_sin_bf16 v5, v1 row_shr:15
370+
// GFX1250: v_sin_bf16_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x1f,0x01,0xff]
371+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
372+
373+
v_sin_bf16 v5, v1 row_ror:1
374+
// GFX1250: v_sin_bf16_dpp v5, v1 row_ror:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x21,0x01,0xff]
375+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
376+
377+
v_sin_bf16 v5, v1 row_ror:15
378+
// GFX1250: v_sin_bf16_dpp v5, v1 row_ror:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x2f,0x01,0xff]
379+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
380+
381+
v_sin_bf16 v5, v1 row_share:0 row_mask:0xf bank_mask:0xf
382+
// GFX1250: v_sin_bf16_dpp v5, v1 row_share:0 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x50,0x01,0xff]
383+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
384+
385+
v_sin_bf16 v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1
386+
// GFX1250: v_sin_bf16_dpp v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1 ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x5f,0x01,0x01]
387+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
388+
389+
v_sin_bf16 v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 fi:0
390+
// GFX1250: v_sin_bf16_dpp v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x60,0x09,0x13]
391+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
392+
393+
v_sin_bf16 v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 bound_ctrl:0 fi:1
394+
// GFX1250: v_sin_bf16_dpp v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 fi:1 ; encoding: [0xfa,0xfc,0xfe,0x7e,0x7f,0x6f,0x35,0x30]
395+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
396+
341397
v_cvt_f32_bf16 v5, v1 quad_perm:[3,2,1,0]
342398
// GFX1250: v_cvt_f32_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xe4,0x0a,0x7e,0x01,0x1b,0x00,0xff]
343399
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU

‎llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s

Lines changed: 60 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -362,6 +362,66 @@ v_exp_bf16 v5.h, v1.h quad_perm:[3,2,1,0]
362362
// GFX1250: v_exp_bf16_dpp v5.h, v1.h quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfa,0x0a,0x7f,0x81,0x1b,0x00,0xff]
363363
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
364364

365+
v_sin_bf16 v5, v1 quad_perm:[3,2,1,0]
366+
// GFX1250: v_sin_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x1b,0x00,0xff]
367+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
368+
369+
v_sin_bf16 v5, v1 quad_perm:[0,1,2,3]
370+
// GFX1250: v_sin_bf16_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0xe4,0x00,0xff]
371+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
372+
373+
v_sin_bf16 v5, v1 row_mirror
374+
// GFX1250: v_sin_bf16_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x40,0x01,0xff]
375+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
376+
377+
v_sin_bf16 v5, v1 row_half_mirror
378+
// GFX1250: v_sin_bf16_dpp v5, v1 row_half_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x41,0x01,0xff]
379+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
380+
381+
v_sin_bf16 v5, v1 row_shl:1
382+
// GFX1250: v_sin_bf16_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x01,0x01,0xff]
383+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
384+
385+
v_sin_bf16 v5, v1 row_shl:15
386+
// GFX1250: v_sin_bf16_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x0f,0x01,0xff]
387+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
388+
389+
v_sin_bf16 v5, v1 row_shr:1
390+
// GFX1250: v_sin_bf16_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x11,0x01,0xff]
391+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
392+
393+
v_sin_bf16 v5, v1 row_shr:15
394+
// GFX1250: v_sin_bf16_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x1f,0x01,0xff]
395+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
396+
397+
v_sin_bf16 v5, v1 row_ror:1
398+
// GFX1250: v_sin_bf16_dpp v5, v1 row_ror:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x21,0x01,0xff]
399+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
400+
401+
v_sin_bf16 v5, v1 row_ror:15
402+
// GFX1250: v_sin_bf16_dpp v5, v1 row_ror:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x2f,0x01,0xff]
403+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
404+
405+
v_sin_bf16 v5, v1 row_share:0 row_mask:0xf bank_mask:0xf
406+
// GFX1250: v_sin_bf16_dpp v5, v1 row_share:0 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x50,0x01,0xff]
407+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
408+
409+
v_sin_bf16 v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1
410+
// GFX1250: v_sin_bf16_dpp v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1 ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x5f,0x01,0x01]
411+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
412+
413+
v_sin_bf16 v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 fi:0
414+
// GFX1250: v_sin_bf16_dpp v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 ; encoding: [0xfa,0xfc,0x0a,0x7e,0x01,0x60,0x09,0x13]
415+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
416+
417+
v_sin_bf16 v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 bound_ctrl:0 fi:1
418+
// GFX1250: v_sin_bf16_dpp v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 fi:1 ; encoding: [0xfa,0xfc,0xfe,0x7e,0x7f,0x6f,0x35,0x30]
419+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
420+
421+
v_sin_bf16 v5.h, v1.h quad_perm:[3,2,1,0]
422+
// GFX1250: v_sin_bf16_dpp v5.h, v1.h quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xfc,0x0a,0x7f,0x81,0x1b,0x00,0xff]
423+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
424+
365425
v_cvt_f32_bf16 v5, v1 quad_perm:[3,2,1,0]
366426
// GFX1250: v_cvt_f32_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xe4,0x0a,0x7e,0x01,0x1b,0x00,0xff]
367427
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU

‎llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -74,6 +74,18 @@ v_exp_bf16 v127, v127 dpp8:[0,0,0,0,0,0,0,0] fi:0
7474
// GFX1250: v_exp_bf16_dpp v127, v127 dpp8:[0,0,0,0,0,0,0,0] ; encoding: [0xe9,0xfa,0xfe,0x7e,0x7f,0x00,0x00,0x00]
7575
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
7676

77+
v_sin_bf16 v5, v1 dpp8:[7,6,5,4,3,2,1,0]
78+
// GFX1250: v_sin_bf16_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding: [0xe9,0xfc,0x0a,0x7e,0x01,0x77,0x39,0x05]
79+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
80+
81+
v_sin_bf16 v5, v1 dpp8:[7,6,5,4,3,2,1,0] fi:1
82+
// GFX1250: v_sin_bf16_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] fi:1 ; encoding: [0xea,0xfc,0x0a,0x7e,0x01,0x77,0x39,0x05]
83+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
84+
85+
v_sin_bf16 v127, v127 dpp8:[0,0,0,0,0,0,0,0] fi:0
86+
// GFX1250: v_sin_bf16_dpp v127, v127 dpp8:[0,0,0,0,0,0,0,0] ; encoding: [0xe9,0xfc,0xfe,0x7e,0x7f,0x00,0x00,0x00]
87+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
88+
7789
v_cvt_f32_bf16 v5, v1 dpp8:[7,6,5,4,3,2,1,0]
7890
// GFX1250: v_cvt_f32_bf16_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding: [0xe9,0xe4,0x0a,0x7e,0x01,0x77,0x39,0x05]
7991
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU

0 commit comments

Comments
 (0)