| 1 | //===- NVGPU.cpp - C Interface for NVGPU dialect ------------------===// |
|---|
| 2 | // |
|---|
| 3 | // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. |
|---|
| 4 | // See https://llvm.org/LICENSE.txt for license information. |
|---|
| 5 | // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception |
|---|
| 6 | // |
|---|
| 7 | //===----------------------------------------------------------------------===// |
|---|
| 8 | |
|---|
| 9 | #include "mlir-c/Dialect/NVGPU.h" |
|---|
| 10 | #include "mlir/CAPI/Registration.h" |
|---|
| 11 | #include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h" |
|---|
| 12 | #include "mlir/IR/BuiltinTypes.h" |
|---|
| 13 | |
|---|
| 14 | using namespace mlir; |
|---|
| 15 | using namespace mlir::nvgpu; |
|---|
| 16 | |
|---|
| 17 | MLIR_DEFINE_CAPI_DIALECT_REGISTRATION(NVGPU, nvgpu, mlir::nvgpu::NVGPUDialect) |
|---|
| 18 | |
|---|
| 19 | bool mlirTypeIsANVGPUTensorMapDescriptorType(MlirType type) { |
|---|
| 20 | return isa<nvgpu::TensorMapDescriptorType>(Val: unwrap(c: type)); |
|---|
| 21 | } |
|---|
| 22 | |
|---|
| 23 | MlirType mlirNVGPUTensorMapDescriptorTypeGet(MlirContext ctx, |
|---|
| 24 | MlirType tensorMemrefType, |
|---|
| 25 | int swizzle, int l2promo, |
|---|
| 26 | int oobFill, int interleave) { |
|---|
| 27 | return wrap(cpp: nvgpu::TensorMapDescriptorType::get( |
|---|
| 28 | context: unwrap(c: ctx), tensor: cast<MemRefType>(Val: unwrap(c: tensorMemrefType)), |
|---|
| 29 | swizzle: TensorMapSwizzleKind(swizzle), l2promo: TensorMapL2PromoKind(l2promo), |
|---|
| 30 | oob: TensorMapOOBKind(oobFill), interleave: TensorMapInterleaveKind(interleave))); |
|---|
| 31 | } |
|---|
| 32 | |
|---|
