Mojo module

tcgen05

This module includes utilities for working with the tensorcore 5th generation (tcgen05) instructions.

Aliases

check_blackwell_constraint

alias check_blackwell_constraint = constrained[_has_blackwell_tcgen05(), "The tcgen05 instructions are only applicable on nVidia Blackwell (sm_100a, sm_101a) hardware.", ?]

Structs

• ​TensorMemory: A wrapper around tensor memory allocated for tcgen05 instructions.

Functions

• ​tcgen05_alloc: Allocates tensor memory for use with tcgen05 instructions.
• ​tcgen05_cp: Copies data from shared memory described by the matrix descriptor s_desc to tensor memory tmem_addr.
• ​tcgen05_dealloc: Deallocates tensor memory allocated by tcgen05_alloc().
• ​tcgen05_fence_after: Orders all the subsequent asynchronous tcgen05 operations.
• ​tcgen05_fence_before: Orders all the prior asynchronous tcgen05 operations.
• ​tcgen05_ld: Loads data from tensor memory into registers.
• ​tcgen05_load_wait: Waits for tensor memory loads to complete.
• ​tcgen05_release_allocation_lock: Releases the allocation lock for the current CTA group.
• ​tcgen05_st: Stores data from registers into tensor memory.
• ​tcgen05_store_wait: Waits for tensor memory stores to complete.

Was this page helpful?

Thank you! We'll create more content like this.

Thank you for helping us improve!

😔 What went wrong?

Some code doesn’t work

It includes inaccurate information

It's missing information I need

It was difficult to understand

Other

Submit