@@ -1077,7 +1077,7 @@ static Value CreateSetGlobal(ConversionPatternRewriter &rewriter,
1077
1077
LogicalResult ConvertGpuModulePattern::matchAndRewrite (
1078
1078
mlir::gpu::GPUModuleOp module_op, OpAdaptor adaptor,
1079
1079
ConversionPatternRewriter &rewriter) const {
1080
- auto data = module_op->getAttrOfType <StringAttr>(getGpuBinaryAttrName ());
1080
+ auto data = module_op->getAttrOfType <StringAttr>(GetGpuBinaryAttrName ());
1081
1081
if (!data)
1082
1082
return rewriter.notifyMatchFailure (module_op, " no device code attribute" );
1083
1083
if (data.size () == 0 ) {
@@ -1087,7 +1087,7 @@ LogicalResult ConvertGpuModulePattern::matchAndRewrite(
1087
1087
}
1088
1088
Location loc = module_op->getLoc ();
1089
1089
auto constants =
1090
- module_op->getAttrOfType <ArrayAttr>(getGpuConstantsAttrName ());
1090
+ module_op->getAttrOfType <ArrayAttr>(GetGpuConstantsAttrName ());
1091
1091
mlir::FunctionType func_type = rewriter.getFunctionType (
1092
1092
rewriter.getType <ContextType>(), rewriter.getType <ModuleType>());
1093
1093
func::FuncOp func_op = rewriter.replaceOpWithNewOp <func::FuncOp>(
@@ -1129,7 +1129,7 @@ LogicalResult ConvertMemrefGlobalPattern::matchAndRewrite(
1129
1129
1130
1130
// If the global is a GPU module symbol, return that.
1131
1131
if (auto module_attr =
1132
- global_op->getAttrOfType <FlatSymbolRefAttr>(getGpuModuleAttrName ())) {
1132
+ global_op->getAttrOfType <FlatSymbolRefAttr>(GetGpuModuleAttrName ())) {
1133
1133
auto once_op = rewriter.create <compiler::OnceOp>(
1134
1134
loc, rewriter.getType <ModuleType>(), ValueRange (context), module_attr);
1135
1135
Value buffer = rewriter.create <ModuleGetGlobalOp>(
@@ -1588,7 +1588,7 @@ void ConvertGpuToTfrtGpuPass::runOnOperation() {
1588
1588
// Rewrite `gpu.module` before rewriting the referenced `memref.global` ops.
1589
1589
if (failed (ConvertGpuModuleOps (getOperation ()))) return signalPassFailure ();
1590
1590
1591
- auto converter = createMemrefToTfrtGpuConverter ();
1591
+ auto converter = CreateMemrefToTfrtGpuConverter ();
1592
1592
ConversionTarget target (getContext ());
1593
1593
1594
1594
// Rewrite `memref.load` before `tfrt_gpu.streamify` is inlined.
@@ -1710,11 +1710,11 @@ static Value MaterializeCast(OpBuilder &builder, Type type, ValueRange values,
1710
1710
return builder.create <CastOp>(loc, type, values).getResult (0 );
1711
1711
}
1712
1712
1713
- StringRef getGpuBinaryAttrName () { return " binary" ; }
1714
- StringRef getGpuConstantsAttrName () { return " constants" ; }
1715
- StringRef getGpuModuleAttrName () { return " gpu_module" ; }
1713
+ StringRef GetGpuBinaryAttrName () { return " binary" ; }
1714
+ StringRef GetGpuConstantsAttrName () { return " constants" ; }
1715
+ StringRef GetGpuModuleAttrName () { return " gpu_module" ; }
1716
1716
1717
- TypeConverter createMemrefToTfrtGpuConverter () {
1717
+ TypeConverter CreateMemrefToTfrtGpuConverter () {
1718
1718
TypeConverter converter;
1719
1719
converter.addConversion ([](Type type) { return type; });
1720
1720
converter.addConversion ([&](BaseMemRefType type) {
0 commit comments