Bitwise matrix operation

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I try to understand more how the new matrix functions work. I found one very interesting bitwise operation that can execute bitwise xor/and operation.

Unfortunately it seems to execute a popc after the xor/and and returns the number of set bit. I would love to get the result of the xor/and +C.

So my first question is: Is there a way to get the result without the POPC instruction?

Second: What is the use case for this? Since it is implemented this way I assume that there is need in the market fore a function like this?

This is the PTX I’m looking at:
mma.sync.aligned.shape.row.col.s32.b1.b1.s32.bitOp.popc d, a, b, c;

Thanks

2

if bitOp.popc is the only available ptx instruction with bitOp, it’s not possible without popcount

xor + popc counts the number of mismatches between two bitmasks. and + popc would count the matches.

Note that according to the c++ programming guide CUDA C++ Programming Guide , some sub-byte instructions are deprecated and removed in Hopper, for example xor+popc. I don’t know how this affects the lower level ptx instructions.

3

Hi- thanks for your information.

4

You can get the full result of ‘and’ by only setting one bit of the same reference mask for each column of a matrix. But you would loose 32x of theoretical computation performance.