Skip to content

HTTPS clone URL

Subversion checkout URL

You can clone with
or
.
Download ZIP
Browse files

GK110: WIP atom-related decoding bug, bftool: allow mixed 32/64 bit i…

…nput
  • Loading branch information...
commit 99d6066236aada1b647fc413118ba15f4061758f 1 parent da2f4e7
@wwa wwa authored
Showing with 64 additions and 59 deletions.
  1. +43 −43 bftool
  2. +21 −16 envydis/gk110.c
View
86 bftool
@@ -1,65 +1,65 @@
#!/usr/bin/perl -sw
#RE tool, decodes hex into bitfield descriptions {pos,len}
# supports disjoint bitfields
-# supports 64 and 32 bit hex (in envydis's endianness)
+# supports mixed 64bit and 32bit (in envydis's endianness) hex words
# modes: and, or, xor (mask, diff)
#Examples:
-# find a diff between to OPs:
-# ./bftool -op=xor 8540dc00001c0002 8540dc00001c3c02
-# U8{0x0000000000003c00} U4{0x00003c00,0x00000000}
-# 10,4
-# the same, but with 32bit input:
-# ./bftool -op=xor 001c0002 8540dc00 001c3c02 8540dc00
-# U8{0x0000000000003c00} U4{0x00003c00,0x00000000}
-# 10,4
+# find a diff between two OPs:
+# ./bftool -op=xor 8540dc00001c0002 8540dc00001c3c02
+# or ./bftool -op=xor 8540dc00001c0002 001c3c02 8540dc00
+# or ./bftool -op=xor 001c0002 8540dc00 001c3c02 8540dc00
+# > x: U8{0x8540dc00001c0002} U4{0x001c0002,0x8540dc00}
+# y: U8{0x8540dc00001c3c02} U4{0x001c3c02,0x8540dc00}
+# z: U8{0x0000000000003c00} U4{0x00003c00,0x00000000}
+# {pos,len} = {10,4} = {0xa,4}
# find a match against a mask:
-# ./bftool -op=and 3fc0000000000000 8540dc00001c0002
+# ./bftool -op=and 3fc0000000000000 8540dc00001c0002
+# or ./bftool -op=and 3fc0000000000000 001c0002 8540dc00
+# ...
# U8{0x0540000000000000} U4{0x00000000,0x05400000}
# 54,1
# 56,1
# 58,1
+use Data::Dumper;
no warnings "portable";
my $op = $op || "nop";
-my $n = scalar(@ARGV);
-my $x = 0;
-my $y = 0;
-die "bad arg count" if not ($n==1 or $n==2 or $n==4);
-if($op eq "nop"){
- die "bad arg count" if $n>2;
- $x = hex($ARGV[0]);
- $x = $x | hex($ARGV[1]) << 32 if($n==2);
+sub readhex{
+ my $A = shift;
+ my $x = shift(@$A);
+ return undef if not defined($x);
+ return hex($x) if(length($x)>8 or scalar(@$A)==0 or length($A->[0])>8);
+ my $y = shift(@$A);
+ return hex($x)|hex($y) << 32;
}
-else{
- if($n==2){
- $x = hex($ARGV[0]);
- $y = hex($ARGV[1]);
- }
- elsif($n==4){
- $x = hex($ARGV[1]) << 32 | hex($ARGV[0]);
- $y = hex($ARGV[3]) << 32 | hex($ARGV[2]);
- }
+my $x = readhex(\@ARGV) || undef;
+my $y = readhex(\@ARGV) || undef;
+my $z = 0;
+if($op eq "nop"){
+ $z = $x;
}
-if($op eq "and"){
- $x = $x & $y;
+elsif($op eq "and"){
+ $z = $x & $y;
}
elsif($op eq "xor"){
- $x = $x ^ $y;
+ $z = $x ^ $y;
}
elsif($op eq "or"){
- $x = $x | $y;
+ $z = $x | $y;
}
-print sprintf("U8{0x%016x} U4{0x%08x,0x%08x}\n",$x,$x&0xFFFFFFFF,($x>>32)&0xFFFFFFFF);
+print sprintf("x: U8{0x%016x} U4{0x%08x,0x%08x}\n",$x,$x&0xFFFFFFFF,($x>>32)&0xFFFFFFFF);
+print sprintf("y: U8{0x%016x} U4{0x%08x,0x%08x}\n",$y,$y&0xFFFFFFFF,($y>>32)&0xFFFFFFFF);
+print sprintf("z: U8{0x%016x} U4{0x%08x,0x%08x}\n",$z,$z&0xFFFFFFFF,($z>>32)&0xFFFFFFFF);
my $pos=0;
-while($x!=0){
- while(($x&0x1) eq 0){
- $pos++;
- $x = $x >> 1;
- }
- my $len=0;
- while(($x&0x1) eq 1){
- $len++;
- $x = $x >> 1;
- }
- print $pos,",",$len,"\n";
+while($z!=0){
+ while(($z&0x1) eq 0){
+ $pos++;
+ $z = $z >> 1;
+ }
+ my $len=0;
+ while(($z&0x1) eq 1){
+ $len++;
+ $z = $z >> 1;
+ }
+ print sprintf("{pos,len} = {%d,%d} = {0x%x,%d}\n",$pos,$len,$pos,$len);
$pos+=$len;
}
View
37 envydis/gk110.c
@@ -902,21 +902,26 @@ static struct insn tablldstd[] = {
* c000000000000000 source type
*/
-//<WIP>
-// observed flag changes:
-// I[0]=I[1]<<29;
-// 00000048: 0e9c080d b7c00e00 lshf b32 $r3 (b64 $r2 $r3) 0x1d [unknown: 00000000 00000200]
-// 00000050: 0e9ffc09 b7c00800 lshf b32 $r2 (b64 0x0 $r2) 0x1d
-// I[0]=I[1]<<33;
-// 00000048: 109c080d b7c00e00 lshf b32 $r3 (b64 $r2 $r3) 0x21 [unknown: 00000000 00000200]
-// 00000050: 109ffc09 b7c00a00 lshf b32 $r2 (b64 0x0 $r2) 0x21 [unknown: 00000000 00000200]
-// ld.global.ca.nc.b32 %r4,[%rd1];
-// 00000020: 7f9c0801 60021084 tex lzero $r0:#:#:# t2d c[0x10] xy## $r2:$r3 0x0
-// ld.global.cs.nc.b32 %r4,[%rd1];
-// 00000020: 7f9c0802 70008484 $r0 $r2 0x0 $r33 ??? [unknown: 00000000 70000084] [unknown instruction]
-// ld.global.cg.nc.b32 %r4,[%rd1];
-// 00000020: 7f9c0802 70008084 $r0 $r2 0x0 $r32 ??? [unknown: 00000000 70000084] [unknown instruction]
-//</WIP>
+/* <WIP>
+observed flag changes:
+I[0]=I[1]<<29;
+ 00000048: 0e9c080d b7c00e00 lshf b32 $r3 (b64 $r2 $r3) 0x1d [unknown: 00000000 00000200]
+ 00000050: 0e9ffc09 b7c00800 lshf b32 $r2 (b64 0x0 $r2) 0x1d
+I[0]=I[1]<<33;
+ 00000048: 109c080d b7c00e00 lshf b32 $r3 (b64 $r2 $r3) 0x21 [unknown: 00000000 00000200]
+ 00000050: 109ffc09 b7c00a00 lshf b32 $r2 (b64 0x0 $r2) 0x21 [unknown: 00000000 00000200]
+
+ld.global.ca.nc.b32 %r4,[%rd1];
+ 00000020: 7f9c0801 60021084 tex lzero $r0:#:#:# t2d c[0x10] xy## $r2:$r3 0x0
+ld.global.cs.nc.b32 %r4,[%rd1];
+ 00000020: 7f9c0802 70008484 $r0 $r2 0x0 $r33 ??? [unknown: 00000000 70000084] [unknown instruction]
+ld.global.cg.nc.b32 %r4,[%rd1];
+ 00000020: 7f9c0802 70008084 $r0 $r2 0x0 $r32 ??? [unknown: 00000000 70000084] [unknown instruction]
+
+atom.global.exch.b32 r4, [ptr], r1;
+ 00000028: 001c0802 6c080000 mul $r0 u32 $r2 s32 0x10000000 [unknown: 00000000 40000000] //OOPS
+ 00000030: 001c0816 7b800000 $r5 $r2 $r0 $r0 ??? [unknown: 00000000 7b800000] [unknown instruction]
+</WIP> */
static struct insn tabm[] = {
{ 0x8400000000000002ull, 0xbfc0000000000003ull, T(sfuop), N("f32"), DST, T(neg33), T(abs31), SRC1 },
@@ -964,7 +969,7 @@ static struct insn tabm[] = {
{ 0x2600000000000002ull, 0x3fc0000000000003ull, N("cvt"), T(sat35), T(cvti2idst), T(neg30), T(abs34), T(cvti2isrc) },
{ 0x25c0000000000002ull, 0x3fc0000000000003ull, N("cvt"), T(frm2a), T(cvti2fdst), T(neg30), T(abs34), T(cvti2fsrc) },
{ 0x27c0000000000002ull, 0x3fc0000000000003ull, N("rshf"), N("b32"), DST, SESTART, T(us64_28), SRC1, SRC3, SEEND, T(shfclamp), T(is2) }, // XXX: check is2 and bits 0x29,0x33(swap srcs ?)
- { 0x2800000000000002ull, 0x3980000000000003ull, N("mul"), DST, T(us32_39), SRC1, T(us32_3a), LIMM },
+ { 0x2800000000000002ull, 0x3980000000000003ull, N("mul"), DST, T(us32_39), SRC1, T(us32_3a), LIMM }, //XXX: wrong! 6c280000011c100a is atom-related
{ 0x7400000000000002ull, 0x7f80000000000003ull, T(lane0e), N("mov"), N("b32"), DST, LIMM },
{ 0x7600000000000002ull, 0x7fc0000000000003ull, N("texgrad"), T(texm), TDST, T(text), TCONST, T(texgrsrc1), T(texgrsrc2) },
{ 0x7700000000000002ull, 0x7fc0000000000003ull, N("texbar"), TEXBARIMM },
Please sign in to comment.
Something went wrong with that request. Please try again.