Skip to content

Commit ec71843

Browse files
victor-edsvictor-eds
authored
[XPU][TritonIntelGPUToLLVM] Do not generate duplicated string constants (#3137)
Do not generate duplicated string constants as result of `tt.assert` or `tt.print` operations. Use a string constant cache to avoid duplication. Closes #3054. --------- Signed-off-by: victor-eds <[email protected]>
1 parent 4d8e9bb commit ec71843

File tree

7 files changed

+161
-54
lines changed

7 files changed

+161
-54
lines changed

‎test/Conversion/intel/tritongpu_to_gen.mlir‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1945,7 +1945,7 @@ module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32} {
19451945
#blocked0 = #ttg.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0], CTAsPerCGA = [1], CTASplitNum = [1], CTAOrder = [0]}>
19461946
module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32} {
19471947
// Test that %u format specifier is used if isSigned is false
1948-
// CHECK: llvm.mlir.global internal constant @printfFormat_0("pid (%u, %u, %u) idx ()int32 tensor: %u\0A\00") {addr_space = 2 : i32}
1948+
// CHECK: llvm.mlir.global internal constant @printfFormat_("pid (%u, %u, %u) idx ()int32 tensor: %u\0A\00") {addr_space = 2 : i32}
19491949
// CHECK-LABEL: print_int32_tensor_issigned_off
19501950
// CHECK: llvm.call @_Z18__spirv_ocl_printf(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) vararg(!llvm.func<i32 (ptr<2>, ...)>) : (!llvm.ptr<2>, i32, i32, i32, i32) -> i32
19511951
tt.func @print_int32_tensor_issigned_off(%arg0 : i32) {
@@ -1958,7 +1958,7 @@ module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32} {
19581958
#blocked0 = #ttg.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0], CTAsPerCGA = [1], CTASplitNum = [1], CTAOrder = [0]}>
19591959
module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32} {
19601960
// Test that %i format specifier is used if isSigned is true
1961-
// CHECK: llvm.mlir.global internal constant @printfFormat_0("pid (%u, %u, %u) idx ()int32 tensor: %i\0A\00") {addr_space = 2 : i32}
1961+
// CHECK: llvm.mlir.global internal constant @printfFormat_("pid (%u, %u, %u) idx ()int32 tensor: %i\0A\00") {addr_space = 2 : i32}
19621962
// CHECK-LABEL: print_int32_tensor_issigned_on
19631963
// CHECK: llvm.call @_Z18__spirv_ocl_printf(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) vararg(!llvm.func<i32 (ptr<2>, ...)>) : (!llvm.ptr<2>, i32, i32, i32, i32) -> i32
19641964
tt.func @print_int32_tensor_issigned_on(%arg0 : i32) {

‎test/Conversion/intel/tritonintelgpu_to_llvm.mlir‎

Lines changed: 85 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,89 @@
11
// RUN: triton-opt %s --convert-triton-intel-gpu-to-llvm | FileCheck %s
22

3+
#blocked = #ttg.blocked<{sizePerThread = [1], threadsPerWarp = [16], warpsPerCTA = [4], order = [0]}>
4+
35
// COM: check that the spirv target env is inserted
46
// CHECK: module attributes {{{.*}}spirv.target_env{{.*}}#spirv.resource_limits<subgroup_size = 16>
5-
module attributes { "ttg.threads-per-warp" = 16 : i32, "ttg.num-warps" = 4 : i32 } { }
7+
module attributes { "ttg.threads-per-warp" = 16 : i32, "ttg.num-warps" = 4 : i32 } {
8+
// As the assert message is shared, a single instance is emitted.
9+
10+
// CHECK-DAG: llvm.mlir.global internal constant @assertFunc_("unknown\00") {addr_space = 1 : i32}
11+
// CHECK-DAG: llvm.mlir.global internal constant @assertFile_("{{.*}}/test/Conversion/intel/tritonintelgpu_to_llvm.mlir\00") {addr_space = 1 : i32}
12+
// CHECK-DAG: llvm.mlir.global internal constant @assertMessage_("assert text\00") {addr_space = 1 : i32}
13+
// CHECK-DAG: llvm.mlir.global internal constant @assertMessage_3("different assert text\00") {addr_space = 1 : i32}
14+
// CHECK-DAG: llvm.func spir_funccc @__assert_fail(!llvm.ptr<4>, !llvm.ptr<4>, i32, !llvm.ptr<4>)
15+
16+
// CHECK: llvm.func spir_kernelcc @assert(%[[VAL_0:.*]]: !llvm.struct<(i1)>, %[[VAL_1:.*]]: !llvm.struct<(i1)>, %[[VAL_2:.*]]: !llvm.struct<(i1)>)
17+
tt.func public @assert(%arg0: tensor<1xi1, #blocked>, %arg1: tensor<1xi1, #blocked>, %arg2: tensor<1xi1, #blocked>) {
18+
// CHECK: %[[VAL_3:.*]] = llvm.extractvalue %[[VAL_0]][0] : !llvm.struct<(i1)>
19+
// CHECK: %[[VAL_4:.*]] = llvm.mlir.constant(false) : i1
20+
// CHECK: %[[VAL_5:.*]] = llvm.mlir.constant(false) : i1
21+
// CHECK: %[[VAL_6:.*]] = llvm.icmp "eq" %[[VAL_3]], %[[VAL_5]] : i1
22+
// CHECK: %[[VAL_7:.*]] = llvm.or %[[VAL_4]], %[[VAL_6]] : i1
23+
// CHECK: llvm.cond_br %[[VAL_7]], ^bb1, ^bb2
24+
// CHECK: ^bb1:
25+
// CHECK: %[[VAL_8:.*]] = llvm.mlir.addressof @assertMessage_ : !llvm.ptr<1>
26+
// CHECK: %[[VAL_9:.*]] = llvm.getelementptr %[[VAL_8]][0] : (!llvm.ptr<1>) -> !llvm.ptr<1>, i8
27+
// CHECK: %[[VAL_10:.*]] = llvm.mlir.addressof @assertFile_ : !llvm.ptr<1>
28+
// CHECK: %[[VAL_11:.*]] = llvm.getelementptr %[[VAL_10]][0] : (!llvm.ptr<1>) -> !llvm.ptr<1>, i8
29+
// CHECK: %[[VAL_12:.*]] = llvm.mlir.addressof @assertFunc_ : !llvm.ptr<1>
30+
// CHECK: %[[VAL_13:.*]] = llvm.getelementptr %[[VAL_12]][0] : (!llvm.ptr<1>) -> !llvm.ptr<1>, i8
31+
// CHECK: %[[VAL_14:.*]] = llvm.mlir.constant
32+
// CHECK: %[[VAL_15:.*]] = llvm.addrspacecast %[[VAL_9]] : !llvm.ptr<1> to !llvm.ptr<4>
33+
// CHECK: %[[VAL_16:.*]] = llvm.addrspacecast %[[VAL_11]] : !llvm.ptr<1> to !llvm.ptr<4>
34+
// CHECK: %[[VAL_17:.*]] = llvm.addrspacecast %[[VAL_13]] : !llvm.ptr<1> to !llvm.ptr<4>
35+
// CHECK: llvm.call spir_funccc @__assert_fail(%[[VAL_15]], %[[VAL_16]], %[[VAL_14]], %[[VAL_17]]) : (!llvm.ptr<4>, !llvm.ptr<4>, i32, !llvm.ptr<4>) -> ()
36+
// CHECK: llvm.br ^bb2
37+
// CHECK: ^bb2:
38+
// CHECK: %[[VAL_18:.*]] = llvm.mlir.constant(1 : i32) : i32
39+
// CHECK: llvm.call spir_funccc @_Z7barrierj(%[[VAL_18]]) {convergent, no_unwind, will_return} : (i32) -> ()
40+
tt.assert %arg0, "assert text" : tensor<1xi1, #blocked>
41+
// CHECK: %[[VAL_19:.*]] = llvm.extractvalue %[[VAL_1]][0] : !llvm.struct<(i1)>
42+
// CHECK: %[[VAL_20:.*]] = llvm.mlir.constant(false) : i1
43+
// CHECK: %[[VAL_21:.*]] = llvm.mlir.constant(false) : i1
44+
// CHECK: %[[VAL_22:.*]] = llvm.icmp "eq" %[[VAL_19]], %[[VAL_21]] : i1
45+
// CHECK: %[[VAL_23:.*]] = llvm.or %[[VAL_20]], %[[VAL_22]] : i1
46+
// CHECK: llvm.cond_br %[[VAL_23]], ^bb3, ^bb4
47+
// CHECK: ^bb3:
48+
// CHECK: %[[VAL_24:.*]] = llvm.mlir.addressof @assertMessage_ : !llvm.ptr<1>
49+
// CHECK: %[[VAL_25:.*]] = llvm.getelementptr %[[VAL_24]][0] : (!llvm.ptr<1>) -> !llvm.ptr<1>, i8
50+
// CHECK: %[[VAL_26:.*]] = llvm.mlir.addressof @assertFile_ : !llvm.ptr<1>
51+
// CHECK: %[[VAL_27:.*]] = llvm.getelementptr %[[VAL_26]][0] : (!llvm.ptr<1>) -> !llvm.ptr<1>, i8
52+
// CHECK: %[[VAL_28:.*]] = llvm.mlir.addressof @assertFunc_ : !llvm.ptr<1>
53+
// CHECK: %[[VAL_29:.*]] = llvm.getelementptr %[[VAL_28]][0] : (!llvm.ptr<1>) -> !llvm.ptr<1>, i8
54+
// CHECK: %[[VAL_30:.*]] = llvm.mlir.constant
55+
// CHECK: %[[VAL_31:.*]] = llvm.addrspacecast %[[VAL_25]] : !llvm.ptr<1> to !llvm.ptr<4>
56+
// CHECK: %[[VAL_32:.*]] = llvm.addrspacecast %[[VAL_27]] : !llvm.ptr<1> to !llvm.ptr<4>
57+
// CHECK: %[[VAL_33:.*]] = llvm.addrspacecast %[[VAL_29]] : !llvm.ptr<1> to !llvm.ptr<4>
58+
// CHECK: llvm.call spir_funccc @__assert_fail(%[[VAL_31]], %[[VAL_32]], %[[VAL_30]], %[[VAL_33]]) : (!llvm.ptr<4>, !llvm.ptr<4>, i32, !llvm.ptr<4>) -> ()
59+
// CHECK: llvm.br ^bb4
60+
// CHECK: ^bb4:
61+
// CHECK: %[[VAL_34:.*]] = llvm.mlir.constant(1 : i32) : i32
62+
// CHECK: llvm.call spir_funccc @_Z7barrierj(%[[VAL_34]]) {convergent, no_unwind, will_return} : (i32) -> ()
63+
tt.assert %arg1, "assert text" : tensor<1xi1, #blocked>
64+
// CHECK: %[[VAL_35:.*]] = llvm.extractvalue %[[VAL_2]][0] : !llvm.struct<(i1)>
65+
// CHECK: %[[VAL_36:.*]] = llvm.mlir.constant(false) : i1
66+
// CHECK: %[[VAL_37:.*]] = llvm.mlir.constant(false) : i1
67+
// CHECK: %[[VAL_38:.*]] = llvm.icmp "eq" %[[VAL_35]], %[[VAL_37]] : i1
68+
// CHECK: %[[VAL_39:.*]] = llvm.or %[[VAL_36]], %[[VAL_38]] : i1
69+
// CHECK: llvm.cond_br %[[VAL_39]], ^bb5, ^bb6
70+
// CHECK: ^bb5:
71+
// CHECK: %[[VAL_40:.*]] = llvm.mlir.addressof @assertMessage_3 : !llvm.ptr<1>
72+
// CHECK: %[[VAL_41:.*]] = llvm.getelementptr %[[VAL_40]][0] : (!llvm.ptr<1>) -> !llvm.ptr<1>, i8
73+
// CHECK: %[[VAL_42:.*]] = llvm.mlir.addressof @assertFile_ : !llvm.ptr<1>
74+
// CHECK: %[[VAL_43:.*]] = llvm.getelementptr %[[VAL_42]][0] : (!llvm.ptr<1>) -> !llvm.ptr<1>, i8
75+
// CHECK: %[[VAL_44:.*]] = llvm.mlir.addressof @assertFunc_ : !llvm.ptr<1>
76+
// CHECK: %[[VAL_45:.*]] = llvm.getelementptr %[[VAL_44]][0] : (!llvm.ptr<1>) -> !llvm.ptr<1>, i8
77+
// CHECK: %[[VAL_46:.*]] = llvm.mlir.constant
78+
// CHECK: %[[VAL_47:.*]] = llvm.addrspacecast %[[VAL_41]] : !llvm.ptr<1> to !llvm.ptr<4>
79+
// CHECK: %[[VAL_48:.*]] = llvm.addrspacecast %[[VAL_43]] : !llvm.ptr<1> to !llvm.ptr<4>
80+
// CHECK: %[[VAL_49:.*]] = llvm.addrspacecast %[[VAL_45]] : !llvm.ptr<1> to !llvm.ptr<4>
81+
// CHECK: llvm.call spir_funccc @__assert_fail(%[[VAL_47]], %[[VAL_48]], %[[VAL_46]], %[[VAL_49]]) : (!llvm.ptr<4>, !llvm.ptr<4>, i32, !llvm.ptr<4>) -> ()
82+
// CHECK: llvm.br ^bb6
83+
// CHECK: ^bb6:
84+
// CHECK: %[[VAL_50:.*]] = llvm.mlir.constant(1 : i32) : i32
85+
// CHECK: llvm.call spir_funccc @_Z7barrierj(%[[VAL_50]]) {convergent, no_unwind, will_return} : (i32) -> ()
86+
tt.assert %arg2, "different assert text" : tensor<1xi1, #blocked>
87+
tt.return
88+
}
89+
}

‎third_party/intel/lib/TritonIntelGPUToLLVM/PrintOpToLLVM.cpp‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,7 @@ struct PrintOpConversion
2121
PatternBenefit benefit)
2222
: ConvertTritonGPUOpToLLVMPattern<triton::PrintOp>(typeConverter,
2323
benefit),
24-
targetInfo(targetInfo) {}
24+
targetInfo(static_cast<const intel::TargetInfo &>(targetInfo)) {}
2525
using ConvertTritonGPUOpToLLVMPattern<
2626
triton::PrintOp>::ConvertTritonGPUOpToLLVMPattern;
2727

@@ -223,17 +223,17 @@ struct PrintOpConversion
223223
llvm::SmallString<64> msgNewline(msg);
224224
msgNewline.push_back('\n');
225225
msgNewline.push_back('\0');
226-
Value msgValue = LLVM::intel::addStringToModule(
227-
UnknownLoc::get(rewriter.getContext()), rewriter, "printfFormat_",
228-
msgNewline, TritonGEN::TritonGENMemorySpace::kUniformConstant);
226+
Value msgValue = targetInfo.getGlobalStringStart(
227+
rewriter.getUnknownLoc(), rewriter, "printfFormat_", msgNewline,
228+
/*addressSpace=*/TritonGEN::kUniformConstant);
229229
targetInfo.printf(rewriter, msgValue, msgNewline.size_in_bytes(), args);
230230
if (formatStrByteCount)
231231
*formatStrByteCount = msgNewline.size_in_bytes();
232232
return msgValue;
233233
}
234234

235235
protected:
236-
const TargetInfoBase &targetInfo;
236+
const intel::TargetInfo &targetInfo;
237237
};
238238

