better naming

Author: Hardcode84

mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h

@@ -29,10 +29,11 @@ void populateAMDGPUToROCDLConversionPatterns(LLVMTypeConverter &converter,
                                              RewritePatternSet &patterns,
                                              amdgpu::Chipset chipset);
 
+namespace amdgpu {
 /// Remap common GPU memory spaces (Workgroup, Private, etc) to LLVM address
 /// spaces.
-void populateCommonAMDGPUTypeAndAttributeConversions(
-    TypeConverter &typeConverter);
+void populateCommonGPUTypeAndAttributeConversions(TypeConverter &typeConverter);
+} // namespace amdgpu
 
 /// Remap AMDGPU memory spaces to LLVM address spaces
 /// by mapping amdgpu::AddressSpace::fat_raw_buffer to ptr addrspace(7),

mlir/include/mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h

@@ -35,10 +35,11 @@ MemRefType getMBarrierMemrefType(MLIRContext *context,
                                  MBarrierGroupType barrierType);
 } // namespace nvgpu
 
+namespace nvgpu {
 /// Remap common GPU memory spaces (Workgroup, Private, etc) to LLVM address
 /// spaces.
-void populateCommonNVGPUTypeAndAttributeConversions(
-    TypeConverter &typeConverter);
+void populateCommonGPUTypeAndAttributeConversions(TypeConverter &typeConverter);
+} // namespace nvgpu
 
 void populateNVGPUToNVVMConversionPatterns(const LLVMTypeConverter &converter,
                                            RewritePatternSet &patterns);

mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp

@@ -2998,7 +2998,7 @@ struct ConvertAMDGPUToROCDLPass
     LLVMTypeConverter converter(ctx);
 
     populateAMDGPUToROCDLConversionPatterns(converter, patterns, *maybeChipset);
-    populateCommonAMDGPUTypeAndAttributeConversions(converter);
+    amdgpu::populateCommonGPUTypeAndAttributeConversions(converter);
     LLVMConversionTarget target(getContext());
     target.addIllegalDialect<::mlir::amdgpu::AMDGPUDialect>();
     target.addLegalDialect<::mlir::LLVM::LLVMDialect>();
@@ -3010,7 +3010,7 @@ struct ConvertAMDGPUToROCDLPass
 };
 } // namespace
 
-void mlir::populateCommonAMDGPUTypeAndAttributeConversions(
+void mlir::amdgpu::populateCommonGPUTypeAndAttributeConversions(
     TypeConverter &typeConverter) {
   populateGpuMemorySpaceAttributeConversions(
       typeConverter, [](gpu::AddressSpace space) {

mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp

@@ -447,7 +447,7 @@ void mlir::configureGpuToNVVMConversionLegality(ConversionTarget &target) {
 }
 
 void mlir::configureGpuToNVVMTypeConverter(LLVMTypeConverter &converter) {
-  populateCommonNVGPUTypeAndAttributeConversions(converter);
+  nvgpu::populateCommonGPUTypeAndAttributeConversions(converter);
 
   // Lowering for MMAMatrixType.
   converter.addConversion([&](gpu::MMAMatrixType type) -> Type {

mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp

@@ -349,7 +349,7 @@ struct LowerGpuOpsToROCDLOpsPass final
     }
 
     LLVMTypeConverter converter(ctx, options);
-    populateCommonAMDGPUTypeAndAttributeConversions(converter);
+    amdgpu::populateCommonGPUTypeAndAttributeConversions(converter);
 
     RewritePatternSet llvmPatterns(ctx);
     LLVMConversionTarget target(getContext());

mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp

@@ -401,7 +401,7 @@ struct ConvertNVGPUToNVVMPass
     RewritePatternSet patterns(&getContext());
     LLVMTypeConverter converter(&getContext(), options);
     IRRewriter rewriter(&getContext());
-    populateCommonNVGPUTypeAndAttributeConversions(converter);
+    nvgpu::populateCommonGPUTypeAndAttributeConversions(converter);
 
     /// device-side async tokens cannot be materialized in nvvm. We just
     /// convert them to a dummy i32 type in order to easily drop them during
@@ -1708,7 +1708,7 @@ struct NVGPURcpOpLowering : public ConvertOpToLLVMPattern<nvgpu::RcpOp> {
 };
 } // namespace
 
-void mlir::populateCommonNVGPUTypeAndAttributeConversions(
+void mlir::nvgpu::populateCommonGPUTypeAndAttributeConversions(
     TypeConverter &typeConverter) {
   // NVVM uses alloca in the default address space to represent private
   // memory allocations, so drop private annotations. NVVM uses address

mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp

@@ -63,7 +63,7 @@ using namespace mlir::transform::gpu;
 void transform::ApplyGPUToNVVMConversionPatternsOp::populatePatterns(
     TypeConverter &typeConverter, RewritePatternSet &patterns) {
   auto &llvmTypeConverter = static_cast<LLVMTypeConverter &>(typeConverter);
-  populateCommonNVGPUTypeAndAttributeConversions(llvmTypeConverter);
+  nvgpu::populateCommonGPUTypeAndAttributeConversions(llvmTypeConverter);
   // Used in GPUToNVVM/WmmaOpsToNvvm.cpp so attaching here for now.
   // TODO: We should have a single to_nvvm_type_converter.
   llvmTypeConverter.addConversion(
@@ -112,7 +112,7 @@ LogicalResult transform::ApplyGPUSubgroupReduceToNVVMConversionPatternsOp::
 void transform::ApplyGPUToROCDLConversionPatternsOp::populatePatterns(
     TypeConverter &typeConverter, RewritePatternSet &patterns) {
   auto &llvmTypeConverter = static_cast<LLVMTypeConverter &>(typeConverter);
-  populateCommonAMDGPUTypeAndAttributeConversions(llvmTypeConverter);
+  amdgpu::populateCommonGPUTypeAndAttributeConversions(llvmTypeConverter);
   FailureOr<amdgpu::Chipset> maybeChipset =
       amdgpu::Chipset::parse(getChipset());
   assert(llvm::succeeded(maybeChipset) && "expected valid chipset");

mlir/lib/Dialect/NVGPU/TransformOps/NVGPUTransformOps.cpp

@@ -49,7 +49,7 @@ void ApplyNVGPUToNVVMConversionPatternsOp::populatePatterns(
   /// device-side async tokens cannot be materialized in nvvm. We just
   /// convert them to a dummy i32 type in order to easily drop them during
   /// conversion.
-  populateCommonNVGPUTypeAndAttributeConversions(llvmTypeConverter);
+  nvgpu::populateCommonGPUTypeAndAttributeConversions(llvmTypeConverter);
   llvmTypeConverter.addConversion([&](DeviceAsyncTokenType type) -> Type {
     return llvmTypeConverter.convertType(
         IntegerType::get(type.getContext(), 32));