[QST] Bitwise Operations with Cutlass datatypes

What is your question?
Hi, i want to perform bitwise operations (& or |) from inside a cutlass conv2d kernel (https://github.com/NVIDIA/cutlass/blob/24f991e87930e1159f1f5a47e329d43bcfbd76b9/include/cutlass/conv/kernel/implicit_gemm_convolution.h#L281), for example similar to this code:

uint32_t mask0 = 0xFFFFFFFF;  
uint32_t mask1 = 0x00000000;  
uint32_t val = reinterpret_cast<uint32_t&>(params.ptr_B[base_index]);
mask0 &= val;    
mask1 &= ~val;  
params.ptr_B[channel_idx] =  reinterpret_cast<ElementB&>(mask0);
params.ptr_B[channel_idx] =  reinterpret_cast<ElementB&>(mask1);

My doubt is which is the proper way to perform this operation knowing that the default datatype of cutlass (as is shown in example 16) is cutlass::half?

With cutlass::half, uint32_t is too big, so I tried with uint16_t, but when I print bit by bit the elements, it does not match the original value, so perhaps something is wrong. I also tried to change the ElementB datatype to float, but it seems is unsupported and similar to tfloat32_t, which gives me the following errors:

/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/conv/threadblock/conv2d_fprop_activation_tile_access_iterator_optimized.h(103): error: static assertion failed with "Require Iterations::kContiguous == 1"
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/conv/threadblock/conv2d_fprop_filter_tile_access_iterator_optimized.h(103): error: static assertion failed with "Require Iterations::kContiguous == 1"
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_policy.h(58): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(119): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(122): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(148): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(149): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(156): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(165): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(166): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(173): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(180): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(190): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(191): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(196): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(197): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(210): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/examples/16_ampere_tensorop_conv2dfprop/ampere_tensorop_conv2dfprop.cu(575): error: no operator "=" matches these operands
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/threadblock/mma_base.h(117): error: name followed by "::" must be a class or namespace name
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/threadblock/mma_base.h(205): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/threadblock/mma_base.h(208): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/conv/kernel/implicit_gemm_convolution.h(92): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/conv/kernel/implicit_gemm_convolution.h(99): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/threadblock/mma_base.h(224): error: too many initializer values
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/threadblock/mma_base.h(225): error: too many initializer values
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(230): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(231): error: incomplete type is not allowed
/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/include/cutlass/gemm/warp/mma_tensor_op_sm70.h(232): error: incomplete type is not allowed
28 errors detected in the compilation of "/mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/examples/16_ampere_tensorop_conv2dfprop/ampere_tensorop_conv2dfprop.cu".

So, how should I proceed ?

Thank you.

Start with examples/16_ampere_tensorop_conv2dfprop/ampere_tensorop_conv2dfprop.cu and the referenced conv/kernel/implicit_gemm_convolution.h entry point. Reproduce the reported errors while comparing the configured ElementB type with the iterator and MMA requirements. Done means establishing which datatypes and bitwise operations are supported in this convolution path, or documenting the required constraints.