Skip to content

Instantly share code, notes, and snippets.

@allanmac

allanmac/natural.cu

Last active December 14, 2015 06:38
Show Gist options
  • Star 0 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
  • Save allanmac/5043569 to your computer and use it in GitHub Desktop.
Save allanmac/5043569 to your computer and use it in GitHub Desktop.
Download ZIP
The `setp` and `selp` instructions are your friends.
Raw
natural.cu
extern "C"
__global__
void natural(const unsigned int b,
const unsigned int c,
const unsigned int y,
const unsigned int z,
const unsigned int id,
unsigned int* const out)
{
const bool flag = (id == 1);
const unsigned int a = flag ? b : c;
const unsigned int x = flag ? y : z;
out[0] = a;
out[1] = x;
}
@allanmac
Copy link
Author

cuobjdump.exe -sass natural.cubin

code for sm_30
    Function : natural
/*0008*/     /*0x10005de428004001*/     MOV R1, c [0x0] [0x44];
/*0010*/     /*0x40001de428004005*/     MOV R0, c [0x0] [0x150];
/*0018*/     /*0x00009de428004005*/     MOV R2, c [0x0] [0x140];
/*0020*/     /*0x2000dde428004005*/     MOV R3, c [0x0] [0x148];
/*0028*/     /*0x0401dc23190ec000*/     ISETP.EQ.AND P0, pt, R0, 0x1, pt;
/*0030*/     /*0x50001de428004005*/     MOV R0, c [0x0] [0x154];
/*0038*/     /*0x10209c0420004005*/     SEL R2, R2, c [0x0] [0x144], P0;
/*0048*/     /*0x3030dc0420004005*/     SEL R3, R3, c [0x0] [0x14c], P0;
/*0050*/     /*0x00009c8590000000*/     ST [R0], R2;
/*0058*/     /*0x1000dc8590000000*/     ST [R0+0x4], R3;
/*0060*/     /*0x00001de780000000*/     EXIT;
/*0068*/     /*0xe0001de74003ffff*/     BRA 0x68;
/*0070*/     /*0x00001de440000000*/     NOP CC.T;
/*0078*/     /*0x00001de440000000*/     NOP CC.T;
    ........................

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment