diff --git a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp index bfe2159ab312..087f971290c9 100644 --- a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp +++ b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp @@ -196,7 +196,7 @@ struct PrintOpConversion static void llPrintf(StringRef msg, ValueRange args, ConversionPatternRewriter &rewriter) { - assert(!msg.empty() && "printf with empty string not support"); + assert(!msg.empty() && "printf with empty string not supported"); Type int8Ptr = ptr_ty(i8_ty); auto *ctx = rewriter.getContext(); @@ -685,29 +685,6 @@ struct AsyncBulkCommitGroupOpConversion } }; -namespace mlir { -namespace LLVM { - -void vprintf(StringRef msg, ValueRange args, - ConversionPatternRewriter &rewriter) { - PrintOpConversion::llPrintf(msg, args, rewriter); -} - -void vprintf_array(Value thread, ArrayRef arr, std::string info, - std::string elem_repr, ConversionPatternRewriter &builder) { - std::string fmt = info + " t-%d "; - std::vector new_arr({thread}); - for (int i = 0; i < arr.size(); ++i) { - fmt += elem_repr + ((i == arr.size() - 1) ? "" : ", "); - new_arr.push_back(arr[i]); - } - - vprintf(fmt, new_arr, builder); -} - -} // namespace LLVM -} // namespace mlir - void populateTritonGPUToLLVMPatterns( TritonGPUToLLVMTypeConverter &typeConverter, RewritePatternSet &patterns, int numWarps, ModuleAxisInfoAnalysis &axisInfoAnalysis, diff --git a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h index eddfb66131f2..8d7a772b1231 100644 --- a/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h +++ b/lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h @@ -37,19 +37,6 @@ namespace ttng = ::mlir::triton::nvidia_gpu; typedef DenseMap TensorPtrMapT; -namespace mlir { -namespace LLVM { - -// Helper function for using printf in LLVM conversion. -void vprintf(StringRef msg, ValueRange args, - ConversionPatternRewriter &rewriter); - -void vprintf_array(Value thread, ArrayRef arr, std::string info, - std::string elem_repr, ConversionPatternRewriter &builder); - -} // namespace LLVM -} // namespace mlir - // FuncOpConversion/FuncOpConversionBase is borrowed from // https://github.com/llvm/llvm-project/blob/fae656b2dd80246c3c6f01e9c77c49560368752c/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp#L276 // since it is not exposed on header files in mlir v14