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
MLIR_DEFINE_CAPI_DIALECT_REGISTRATION(NVGPU,nvgpu,mlir::nvgpu::NVGPUDialect)17 MLIR_DEFINE_CAPI_DIALECT_REGISTRATION(NVGPU, nvgpu, mlir::nvgpu::NVGPUDialect)
18
19 bool mlirTypeIsANVGPUTensorMapDescriptorType(MlirType type) {
20 return isa<nvgpu::TensorMapDescriptorType>(unwrap(type));
21 }
22
mlirNVGPUTensorMapDescriptorTypeGet(MlirContext ctx,MlirType tensorMemrefType,int swizzle,int l2promo,int oobFill,int interleave)23 MlirType mlirNVGPUTensorMapDescriptorTypeGet(MlirContext ctx,
24 MlirType tensorMemrefType,
25 int swizzle, int l2promo,
26 int oobFill, int interleave) {
27 return wrap(nvgpu::TensorMapDescriptorType::get(
28 unwrap(ctx), cast<MemRefType>(unwrap(tensorMemrefType)),
29 TensorMapSwizzleKind(swizzle), TensorMapL2PromoKind(l2promo),
30 TensorMapOOBKind(oobFill), TensorMapInterleaveKind(interleave)));
31 }
32