CUDA 13.3 : the CUDA Tile C++ API has just been released.
But the documentation is terrible, it is like Cutlass, an unpalatable list of templates for concepts, but it really lacks pragmatic examples of how to do things.
For instance, let’s say I do not want to use ct::partition_view : there seems to be no example of how to manually create “loadable tiles”. Something like an API that explains what constructor/factory can be used, with what kind of paremeter will express an ROI…
I also need to understand masks for loading with masks and handled OOB gracefully, but here again, absolutely nothing will help having an overview of how such masks are created from realistic data.
Here is an example of what I would like to achieve :
- define a view of a strided 2D memory
- define a tile at an explicit location so that I can handle overlapping tiles
- define a mask so that I can align my tile to efficient left and right boundaries different from “interesting content”, and load it with a handy OOB behaviour (like 0 padding)
- define a tile subview to perform some operation, on the “interesting content”
After a few hours I am still not able to understand how to achieve those steps (except the first one : layouts are OK)
The API is very difficult to navigate. It does not seem to be made for humans. Please do not make it like Cutlass !
in addition to the API reference, some additional material is listed here.
This is not really a list of “correct syntax for basic tasks”.
- How do I instanciate a tile around a memory address without partition_view ?
- Even the implicit cast of a scalar into a tile is not clearly stated in the API doc (it works, but it is not documented for humans)
- How do I instanciate a tile from a tensor_span, with explicit left/right/top/right padding ?
- How do I extend a tensor_span (to add padding) or make a sub-tensor_span from an existing one (Region Of Interest) ?
- How do I build a mask to load from memory with explicit left/right/top/right OOB ?
- What is the type of a mask ? A Lambda, a bitmap, a tile with bool scalar type ?
Those are the kind of questions I cannot answer while navigating through the doc or the API headers.
[edit]
Currently, the most useful resource for those questions seems to be https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-tile-kernels.html#gather-and-scatter-boundary-handling
Update of new things I cannot find in the doc :
- When using ct::matmul, are there hardware “best shapes” to reach optimal performance ?
- How do I share data between tiles (for example something expensive to compute that could be computed once, or loaded once from global memory) ? With regular cuda, shared memory and __syncthreads() is the simplest way to do. I could not find a tile API for that.
(neither in https://docs.nvidia.com/cuda/cuda-tile-cpp-api-reference/general_operations.html#thread-functions nor in Memory Operations — CUDA Tile C++ API Reference 13.3 documentation)