[mlir][nvvm] Add attribute to nvvm.cpAsyncOp to control l1 bypass
authorThomas Raoux <thomasraoux@google.com>
Mon, 9 May 2022 15:46:28 +0000 (15:46 +0000)
committerThomas Raoux <thomasraoux@google.com>
Mon, 9 May 2022 19:34:48 +0000 (19:34 +0000)
commit09fc685ce6808ae34de8e235bad686252eef3812
treeca6aa7e54de55ed36bb1353b85084b16ab309170
parent266ea446ab747671eb6c736569c3c9c5f3c53d11
[mlir][nvvm] Add attribute to nvvm.cpAsyncOp to control l1 bypass

Add attribute to be able to generate the intrinsic version of async copy
generating a copy with l1 bypass. This correspond to
cp.async.cg.shared.global in ptx.

Differential Revision: https://reviews.llvm.org/D125241
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
mlir/test/Dialect/LLVMIR/invalid.mlir
mlir/test/Dialect/LLVMIR/nvvm.mlir
mlir/test/Target/LLVMIR/nvvmir.mlir