239239
} // namespace

‎third_party/intel/lib/TritonIntelGPUToLLVM/TargetInfo.cpp‎

Lines changed: 57 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -229,9 +229,9 @@ void TargetInfo::printf(RewriterBase &rewriter, StringRef msg,
229229
llvm::SmallString<64> msgNewline(msg);
230230
msgNewline.push_back('\n');
231231
msgNewline.push_back('\0');
232-
Value msgValue = LLVM::intel::addStringToModule(
233-
UnknownLoc::get(rewriter.getContext()), rewriter, "printfFormat_",
234-
msgNewline, /*AddressSpace*/ TritonGEN::kUniformConstant);
232+
Value msgValue = getGlobalStringStart(
233+
rewriter.getUnknownLoc(), rewriter, "printfFormat_", msgNewline,
234+
/*addressSpace=*/TritonGEN::kUniformConstant);
235235
printf(rewriter, msgValue, msgNewline.size_in_bytes(), args);
236236
}
237237

@@ -271,12 +271,15 @@ void TargetInfo::assertFail(RewriterBase &rewriter, Location loc,
271271
messageString.push_back('\0');
272272
fileString.push_back('\0');
273273
funcString.push_back('\0');
274-
Value messageStringVal = LLVM::intel::addStringToModule(
275-
loc, rewriter, "assertMessage_", messageString, addrSpace);
276-
Value fileStringVal = LLVM::intel::addStringToModule(
277-
loc, rewriter, "assertFile_", fileString, addrSpace);
278-
Value funcStringVal = LLVM::intel::addStringToModule(
279-
loc, rewriter, "assertFunc_", funcString, addrSpace);
274+
Value messageStringVal =
275+
getGlobalStringStart(loc, rewriter, "assertMessage_", messageString,
276+
/*addressSpace=*/TritonGEN::kCrossWorkgroup);
277+
Value fileStringVal =
278+
getGlobalStringStart(loc, rewriter, "assertFile_", fileString,
279+
/*addressSpace=*/TritonGEN::kCrossWorkgroup);
280+
Value funcStringVal =
281+
getGlobalStringStart(loc, rewriter, "assertFunc_", funcString,
282+
/*addressSpace=*/TritonGEN::kCrossWorkgroup);
280283
Value lineNumber = i32_val(line);
281284

282285
auto *ctx = rewriter.getContext();
@@ -312,4 +315,49 @@ Value TargetInfo::getStackPointer(RewriterBase &rewriter,
312315
return funcOp.getArgument(funcOp.getNumArguments() - 1);
313316
}
314317

318+
Value TargetInfo::getGlobalStringStart(Location loc, RewriterBase &rewriter,
319+
StringRef name, StringRef value,
320+
unsigned addressSpace) const {
321+
LLVM::GlobalOp global =
322+
getGlobalString(loc, rewriter, name, value, addressSpace);
323+
MLIRContext *ctx = rewriter.getContext();
324+
Type globalPtrType = ptr_ty(ctx, addressSpace);
325+
Value globalPtr = rewriter.create<LLVM::AddressOfOp>(loc, global);
326+
return gep(globalPtrType, i8_ty, globalPtr, LLVM::GEPArg{0});
327+
}
328+
329+
LLVM::GlobalOp TargetInfo::getGlobalString(Location loc, RewriterBase &rewriter,
330+
StringRef name, StringRef value,
331+
unsigned addressSpace) const {
332+
StringAttr valueAttr = rewriter.getStringAttr(value);
333+
std::pair<unsigned, StringAttr> cacheKey{addressSpace, valueAttr};
334+
auto pos = globals.find(cacheKey);
335+
if (pos != globals.end())
336+
return pos->second;
337+
338+
ModuleOp moduleOp = rewriter.getInsertionPoint()->getParentOfType<ModuleOp>();
339+
340+
llvm::SmallString<64> contentStr(value);
341+
size_t contentSize = contentStr.size_in_bytes();
342+
auto globalType = LLVM::LLVMArrayType::get(i8_ty, contentSize);
343+
344+
auto createGlobal = [&](StringRef name) {
345+
RewriterBase::InsertionGuard guard(rewriter);
346+
rewriter.setInsertionPointToStart(moduleOp.getBody());
347+
return rewriter.create<LLVM::GlobalOp>(
348+
rewriter.getUnknownLoc(), globalType,
349+
/*isConstant=*/true, LLVM::Linkage::Internal, name, valueAttr,
350+
/*alignment=*/0, addressSpace);
351+
};
352+
353+
LLVM::GlobalOp global =
354+
moduleOp.lookupSymbol(name)
355+
? createGlobal(Twine{name}.concat(Twine{globals.size()}).str())
356+
: createGlobal(name);
357+
358+
globals.try_emplace(cacheKey, global);
359+
360+
return global;
361+
}
362+
315363
} // namespace mlir::triton::intel

