[NFC] Make NVGPU casing consistent (#91903)

Author: tyb0807

mlir/lib/Bindings/Python/DialectNVGPU.cpp

@@ -1,4 +1,4 @@
-//===--- DialectNvgpu.cpp - Pybind module for Nvgpu dialect API support ---===//
+//===--- DialectNVGPU.cpp - Pybind module for NVGPU dialect API support ---===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -17,7 +17,7 @@ using namespace mlir;
 using namespace mlir::python;
 using namespace mlir::python::adaptors;
 
-static void populateDialectNvgpuSubmodule(const pybind11::module &m) {
+static void populateDialectNVGPUSubmodule(const pybind11::module &m) {
   auto nvgpuTensorMapDescriptorType = mlir_type_subclass(
       m, "TensorMapDescriptorType", mlirTypeIsANVGPUTensorMapDescriptorType);
 
@@ -34,8 +34,8 @@ static void populateDialectNvgpuSubmodule(const pybind11::module &m) {
       py::arg("ctx") = py::none());
 }
 
-PYBIND11_MODULE(_mlirDialectsNvgpu, m) {
+PYBIND11_MODULE(_mlirDialectsNVGPU, m) {
   m.doc() = "MLIR NVGPU dialect.";
 
-  populateDialectNvgpuSubmodule(m);
+  populateDialectNVGPUSubmodule(m);
 }

mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp

@@ -441,7 +441,7 @@ struct PrepareContractToGPUMMA
   }
 };
 
-// Fold transpose op into the transfer read op. Nvgpu mma.sync op only supports
+// Fold transpose op into the transfer read op. NVGPU mma.sync op only supports
 // row-, column-, and row-major layout for matrixA, matrixB, and matrixC,
 // respectively. We can fold the transpose operation when loading the data from
 // Shared Memory to registers.

mlir/lib/Dialect/MemRef/Transforms/FoldMemRefAliasOps.cpp

@@ -362,7 +362,7 @@ class SubViewOfSubViewFolder : public OpRewritePattern<memref::SubViewOp> {
 
 /// Folds nvgpu.device_async_copy subviews into the copy itself. This pattern
 /// is folds subview on src and dst memref of the copy.
-class NvgpuAsyncCopyOpSubViewOpFolder final
+class NVGPUAsyncCopyOpSubViewOpFolder final
     : public OpRewritePattern<nvgpu::DeviceAsyncCopyOp> {
 public:
   using OpRewritePattern<nvgpu::DeviceAsyncCopyOp>::OpRewritePattern;
@@ -694,7 +694,7 @@ LogicalResult StoreOpOfCollapseShapeOpFolder<OpTy>::matchAndRewrite(
   return success();
 }
 
-LogicalResult NvgpuAsyncCopyOpSubViewOpFolder::matchAndRewrite(
+LogicalResult NVGPUAsyncCopyOpSubViewOpFolder::matchAndRewrite(
     nvgpu::DeviceAsyncCopyOp copyOp, PatternRewriter &rewriter) const {
 
   LLVM_DEBUG(DBGS() << "copyOp       : " << copyOp << "\n");
@@ -769,7 +769,7 @@ void memref::populateFoldMemRefAliasOpPatterns(RewritePatternSet &patterns) {
                LoadOpOfCollapseShapeOpFolder<memref::LoadOp>,
                StoreOpOfCollapseShapeOpFolder<affine::AffineStoreOp>,
                StoreOpOfCollapseShapeOpFolder<memref::StoreOp>,
-               SubViewOfSubViewFolder, NvgpuAsyncCopyOpSubViewOpFolder>(
+               SubViewOfSubViewFolder, NVGPUAsyncCopyOpSubViewOpFolder>(
       patterns.getContext());
 }
 

mlir/python/CMakeLists.txt

@@ -538,7 +538,7 @@ declare_mlir_python_extension(MLIRPythonExtension.Dialects.Quant.Pybind
 )
 
 declare_mlir_python_extension(MLIRPythonExtension.Dialects.NVGPU.Pybind
-  MODULE_NAME _mlirDialectsNvgpu
+  MODULE_NAME _mlirDialectsNVGPU
   ADD_TO_PARENT MLIRPythonSources.Dialects.nvgpu
   ROOT_DIR "${PYTHON_SOURCE_DIR}"
   SOURCES

mlir/python/mlir/dialects/nvgpu.py

@@ -4,4 +4,4 @@
 
 from ._nvgpu_ops_gen import *
 from ._nvgpu_enum_gen import *
-from .._mlir_libs._mlirDialectsNvgpu import *
+from .._mlir_libs._mlirDialectsNVGPU import *

mlir/test/lib/Dialect/NVGPU/TestNVGPUTransforms.cpp

@@ -68,9 +68,9 @@ struct TestMmaSyncF32ToTF32Patterns
 
 namespace mlir {
 namespace test {
-void registerTestNvgpuLowerings() {
+void registerTestNVGPULowerings() {
   PassRegistration<TestMmaSyncF32ToTF32Patterns>();
 }
 
 } // namespace test
-} // namespace mlir
+} // namespace mlir

mlir/tools/mlir-opt/mlir-opt.cpp

@@ -143,7 +143,7 @@ void registerTestTransformDialectEraseSchedulePass();
 void registerTestWrittenToPass();
 void registerTestVectorLowerings();
 void registerTestVectorReductionToSPIRVDotProd();
-void registerTestNvgpuLowerings();
+void registerTestNVGPULowerings();
 #if MLIR_ENABLE_PDL_IN_PATTERNMATCH
 void registerTestDialectConversionPasses();
 void registerTestPDLByteCodePass();
@@ -270,7 +270,7 @@ void registerTestPasses() {
   mlir::test::registerTestTransformDialectEraseSchedulePass();
   mlir::test::registerTestVectorLowerings();
   mlir::test::registerTestVectorReductionToSPIRVDotProd();
-  mlir::test::registerTestNvgpuLowerings();
+  mlir::test::registerTestNVGPULowerings();
   mlir::test::registerTestWrittenToPass();
 #if MLIR_ENABLE_PDL_IN_PATTERNMATCH
   mlir::test::registerTestDialectConversionPasses();