Jing's contribution: prototype of mixed precision gemm FP16/BF16xint4 GEMM

[aosewski]

Proposed changes

Prototype of Mixed Precision GEMM FP16/BF16xint4. Contribution from PR #1572

Checklist

Please put an x into the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask.

• I have added tests relevant to the introduced functionality, and the unit tests are passing locally
• I have added inline documentation which enables the maintainers with understanding the motivation
• I have removed the stale documentation which is no longer relevant after this pull request
• (If this change is user-facing) I have added release notes which provide the end users with a brief summary of the improvement from this pull request
• I have run clang-format on all changed files
• Any dependent changes have been merged

Discussion

This is just a prototype and as is it needs some polishing and further improvements. Especially there's a need to:

• Move PackSize into custom data type as a attribute
• Add is_pack_type in trait
• Refactor coordinate and copy function for sub-byte data type
• Refactor dynamic/static buffer/tensor for sub-byte data type

Those additional changes probably would be added in next PR.

[spolifroni-amd]

doesn't seem to be anything to do for docs.

[zjing14]

@aosewski Could post logs of CI failure?

[mtgu0705]

@zjing14, attached please find the CI failed log:

[2024-12-18T16:17:18.860Z] In file included from /var/jenkins/workspace/MLLIBS_composable_kernel_PR-1762/test/gemm/instance/gemm_f16_nt_instance.cpp:9:
[2024-12-18T16:17:18.861Z] In file included from /var/jenkins/workspace/MLLIBS_composable_kernel_PR-1762/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle.hpp:15:
[2024-12-18T16:17:18.861Z] In file included from /var/jenkins/workspace/MLLIBS_composable_kernel_PR-1762/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp:13:
[2024-12-18T16:17:18.861Z] In file included from /var/jenkins/workspace/MLLIBS_composable_kernel_PR-1762/include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp:10:
[2024-12-18T16:17:18.861Z] /var/jenkins/workspace/MLLIBS_composable_kernel_PR-1762/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp:378:27: error: static assertion failed due to requirement 'is_same_v<_Float16, ck::pk_i4_t>': transpose is not allowed for pk_i4_t
[2024-12-18T16:17:18.861Z] 378 | static_assert(is_same_v<remove_cvref_t, pk_i4_t>,
[2024-12-18T16:17:18.861Z] | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[2024-12-18T16:17:18.861Z] /var/jenkins/workspace/MLLIBS_composable_kernel_PR-1762/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp:460:9: note: in instantiation of function template specialization

[mtgu0705]

@zjing14 Also attached the fully log file:
failed.log.txt