projects / platform / upstream / llvm.git / commit
? search:
summary | shortlog | log | commit | commitdiff | tree
(parent: 5fdf4d5) | patch
[mlir][NVGPU] Support cache all (.ca) in nvgpu.device_async_copy
authorNicolas Vasilache <nicolas.vasilache@gmail.com>
Tue, 18 Apr 2023 11:28:45 +0000 (04:28 -0700)
committerNicolas Vasilache <nicolas.vasilache@gmail.com>
Tue, 18 Apr 2023 12:00:53 +0000 (05:00 -0700)
commit95cb9862a8dcd3b8e9cdf0a27b5eafb910c9e983
treed64d974a3cf60abeed14380ccb3aceb88429f187tree | snapshot
parent5fdf4d53819e613f5c5be0ca0ec12444c17812d7commit | diff
[mlir][NVGPU] Support cache all (.ca) in nvgpu.device_async_copy

This patch adds support for cache all (.ca) in conversion from nvgpu-to-nvvm for inline asm `cp.async`.

For sizes other than 16 bytes cp.async cache global is not allowed and cache all is required to generate a valid ptx.

Differential revision: https://reviews.llvm.org/D148604

Authored-by: Manish Gupta <manigupta@google.com>
mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp diff | blob | history
mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir diff | blob | history
Domain: System / Toolchain;