Last active
December 14, 2015 06:38
-
-
Save allanmac/5043569 to your computer and use it in GitHub Desktop.
The `setp` and `selp` instructions are your friends.
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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; | |
} |
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
Compiled with
nvcc -m 32 -arch sm_30 -ptx natural.cu
: