Skip to content

Commit 4a9ce65

Browse files
vlaindicvlaindic
authored
[AFAR VII] Using v_rcp_f32 instead of v_fmac_f32 in exec_mask_manipulation.cpp (#47)
use v_rcp_f32 instead of v_fmac_f32 [ROCm/rocprofiler-sdk commit: 61ce79c]
1 parent e0d71c3 commit 4a9ce65

File tree

2 files changed

+114
-114
lines changed

2 files changed

+114
-114
lines changed

‎tests/bin/pc-sampling/exec-mask-manipulation/exec_mask_manipulation.cpp‎

Lines changed: 101 additions & 101 deletions
Original file line numberDiff line numberDiff line change
@@ -278,7 +278,7 @@ kernel3(const float c)
278278
{
279279
double a = threadIdx.x;
280280
float i = 0;
281-
float d = 0;
281+
float d = threadIdx.x;
282282
float e = 0;
283283
int tid_even = threadIdx.x % 2;
284284
for(int j = 0; j < ITER_NUM; j++)
@@ -388,106 +388,106 @@ kernel3(const float c)
388388
}
389389
else
390390
{
391-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
392-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
393-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
394-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
395-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
396-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
397-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
398-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
399-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
400-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
401-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
402-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
403-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
404-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
405-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
406-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
407-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
408-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
409-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
410-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
411-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
412-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
413-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
414-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
415-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
416-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
417-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
418-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
419-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
420-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
421-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
422-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
423-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
424-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
425-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
426-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
427-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
428-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
429-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
430-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
431-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
432-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
433-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
434-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
435-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
436-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
437-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
438-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
439-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
440-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
441-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
442-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
443-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
444-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
445-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
446-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
447-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
448-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
449-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
450-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
451-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
452-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
453-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
454-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
455-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
456-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
457-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
458-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
459-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
460-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
461-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
462-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
463-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
464-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
465-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
466-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
467-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
468-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
469-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
470-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
471-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
472-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
473-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
474-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
475-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
476-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
477-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
478-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
479-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
480-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
481-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
482-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
483-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
484-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
485-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
486-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
487-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
488-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
489-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
490-
asm volatile("v_fmac_f32 %0, %0, %1\n" : "+v"(d) : "v"(e));
391+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
392+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
393+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
394+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
395+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
396+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
397+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
398+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
399+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
400+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
401+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
402+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
403+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
404+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
405+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
406+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
407+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
408+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
409+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
410+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
411+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
412+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
413+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
414+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
415+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
416+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
417+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
418+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
419+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
420+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
421+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
422+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
423+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
424+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
425+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
426+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
427+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
428+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
429+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
430+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
431+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
432+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
433+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
434+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
435+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
436+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
437+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
438+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
439+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
440+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
441+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
442+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
443+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
444+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
445+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
446+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
447+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
448+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
449+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
450+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
451+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
452+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
453+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
454+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
455+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
456+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
457+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
458+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
459+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
460+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
461+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
462+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
463+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
464+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
465+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
466+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
467+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
468+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
469+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
470+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
471+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
472+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
473+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
474+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
475+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
476+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
477+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
478+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
479+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
480+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
481+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
482+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
483+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
484+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
485+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
486+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
487+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
488+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
489+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
490+
asm volatile("v_rcp_f32 %0, %0\n" : "+v"(d), "=s"(e) : "s"(c));
491491
}
492492
}
493493
}

‎tests/rocprofv3/pc-sampling/host-trap/exec-mask-manipulation/validate.py‎

Lines changed: 13 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -162,11 +162,11 @@ def exec_mask_manipulation_validate_csv(df, all_sampled=False):
162162
all_source_lines_samples=all_sampled,
163163
)
164164

165-
# assert that v_fmac_f32 instructions are properly decoded
166-
# the v_fmac_f32 is executed by odd SIMD threads
165+
# assert that v_rcp_f32 instructions are properly decoded
166+
# the v_rcp_f32 is executed by odd SIMD threads
167167
validate_instruction_decoding(
168168
kernel_65_df,
169-
"v_fmac_f32",
169+
"v_rcp_f32",
170170
exec_mask_uint64=np.uint64(int("AAAAAAAAAAAAAAAA", 16)),
171171
source_code_lines_range=(391, 490),
172172
all_source_lines_samples=all_sampled,
@@ -210,15 +210,15 @@ def validate_json_exec_mask_manipulation(data_json, all_sampled=False):
210210
# execution mask where even SIMD lanes are active
211211
# correspond to the v_rcp_f64 instructions of the last kernel
212212
odd_simds_active_exec_mask = np.uint64(int("AAAAAAAAAAAAAAAA", 16))
213-
# start and end source code lines of the v_fmac_f32 0 instructions of the last kernel
214-
v_fmac_f32_start_line_num, v_fmac_f32_end_line_num = 391, 490
213+
# start and end source code lines of the v_rcp_f32 0 instructions of the last kernel
214+
v_rcp_f32_start_line_num, v_rcp_f32_end_line_num = 391, 490
215215

216216
# sampled wave_ids of the last kernel
217217
kernel65_sampled_wave_in_grp = set()
218218
# sampled source lines of the last kernel matching v_rcp_f64 instructions
219219
kernel65_v_rcp_64_sampled_source_line_set = set()
220220
# sampled source lines of the last kernel matching v_rcp_f64 instructions
221-
kernel65_v_fmac_f32_sampled_source_line_set = set()
221+
kernel65_v_rcp_f32_sampled_source_line_set = set()
222222
# sampled correlation IDs
223223
sampled_cids_set = set()
224224
# pairs of sampled SIMD ids and waveslot IDs
@@ -317,14 +317,14 @@ def validate_json_exec_mask_manipulation(data_json, all_sampled=False):
317317
and line_num <= v_rcp_f64_end_line_num
318318
)
319319
kernel65_v_rcp_64_sampled_source_line_set.add(line_num)
320-
elif inst.startswith("v_fmac_f32"):
320+
elif inst.startswith("v_rcp_f32"):
321321
# odd SIMD lanes active
322322
assert np.uint64(exec_mask) == odd_simds_active_exec_mask
323323
assert (
324-
line_num >= v_fmac_f32_start_line_num
325-
and line_num <= v_fmac_f32_end_line_num
324+
line_num >= v_rcp_f32_start_line_num
325+
and line_num <= v_rcp_f32_end_line_num
326326
)
327-
kernel65_v_fmac_f32_sampled_source_line_set.add(line_num)
327+
kernel65_v_rcp_f32_sampled_source_line_set.add(line_num)
328328

329329
if all_sampled:
330330
# All cids that belongs to the range [1, 65] should be samples
@@ -337,9 +337,9 @@ def validate_json_exec_mask_manipulation(data_json, all_sampled=False):
337337
assert len(kernel65_v_rcp_64_sampled_source_line_set) == (
338338
v_rcp_f64_end_line_num - v_rcp_f64_start_line_num + 1
339339
)
340-
# all source lines matches v_fmac_f32 instructions of the last kernel should be sampled
341-
assert len(kernel65_v_fmac_f32_sampled_source_line_set) == (
342-
v_fmac_f32_end_line_num - v_fmac_f32_start_line_num + 1
340+
# all source lines matches v_rcp_f32 instructions of the last kernel should be sampled
341+
assert len(kernel65_v_rcp_f32_sampled_source_line_set) == (
342+
v_rcp_f32_end_line_num - v_rcp_f32_start_line_num + 1
343343
)
344344

345345
# all chiplets must be sampled

0 commit comments

Comments
 (0)