From e1ebeed6dc27b67d279be87f99fb6829c10ac7b0 Mon Sep 17 00:00:00 2001 From: Kyle Wang Date: Tue, 26 Nov 2024 02:19:39 +0800 Subject: [PATCH] [AMD] NFC: Unified header guard in third_party/amd (#5244) This commit unified the names of header guards in third_party/amd. --- third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h | 6 +++--- third_party/amd/include/TritonAMDGPUToLLVM/GCNAsmFormat.h | 6 +++--- third_party/amd/include/TritonAMDGPUToLLVM/Passes.h | 6 +++--- .../include/TritonAMDGPUToLLVM/PatternTritonAMDGPUToLLVM.h | 6 +++--- third_party/amd/include/TritonAMDGPUToLLVM/TargetUtils.h | 6 +++--- third_party/amd/include/TritonAMDGPUTransforms/MfmaGroup.h | 6 +++--- third_party/amd/include/TritonAMDGPUTransforms/Passes.h | 6 +++--- .../include/TritonAMDGPUTransforms/TritonGPUConversion.h | 6 +++--- third_party/amd/lib/TritonAMDGPUToLLVM/BufferOpsEmitter.h | 6 +++--- .../ConvertLayoutOpToLLVM/SharedToDotOperandHelper.h | 6 +++--- .../ConvertLayoutOpToLLVM/SharedToDotOperandMFMA.cpp | 4 ++-- third_party/amd/lib/TritonAMDGPUToLLVM/DotOpToLLVM/MFMA.cpp | 6 ++++-- third_party/amd/lib/TritonAMDGPUToLLVM/OptimizeLDSUtility.h | 6 +++--- .../amd/lib/TritonAMDGPUToLLVM/PatternTritonGPUOpToLLVM.h | 6 +++--- third_party/amd/lib/TritonAMDGPUToLLVM/SchedInstructions.h | 6 +++--- third_party/amd/lib/TritonAMDGPUToLLVM/TargetInfo.h | 6 +++--- third_party/amd/lib/TritonAMDGPUToLLVM/Utility.h | 6 +++--- .../amd/lib/TritonAMDGPUTransforms/AccelerateAMDMatmul.cpp | 3 ++- .../amd/lib/TritonAMDGPUTransforms/CanonicalizePointers.cpp | 2 +- 19 files changed, 54 insertions(+), 51 deletions(-) diff --git a/third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h b/third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h index 6dbb0435e20c..59b11f0f7f96 100644 --- a/third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h +++ b/third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h @@ -21,8 +21,8 @@ * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef TRITON_DIALECT_AMDGPU_IR_DIALECT_H_ -#define TRITON_DIALECT_AMDGPU_IR_DIALECT_H_ +#ifndef TRITON_THIRD_PARTY_AMD_INCLUDE_DIALECT_TRITONAMDGPU_IR_DIALECT_H_ +#define TRITON_THIRD_PARTY_AMD_INCLUDE_DIALECT_TRITONAMDGPU_IR_DIALECT_H_ #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" @@ -47,4 +47,4 @@ namespace amdgpu {} // namespace amdgpu } // namespace triton } // namespace mlir -#endif // TRITON_DIALECT_TRITONGPU_IR_DIALECT_H_ +#endif // TRITON_THIRD_PARTY_AMD_INCLUDE_DIALECT_TRITONAMDGPU_IR_DIALECT_H_ diff --git a/third_party/amd/include/TritonAMDGPUToLLVM/GCNAsmFormat.h b/third_party/amd/include/TritonAMDGPUToLLVM/GCNAsmFormat.h index ac37aab817fa..ca356a31c0ee 100644 --- a/third_party/amd/include/TritonAMDGPUToLLVM/GCNAsmFormat.h +++ b/third_party/amd/include/TritonAMDGPUToLLVM/GCNAsmFormat.h @@ -20,8 +20,8 @@ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef TRITON_CONVERSION_TRITON_GPU_TO_LLVM_GCN_FORMAT_H_ -#define TRITON_CONVERSION_TRITON_GPU_TO_LLVM_GCN_FORMAT_H_ +#ifndef TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_GCNASMFORMAT_H_ +#define TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_GCNASMFORMAT_H_ #include "mlir/IR/Value.h" #include "triton/Dialect/Triton/IR/Dialect.h" @@ -400,4 +400,4 @@ struct GCNMemInstr : public GCNInstrBase { } // namespace triton } // namespace mlir -#endif // TRITON_CONVERSION_TRITON_GPU_TO_LLVM_ASM_FORMAT_H_ +#endif // TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_GCNASMFORMAT_H_ diff --git a/third_party/amd/include/TritonAMDGPUToLLVM/Passes.h b/third_party/amd/include/TritonAMDGPUToLLVM/Passes.h index f592694295a8..f60ba96161ef 100644 --- a/third_party/amd/include/TritonAMDGPUToLLVM/Passes.h +++ b/third_party/amd/include/TritonAMDGPUToLLVM/Passes.h @@ -1,5 +1,5 @@ -#ifndef TRITONAMDGPU_CONVERSION_PASSES_H -#define TRITONAMDGPU_CONVERSION_PASSES_H +#ifndef TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_PASSES_H_ +#define TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_PASSES_H_ #include "mlir/Conversion/LLVMCommon/TypeConverter.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" @@ -49,4 +49,4 @@ createTritonAMDGPULowerInstructionSchedHintsPass(StringRef arch, } // namespace mlir -#endif +#endif // TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_PASSES_H_ diff --git a/third_party/amd/include/TritonAMDGPUToLLVM/PatternTritonAMDGPUToLLVM.h b/third_party/amd/include/TritonAMDGPUToLLVM/PatternTritonAMDGPUToLLVM.h index 90922e802988..01975f680fd9 100644 --- a/third_party/amd/include/TritonAMDGPUToLLVM/PatternTritonAMDGPUToLLVM.h +++ b/third_party/amd/include/TritonAMDGPUToLLVM/PatternTritonAMDGPUToLLVM.h @@ -1,5 +1,5 @@ -#ifndef TRITONAMDGPU_TO_LLVM_PATTERNS_AMDGPU_OP_TO_LLVM_H -#define TRITONAMDGPU_TO_LLVM_PATTERNS_AMDGPU_OP_TO_LLVM_H +#ifndef TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_PATTERNTRITONAMDGPUTOLLVM_H_ +#define TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_PATTERNTRITONAMDGPUTOLLVM_H_ #include "mlir/Conversion/LLVMCommon/TypeConverter.h" @@ -11,4 +11,4 @@ void populateExtractSliceOpToLLVMPatterns( } -#endif +#endif // TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_PATTERNTRITONAMDGPUTOLLVM_H_ diff --git a/third_party/amd/include/TritonAMDGPUToLLVM/TargetUtils.h b/third_party/amd/include/TritonAMDGPUToLLVM/TargetUtils.h index 9e174d545dd9..6cf5548053b5 100644 --- a/third_party/amd/include/TritonAMDGPUToLLVM/TargetUtils.h +++ b/third_party/amd/include/TritonAMDGPUToLLVM/TargetUtils.h @@ -1,5 +1,5 @@ -#ifndef TRITON_CONVERSION_TRITONGPU_TO_LLVM_TARGETUTILS_H -#define TRITON_CONVERSION_TRITONGPU_TO_LLVM_TARGETUTILS_H +#ifndef TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_TARGETUTILS_H_ +#define TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_TARGETUTILS_H_ #include "llvm/ADT/StringRef.h" @@ -32,4 +32,4 @@ enum class DppCtrl : uint32_t { } // namespace mlir::triton::AMD -#endif // TRITON_CONVERSION_TRITONGPU_TO_LLVM_TARGETUTILS_H +#endif // TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTOLLVM_TARGETUTILS_H_ diff --git a/third_party/amd/include/TritonAMDGPUTransforms/MfmaGroup.h b/third_party/amd/include/TritonAMDGPUTransforms/MfmaGroup.h index 121bb617265f..57c8b6a58724 100644 --- a/third_party/amd/include/TritonAMDGPUTransforms/MfmaGroup.h +++ b/third_party/amd/include/TritonAMDGPUTransforms/MfmaGroup.h @@ -1,5 +1,5 @@ -#ifndef TRITON_DIALECT_TRITONAMDGPU_TRANSFORMS_MFMAGROUP_H_ -#define TRITON_DIALECT_TRITONAMDGPU_TRANSFORMS_MFMAGROUP_H_ +#ifndef TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_MFMAGROUP_H_ +#define TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_MFMAGROUP_H_ #include "mlir/Dialect/LLVMIR/ROCDLDialect.h" #include "llvm/ADT/DenseMap.h" @@ -91,4 +91,4 @@ class MfmaInsn { }; } // namespace mlir -#endif // TRITON_DIALECT_TRITONAMDGPU_TRANSFORMS_MFMAGROUP_H_ +#endif // TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_MFMAGROUP_H_ diff --git a/third_party/amd/include/TritonAMDGPUTransforms/Passes.h b/third_party/amd/include/TritonAMDGPUTransforms/Passes.h index 636743d305f9..0a8d51bc8f44 100644 --- a/third_party/amd/include/TritonAMDGPUTransforms/Passes.h +++ b/third_party/amd/include/TritonAMDGPUTransforms/Passes.h @@ -1,5 +1,5 @@ -#ifndef TRITON_DIALECT_TRITONAMDGPU_TRANSFORMS_PASSES_H_ -#define TRITON_DIALECT_TRITONAMDGPU_TRANSFORMS_PASSES_H_ +#ifndef TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_PASSES_H_ +#define TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_PASSES_H_ #include "mlir/Pass/Pass.h" #include "third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h" @@ -32,4 +32,4 @@ std::unique_ptr createTritonAMDGPUConvertToBufferOpsPass(); #include "TritonAMDGPUTransforms/Passes.h.inc" } // namespace mlir -#endif +#endif // TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_PASSES_H_ diff --git a/third_party/amd/include/TritonAMDGPUTransforms/TritonGPUConversion.h b/third_party/amd/include/TritonAMDGPUTransforms/TritonGPUConversion.h index fbfa235fc6bb..0e8b7a624010 100644 --- a/third_party/amd/include/TritonAMDGPUTransforms/TritonGPUConversion.h +++ b/third_party/amd/include/TritonAMDGPUTransforms/TritonGPUConversion.h @@ -4,8 +4,8 @@ // //===----------------------------------------------------------------------===// -#ifndef TRITON_DIALECT_TRITONGPU_TRANSFORMS_TRITONGPUCONVERSION_H_ -#define TRITON_DIALECT_TRITONGPU_TRANSFORMS_TRITONGPUCONVERSION_H_ +#ifndef TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_TRITONGPUCONVERSION_H_ +#define TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_TRITONGPUCONVERSION_H_ #include "mlir/Transforms/DialectConversion.h" @@ -35,4 +35,4 @@ class TritonGPUConversionTarget : public ConversionTarget { } // namespace mlir -#endif // TRITON_DIALECT_TRITONGPU_TRANSFORMS_TRITONGPUCONVERSION_H_ +#endif // TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_TRITONGPUCONVERSION_H_ diff --git a/third_party/amd/lib/TritonAMDGPUToLLVM/BufferOpsEmitter.h b/third_party/amd/lib/TritonAMDGPUToLLVM/BufferOpsEmitter.h index ad6d46ff78a0..0bef7a644729 100644 --- a/third_party/amd/lib/TritonAMDGPUToLLVM/BufferOpsEmitter.h +++ b/third_party/amd/lib/TritonAMDGPUToLLVM/BufferOpsEmitter.h @@ -1,5 +1,5 @@ -#ifndef TRITON_CONVERSION_TRITONAMDGPU_TO_LLVM_BUFFER_OPS_EMITTER_H -#define TRITON_CONVERSION_TRITONAMDGPU_TO_LLVM_BUFFER_OPS_EMITTER_H +#ifndef TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_BUFFEROPSEMITTER_H_ +#define TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_BUFFEROPSEMITTER_H_ #include "TargetInfo.h" #include "TritonAMDGPUToLLVM/GCNAsmFormat.h" @@ -90,4 +90,4 @@ struct BufferEmitter { } // namespace mlir::LLVM::AMD -#endif // TRITON_CONVERSION_TRITONAMDGPU_TO_LLVM_BUFFER_OPS_EMITTER_H +#endif // TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_BUFFEROPSEMITTER_H_ diff --git a/third_party/amd/lib/TritonAMDGPUToLLVM/ConvertLayoutOpToLLVM/SharedToDotOperandHelper.h b/third_party/amd/lib/TritonAMDGPUToLLVM/ConvertLayoutOpToLLVM/SharedToDotOperandHelper.h index 0db193e1c102..b2c6759fcb59 100644 --- a/third_party/amd/lib/TritonAMDGPUToLLVM/ConvertLayoutOpToLLVM/SharedToDotOperandHelper.h +++ b/third_party/amd/lib/TritonAMDGPUToLLVM/ConvertLayoutOpToLLVM/SharedToDotOperandHelper.h @@ -1,5 +1,5 @@ -#ifndef TRITON_CONVERSION_TRITONGPU_TO_LLVM_SHARED_TO_DOT_OPERAND_MATRIXCORE_H -#define TRITON_CONVERSION_TRITONGPU_TO_LLVM_SHARED_TO_DOT_OPERAND_MATRIXCORE_H +#ifndef TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_CONVERTLAYOUTOPTOLLVM_SHAREDTODOTOPERANDHELPER_H_ +#define TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_CONVERTLAYOUTOPTOLLVM_SHAREDTODOTOPERANDHELPER_H_ #include "Utility.h" @@ -61,4 +61,4 @@ llvm::SmallVector computeOffsetsBType( } // namespace mlir::triton::AMD -#endif +#endif // TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_CONVERTLAYOUTOPTOLLVM_SHAREDTODOTOPERANDHELPER_H_ diff --git a/third_party/amd/lib/TritonAMDGPUToLLVM/ConvertLayoutOpToLLVM/SharedToDotOperandMFMA.cpp b/third_party/amd/lib/TritonAMDGPUToLLVM/ConvertLayoutOpToLLVM/SharedToDotOperandMFMA.cpp index c79df66c482f..460bb37ab583 100644 --- a/third_party/amd/lib/TritonAMDGPUToLLVM/ConvertLayoutOpToLLVM/SharedToDotOperandMFMA.cpp +++ b/third_party/amd/lib/TritonAMDGPUToLLVM/ConvertLayoutOpToLLVM/SharedToDotOperandMFMA.cpp @@ -93,9 +93,9 @@ llvm::SmallVector> computeTensorElemMappingInBlock( Value laneVOffset = urem(laneId, nonKDim); Value laneHOffset; - if (iNonKDim == 32) + if (iNonKDim == 32) { laneHOffset = select(icmp_uge(laneId, _32), i32_val(numOfElems), _0); - else { + } else { // In this configuration warp contains 16 copies of same data if ((iKDim == 1 || iKDim == 4) && iNonKDim == 4) { laneHOffset = i32_val(0); diff --git a/third_party/amd/lib/TritonAMDGPUToLLVM/DotOpToLLVM/MFMA.cpp b/third_party/amd/lib/TritonAMDGPUToLLVM/DotOpToLLVM/MFMA.cpp index 1eed112c30c0..12e6100380c9 100644 --- a/third_party/amd/lib/TritonAMDGPUToLLVM/DotOpToLLVM/MFMA.cpp +++ b/third_party/amd/lib/TritonAMDGPUToLLVM/DotOpToLLVM/MFMA.cpp @@ -294,8 +294,9 @@ struct DotOpMFMAConversionHelper { // rocdl.mfma.f32.32x32x8bf16.1k calls for input of i16 type auto cast = bitcast(val, i16_ty); vec = insert_element(vecTy, vec, cast, i32_val(elemId)); - } else + } else { vec = insert_element(vecTy, vec, val, i32_val(elemId)); + } } if (type.getIntOrFloatBitWidth() == 8) { if (4 == kBase) @@ -303,8 +304,9 @@ struct DotOpMFMAConversionHelper { results.push_back(bitcast(vec, i32_ty)); if (8 == kBase) results.push_back(bitcast(vec, i64_ty)); - } else + } else { results.push_back(vec); + } } return results; } diff --git a/third_party/amd/lib/TritonAMDGPUToLLVM/OptimizeLDSUtility.h b/third_party/amd/lib/TritonAMDGPUToLLVM/OptimizeLDSUtility.h index 2bd2a977fa7d..701f03b129f3 100644 --- a/third_party/amd/lib/TritonAMDGPUToLLVM/OptimizeLDSUtility.h +++ b/third_party/amd/lib/TritonAMDGPUToLLVM/OptimizeLDSUtility.h @@ -1,5 +1,5 @@ -#ifndef TRITON_CONVERSION_TRITONAMDGPU_TO_LLVM_OPTIMIZE_LDS_UTILITY_H -#define TRITON_CONVERSION_TRITONAMDGPU_TO_LLVM_OPTIMIZE_LDS_UTILITY_H +#ifndef TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_OPTIMIZELDSUTILITY_H_ +#define TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_OPTIMIZELDSUTILITY_H_ #include "triton/Dialect/TritonGPU/IR/Dialect.h" @@ -47,4 +47,4 @@ estimateResourcesForReplacement(OpBuilder builder, } // namespace mlir::triton::AMD -#endif // TRITON_CONVERSION_TRITONAMDGPU_TO_LLVM_OPTIMIZE_LDS_UTILITY_H +#endif // TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_OPTIMIZELDSUTILITY_H_ diff --git a/third_party/amd/lib/TritonAMDGPUToLLVM/PatternTritonGPUOpToLLVM.h b/third_party/amd/lib/TritonAMDGPUToLLVM/PatternTritonGPUOpToLLVM.h index 1fdf3bdaa1cd..b217fc495643 100644 --- a/third_party/amd/lib/TritonAMDGPUToLLVM/PatternTritonGPUOpToLLVM.h +++ b/third_party/amd/lib/TritonAMDGPUToLLVM/PatternTritonGPUOpToLLVM.h @@ -1,5 +1,5 @@ -#ifndef TRITON_CONVERSION_TRITONAMDPU_TO_LLVM_PATTERNS_TRITON_GPU_OP_TO_LLVM_H -#define TRITON_CONVERSION_TRITONAMDPU_TO_LLVM_PATTERNS_TRITON_GPU_OP_TO_LLVM_H +#ifndef TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_PATTERNTRITONGPUOPTOLLVM_H_ +#define TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_PATTERNTRITONGPUOPTOLLVM_H_ #include "TargetInfo.h" #include "mlir/Conversion/LLVMCommon/TypeConverter.h" @@ -41,4 +41,4 @@ void populateUpcastMXFPToLLVMPatterns(LLVMTypeConverter &typeConverter, } // namespace mlir::triton::AMD -#endif +#endif // TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_PATTERNTRITONGPUOPTOLLVM_H_ diff --git a/third_party/amd/lib/TritonAMDGPUToLLVM/SchedInstructions.h b/third_party/amd/lib/TritonAMDGPUToLLVM/SchedInstructions.h index 45985fe808f2..b1836026032a 100644 --- a/third_party/amd/lib/TritonAMDGPUToLLVM/SchedInstructions.h +++ b/third_party/amd/lib/TritonAMDGPUToLLVM/SchedInstructions.h @@ -1,5 +1,5 @@ -#ifndef TRITON_CONVERSION_TRITONAMDGPU_TO_LLVM_SCHED_INSTRUCTIONS_H -#define TRITON_CONVERSION_TRITONAMDGPU_TO_LLVM_SCHED_INSTRUCTIONS_H +#ifndef TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_SCHEDINSTRUCTIONS_H_ +#define TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_SCHEDINSTRUCTIONS_H_ #include "mlir/IR/Types.h" #include "third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h" @@ -23,4 +23,4 @@ void storeOpConversionCallback(triton::gpu::LocalStoreOp op, size_t llvmOpCount, triton::DotOp getSingleDotOpIfExists(scf::ForOp forOp); } // namespace mlir::triton -#endif +#endif // TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_SCHEDINSTRUCTIONS_H_ diff --git a/third_party/amd/lib/TritonAMDGPUToLLVM/TargetInfo.h b/third_party/amd/lib/TritonAMDGPUToLLVM/TargetInfo.h index 0ce38d4d7660..195bef8a06e9 100644 --- a/third_party/amd/lib/TritonAMDGPUToLLVM/TargetInfo.h +++ b/third_party/amd/lib/TritonAMDGPUToLLVM/TargetInfo.h @@ -1,5 +1,5 @@ -#ifndef TRITON_CONVERSION_TRITONGPU_TO_LLVM_TARGETINFOAMD_H -#define TRITON_CONVERSION_TRITONGPU_TO_LLVM_TARGETINFOAMD_H +#ifndef TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_TARGETINFO_H_ +#define TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_TARGETINFO_H_ #include "TritonAMDGPUToLLVM/TargetUtils.h" #include "triton/Conversion/TritonGPUToLLVM/TargetInfoBase.h" @@ -68,4 +68,4 @@ class TargetInfo : public mlir::triton::TargetInfoBase { }; } // namespace mlir::triton::AMD -#endif // TRITON_CONVERSION_TRITONGPU_TO_LLVM_TARGETINFOAMD_H +#endif // TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_TARGETINFO_H_ diff --git a/third_party/amd/lib/TritonAMDGPUToLLVM/Utility.h b/third_party/amd/lib/TritonAMDGPUToLLVM/Utility.h index d150531848e3..cba2db5a896b 100644 --- a/third_party/amd/lib/TritonAMDGPUToLLVM/Utility.h +++ b/third_party/amd/lib/TritonAMDGPUToLLVM/Utility.h @@ -1,5 +1,5 @@ -#ifndef TRITON_CONVERSION_TRITONAMDGPU_TO_LLVM_UTILITY_H -#define TRITON_CONVERSION_TRITONAMDGPU_TO_LLVM_UTILITY_H +#ifndef TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_UTILITY_H_ +#define TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_UTILITY_H_ #include "TritonAMDGPUToLLVM/GCNAsmFormat.h" #include "TritonAMDGPUToLLVM/TargetUtils.h" @@ -49,4 +49,4 @@ void llStore(RewriterBase &rewriter, Location loc, Value ptr, Value val, triton::CacheModifier cm = triton::CacheModifier::NONE); } // namespace mlir::LLVM::AMD -#endif +#endif // TRITON_THIRD_PARTY_AMD_LIB_TRITONAMDGPUTOLLVM_UTILITY_H_ diff --git a/third_party/amd/lib/TritonAMDGPUTransforms/AccelerateAMDMatmul.cpp b/third_party/amd/lib/TritonAMDGPUTransforms/AccelerateAMDMatmul.cpp index b2602956a686..e8d3607c46ab 100644 --- a/third_party/amd/lib/TritonAMDGPUTransforms/AccelerateAMDMatmul.cpp +++ b/third_party/amd/lib/TritonAMDGPUTransforms/AccelerateAMDMatmul.cpp @@ -68,8 +68,9 @@ warpsPerTile(Operation *dotOp, ArrayRef shape, int numWarps, tensorShape[1] / shapePerWarp.second / ret[1]) { if (ret[0] < tensorShape[0] / shapePerWarp.first) { ret[0] *= 2; - } else + } else { ret[1] *= 2; + } } else { ret[1] *= 2; } diff --git a/third_party/amd/lib/TritonAMDGPUTransforms/CanonicalizePointers.cpp b/third_party/amd/lib/TritonAMDGPUTransforms/CanonicalizePointers.cpp index a5b32abfefd7..fa7d54e0fbee 100644 --- a/third_party/amd/lib/TritonAMDGPUTransforms/CanonicalizePointers.cpp +++ b/third_party/amd/lib/TritonAMDGPUTransforms/CanonicalizePointers.cpp @@ -107,7 +107,7 @@ class PointerCanonicalizer { // Utility copy functions FatPtr copy(Value newBasePtr, Value newOffset) { return FatPtr{newBasePtr, newOffset, canNarrow}; - }; + } FatPtr copyWithBase(Value newOffset) { return FatPtr{basePtr, newOffset, canNarrow}; }