How to retrieve ptx code from the Binary file without using cuobjdump ?

I know that PTX file is embedded in the compiled binaries and is encoded text. What is the best way to extract this text out of the cuda binary by decoding it, without using cuobjdump ?

Note that PTX is only embedded if the compilation switches call for it.
It’s encoded, perhaps, but its not really compressed, from what I can see.

Here’s a sample output from cuobjdump:

$ cuobjdump -ptx t1215

Fatbin elf code:
================
arch = sm_20
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit

Fatbin elf code:
================
arch = sm_20
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit

Fatbin ptx code:
================
arch = sm_20
code version = [5,0]
producer = cuda
host = linux
compile_size = 64bit
compressed








.version 5.0
.target sm_20
.address_size 64


.extern .func (.param .b32 func_retval0) vprintf
(
.param .b64 vprintf_param_0,
.param .b64 vprintf_param_1
)
;
.global .align 1 .b8 $str[7] = {104, 101, 108, 108, 111, 10, 0};

.visible .entry _Z8mykernelv(

)
{
.reg .b32 %r<2>;
.reg .b64 %rd<4>;


mov.u64 %rd1, $str;
cvta.global.u64 %rd2, %rd1;
mov.u64 %rd3, 0;

        {
.reg .b32 temp_param_reg;

        .param .b64 param0;
st.param.b64    [param0+0], %rd2;
.param .b64 param1;
st.param.b64    [param1+0], %rd3;
.param .b32 retval0;
call.uni (retval0),
vprintf,
(
param0,
param1
);
ld.param.b32    %r1, [retval0+0];


        }
        ret;
}

Here’s output just from opening the binary with a text editor:

\.version 5.0
.target sm_20
.address_size 64.^@ð^Kextern .func (.param .b32 ^R^@õ^E_retval0) vprintf
(
$^@$64^V^@^Q_^S^@?_0,^]^@^Hð^_1
)
;
.global .align 1 .b8 $str[7] = {104, 101^E^@^T8^E^@^Q1^O^@P, 0};à^@ò^Visible .entry _Z8mykernelv(

)
{
.regÄ^@d%r<2>;^Q^@ò^C64 %rd<4>;


mov.u^R^@!1,<83>^@c;
cvta¢^@^D^\^@!2,"^@^X;0^@<88>3, 0;

        j^@Ctempê^@0reg^^^@^H^K^A^A^K^@R0;
st^V^@^@^U^@"   [^V^@1+0]d^@=2;
3^@^_13^@^B^T13^@83;
²^A^C­^AÄ;
call.uni (À^A3,
<8f>^AR,
(
z^@",   ^@t1
);
ldh^@<93>32     %r1, [=^@0+0]1^AÐ       }
        ret;
}

Yes, You are right by compressed I meant encoded. I just used the work compressed because the cuobjdump calls it compressed.

So my question is basically how to decode this encoded ptx?

Because half of it is pretty much readable but rest is symbols. What encoding is that?

It would be interesting to know about the compression method. It’s pointless to keep it secret since you can read the PTX code with cuobjdump anyway.