Skip to content

[BACKEND] Delete dead vprintf and vprintf_array functions #2531

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
jlebar merged 2 commits into from
Oct 25, 2023
Merged

[BACKEND] Delete dead vprintf and vprintf_array functions #2531

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
da2abfa
Delete dead vprintf and vprintf_array functions.
jlebar Oct 19, 2023
2231c7d
Merge branch 'main' into jlebar/dead-vprintf
ptillet Oct 25, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
25 changes: 1 addition & 24 deletions lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVM.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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();
Expand Down Expand Up @@ -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<Value> arr, std::string info,
std::string elem_repr, ConversionPatternRewriter &builder) {
std::string fmt = info + " t-%d ";
std::vector<Value> 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,
Expand Down
13 changes: 0 additions & 13 deletions lib/Conversion/TritonGPUToLLVM/TritonGPUToLLVMBase.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,19 +37,6 @@ namespace ttng = ::mlir::triton::nvidia_gpu;

typedef DenseMap<Operation *, triton::MakeTensorPtrOp> 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<Value> 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
Expand Down