‎third_party/intel/lib/TritonIntelGPUToLLVM/TargetInfo.h‎

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,8 @@
1111

1212
#include "triton/Conversion/TritonGPUToLLVM/TargetInfoBase.h"
1313

14+
#include <mlir/Dialect/LLVMIR/LLVMDialect.h>
15+
1416
namespace mlir::triton::intel {
1517
class TargetInfo : public mlir::triton::TargetInfoBase {
1618
public:
@@ -69,7 +71,17 @@ class TargetInfo : public mlir::triton::TargetInfoBase {
6971
Value getStackPointer(RewriterBase &rewriter,
7072
FunctionOpInterface funcOp) const override;
7173

74+
Value getGlobalStringStart(Location loc, RewriterBase &rewriter,
75+
StringRef name, StringRef value,
76+
unsigned addressSpace) const;
77+
7278
private:
79+
LLVM::GlobalOp getGlobalString(Location loc, RewriterBase &rewriter,
80+
StringRef name, StringRef value,
81+
unsigned addressSpace) const;
82+
83+
mutable llvm::DenseMap<std::pair<unsigned, StringAttr>, LLVM::GlobalOp>
84+
globals;
7385
};
7486
} // namespace mlir::triton::intel
7587
#endif // TRITON_CONVERSION_TRITONGPU_TO_LLVM_TARGETINFOINTEL_H

‎third_party/intel/lib/TritonIntelGPUToLLVM/Utility.cpp‎

Lines changed: 0 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -83,40 +83,6 @@ Value shuffleIdx(Location loc, RewriterBase &rewriter, Value val, Value i) {
8383
return shuffleCommon(loc, rewriter, val, i, mlir::gpu::ShuffleMode::IDX);
8484
}
8585

86-
Value addStringToModule(Location loc, RewriterBase &rewriter, StringRef key,
87-
StringRef content, unsigned addressSpace) {
88-
auto moduleOp = rewriter.getBlock()->getParent()->getParentOfType<ModuleOp>();
89-
auto ctx = moduleOp.getContext();
90-
unsigned stringNumber = 0;
91-
SmallString<16> stringConstName;
92-
do {
93-
stringConstName.clear();
94-
(key + Twine(stringNumber++)).toStringRef(stringConstName);
95-
} while (moduleOp.lookupSymbol(stringConstName));
96-
97-
llvm::SmallString<64> contentStr(content);
98-
size_t contentSize = contentStr.size_in_bytes();
99-
auto globalType = LLVM::LLVMArrayType::get(i8_ty, contentSize);
100-
101-
LLVM::GlobalOp global;
102-
{
103-
RewriterBase::InsertionGuard guard(rewriter);
104-
rewriter.setInsertionPointToStart(moduleOp.getBody());
105-
global = rewriter.create<LLVM::GlobalOp>(
106-
UnknownLoc::get(ctx), globalType,
107-
/*isConstant=*/true, LLVM::Linkage::Internal, stringConstName,
108-
rewriter.getStringAttr(contentStr), /*alignment=*/0, addressSpace);
109-
}
110-
111-
Value zero = i32_val(0);
112-
Type globalPtrType = LLVM::LLVMPointerType::get(ctx, global.getAddrSpace());
113-
Value globalPtr = rewriter.create<LLVM::AddressOfOp>(
114-
UnknownLoc::get(ctx), globalPtrType, global.getSymName());
115-
Value stringStart = gep(ptr_ty(ctx, global.getAddrSpace()), i8_ty, globalPtr,
116-
SmallVector<Value>({zero}));
117-
return stringStart;
118-
}
119-
12086
// declare __spirv_ocl_printf(i8*, ...) as external function
12187
LLVM::LLVMFuncOp getSpirvPrintfDeclaration(RewriterBase &rewriter) {
12288
auto moduleOp = rewriter.getBlock()->getParent()->getParentOfType<ModuleOp>();

‎third_party/intel/lib/TritonIntelGPUToLLVM/Utility.h‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -78,9 +78,6 @@ Value shuffleUp(Location loc, RewriterBase &rewriter, Value val, int i);
7878
Value shuffleIdx(Location loc, RewriterBase &rewriter, Value val, int i);
7979
Value shuffleIdx(Location loc, RewriterBase &rewriter, Value val, Value i);
8080

81-
Value addStringToModule(Location loc, RewriterBase &rewriter, StringRef key,
82-
StringRef content, unsigned addressSpace);
83-
8481
LLVM::LLVMFuncOp getSpirvPrintfDeclaration(RewriterBase &rewriter);
8582

8683
static Value getModuleWarpSize(RewriterBase &rewriter, Location loc) {

0 commit comments

Comments
 (0)