| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
[mlir][nvgpu] Add tma.create.descriptor to create tensor map descriptor The Op creates a tensor map descriptor object representing tiled memory region. The descriptor is used by Tensor Memory Access (TMA). The tensor is the source tensor to be tiled. The boxDimensions is the size of the tiled memory region in each dimension. The pattern here lowers tma.create.descriptor to a runtime function call that eventually calls calls CUDA Driver's cuTensorMapEncodeTiled. For more information see below: https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TENSOR__MEMORY.html Depends on D155453 Reviewed By: nicolasvasilache Differential Revision: https://reviews.llvm.org/D155680 | 2 年前 | |
[mlir][nvgpu] Implement nvgpu.device_async_copy by NVVMToLLVM Pass nvgpu.device_async_copy is lowered into cp.async PTX instruction. However, NVPTX backend does not support its all mode especially when zero padding is needed. Therefore, current MLIR implementation genereates inline assembly for that. This work simplifies PTX generation for nvgpu.device_async_copy, and implements it by NVVMToLLVM Pass. Depends on D154060 Reviewed By: nicolasvasilache, manishucsd Differential Revision: https://reviews.llvm.org/D154345 | 2 年前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 2 年前 | ||
| 2 年前 |