The PTX docs give the following syntax:
prefetch{.tensormap_space}.tensormap [a]; // prefetch the tensormap
.tensormap_space = { .const, .param };
This indicates that only .const
and .param
spaces can be used. In the documentation text, it says
If the
.tensormap
qualifier is specified then theprefetch
instruction brings the cache line containing the specified address in the.const
or.param
memory state space for subsequent use by thecp.async.bulk.tensor
instruction.
Again, it sounds like only .const
or .param
can be used.
However, under the Examples section we have:
prefetch.global.tensormap [ptr];
I would appreciate clarification – is tensormap prefetch from global
space allowed or not?