Cp.async.cg.shared.global : L1 or L2 cache?

Does this instruction read from L1 or L2 cache?

What is the difference between .cg and .volatile? is this the correct description: .cg dictates on this read, if we have cache misses, where I would be caching the data (only l2) and .volatile dictates where would this read happen on (bypass l1)?

Paragraph 8 of the description for cp.async mentions:

" As described in Cache Operators, the .cg qualifier indicates caching of data only at global level cache L2 and not at L1 whereas .ca qualifier indicates caching of data at all levels including L1 cache. Cache operator are treated as performance hints only."

1 Like