Skip to content

Commit 70261a5

Browse files
silee2silee2
committed
Update LLVM version.
1 parent 1918bc0 commit 70261a5

26 files changed

+77
-60
lines changed

‎benchmarks/pipelines/linalg-to-gpu.pp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -41,6 +41,6 @@
4141
convert-gpux-to-llvm
4242
expand-strided-metadata
4343
lower-affine
44-
convert-memref-to-llvm
44+
finalize-memref-to-llvm
4545
reconcile-unrealized-casts)
4646
// End

‎build_tools/llvm_version.txt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1 +1 @@
1-
922c8891d9916fa094147c5c793291f97a83c84c
1+
eecb8c5f06149baf970fa0943e9fb9a6afe00207

‎lib/Conversion/GPUXToLLVM/GPUXToLLVMPass.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -346,7 +346,7 @@ class ConvertLaunchFuncOpToGpuRuntimeCallPattern
346346
return mlir::LLVM::createGlobalString(
347347
loc, builder, globalName,
348348
mlir::StringRef(kernelName.data(), kernelName.size()),
349-
mlir::LLVM::Linkage::Internal);
349+
mlir::LLVM::Linkage::Internal, false);
350350
}
351351

352352
mlir::LogicalResult
@@ -375,7 +375,7 @@ class ConvertLaunchFuncOpToGpuRuntimeCallPattern
375375
nameBuffer.append(kGpuBinaryStorageSuffix);
376376
mlir::Value data = mlir::LLVM::createGlobalString(
377377
loc, rewriter, nameBuffer.str(), binaryAttr.getValue(),
378-
mlir::LLVM::Linkage::Internal);
378+
mlir::LLVM::Linkage::Internal, false);
379379

380380
auto size = rewriter.create<mlir::LLVM::ConstantOp>(
381381
loc, llvmIndexType,
@@ -593,7 +593,7 @@ void GPUXToLLVMPass::runOnOperation() {
593593
mlir::arith::populateArithToLLVMConversionPatterns(converter, patterns);
594594
mlir::cf::populateControlFlowToLLVMConversionPatterns(converter, patterns);
595595
mlir::populateVectorToLLVMConversionPatterns(converter, patterns);
596-
mlir::populateMemRefToLLVMConversionPatterns(converter, patterns);
596+
mlir::populateFinalizeMemRefToLLVMConversionPatterns(converter, patterns);
597597
mlir::populateFuncToLLVMConversionPatterns(converter, patterns);
598598
mlir::populateAsyncStructuralTypeConversionsAndLegality(converter, patterns,
599599
target);

‎lib/Transforms/InsertGPUAllocs.cpp

Lines changed: 32 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -305,16 +305,25 @@ class InsertGPUAllocsPass final
305305
filter.insert(copy);
306306
}
307307

308-
if (allocType != memrefType)
309-
allocResult = builder.create<mlir::memref::CastOp>(loc, memrefType,
310-
allocResult);
311-
312-
op.replaceAllUsesExcept(allocResult, filter);
313-
builder.setInsertionPoint(term);
314-
if (access.hostRead && access.deviceWrite) {
315-
builder.create<mlir::memref::CopyOp>(loc, allocResult, op);
308+
if (allocType != memrefType) {
309+
mlir::Value castedAllocResult = builder.create<mlir::memref::CastOp>(
310+
loc, memrefType, allocResult);
311+
312+
op.replaceAllUsesExcept(castedAllocResult, filter);
313+
builder.setInsertionPoint(term);
314+
if (access.hostRead && access.deviceWrite) {
315+
builder.create<mlir::memref::CopyOp>(loc, castedAllocResult, op);
316+
}
317+
builder.create<mlir::gpu::DeallocOp>(loc, std::nullopt,
318+
castedAllocResult);
319+
} else {
320+
op.replaceAllUsesExcept(allocResult, filter);
321+
builder.setInsertionPoint(term);
322+
if (access.hostRead && access.deviceWrite) {
323+
builder.create<mlir::memref::CopyOp>(loc, allocResult, op);
324+
}
325+
builder.create<mlir::gpu::DeallocOp>(loc, std::nullopt, allocResult);
316326
}
317-
builder.create<mlir::gpu::DeallocOp>(loc, std::nullopt, allocResult);
318327
} else if (m_clientAPI == "vulkan") {
319328
auto gpuAlloc =
320329
builder.create<mlir::memref::AllocOp>(loc, allocType, dims);
@@ -325,14 +334,21 @@ class InsertGPUAllocsPass final
325334
filter.insert(copy);
326335
}
327336

328-
if (allocType != memrefType)
329-
allocResult = builder.create<mlir::memref::CastOp>(loc, memrefType,
330-
allocResult);
337+
if (allocType != memrefType) {
338+
mlir::Value castedAllocResult = builder.create<mlir::memref::CastOp>(
339+
loc, memrefType, allocResult);
331340

332-
op.replaceAllUsesExcept(allocResult, filter);
333-
builder.setInsertionPoint(term);
334-
if (access.hostRead && access.deviceWrite) {
335-
builder.create<mlir::memref::CopyOp>(loc, allocResult, op);
341+
op.replaceAllUsesExcept(castedAllocResult, filter);
342+
builder.setInsertionPoint(term);
343+
if (access.hostRead && access.deviceWrite) {
344+
builder.create<mlir::memref::CopyOp>(loc, castedAllocResult, op);
345+
}
346+
} else {
347+
op.replaceAllUsesExcept(allocResult, filter);
348+
builder.setInsertionPoint(term);
349+
if (access.hostRead && access.deviceWrite) {
350+
builder.create<mlir::memref::CopyOp>(loc, allocResult, op);
351+
}
336352
}
337353
}
338354
};

‎test/Gen/PlaidML/linalg-to-cpu.pp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
builtin.module(convert-tensor-to-linalg
33
arith-bufferize
44
func.func(empty-tensor-to-alloc-tensor
5-
eliminate-empty-tensors
5+
//eliminate-empty-tensors
66
scf-bufferize
77
shape-bufferize
88
linalg-bufferize
@@ -22,7 +22,7 @@
2222
convert-index-to-llvm
2323
expand-strided-metadata
2424
lower-affine
25-
convert-memref-to-llvm
25+
finalize-memref-to-llvm
2626
lower-affine
2727
convert-func-to-llvm
2828
reconcile-unrealized-casts)

‎test/Gen/PlaidML/linalg-to-llvm.pp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
builtin.module(convert-tensor-to-linalg
44
arith-bufferize
55
func.func(empty-tensor-to-alloc-tensor
6-
eliminate-empty-tensors
6+
//eliminate-empty-tensors
77
scf-bufferize
88
shape-bufferize
99
linalg-bufferize
@@ -41,6 +41,6 @@
4141
convert-gpux-to-llvm
4242
expand-strided-metadata
4343
lower-affine
44-
convert-memref-to-llvm
44+
finalize-memref-to-llvm
4545
reconcile-unrealized-casts)
4646
// End

‎test/Integration/Dialect/Linalg/OpenCL/linalg-to-gpux-opencl.pp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
builtin.module(convert-tensor-to-linalg
44
arith-bufferize
55
func.func(empty-tensor-to-alloc-tensor
6-
eliminate-empty-tensors
6+
//eliminate-empty-tensors
77
scf-bufferize
88
shape-bufferize
99
linalg-bufferize
@@ -37,6 +37,6 @@
3737
convert-gpu-to-gpux
3838
convert-func-to-llvm
3939
convert-gpux-to-llvm
40-
convert-memref-to-llvm
40+
finalize-memref-to-llvm
4141
reconcile-unrealized-casts)
4242
// End

‎test/Integration/Dialect/Linalg/Vulkan/linalg-to-gpu-vulkan.pp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5,7 +5,7 @@
55
arith-bufferize
66
func.func(
77
empty-tensor-to-alloc-tensor
8-
eliminate-empty-tensors
8+
//eliminate-empty-tensors
99
scf-bufferize
1010
shape-bufferize
1111
linalg-bufferize

‎test/Jax/gordon/linalg-to-cpu.pp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,7 @@
44
convert-elementwise-to-linalg
55
arith-bufferize
66
func.func(empty-tensor-to-alloc-tensor
7-
eliminate-empty-tensors
7+
//eliminate-empty-tensors
88
scf-bufferize
99
shape-bufferize
1010
linalg-bufferize
@@ -22,7 +22,7 @@
2222
convert-index-to-llvm
2323
expand-strided-metadata
2424
lower-affine
25-
convert-memref-to-llvm
25+
finalize-memref-to-llvm
2626
convert-func-to-llvm
2727
reconcile-unrealized-casts)
2828
// End

‎test/Jax/gordon/linalg-to-llvm.pp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
builtin.module(convert-tensor-to-linalg
44
arith-bufferize
55
func.func(empty-tensor-to-alloc-tensor
6-
eliminate-empty-tensors
6+
//eliminate-empty-tensors
77
scf-bufferize
88
shape-bufferize
99
linalg-bufferize
@@ -41,6 +41,6 @@
4141
convert-gpux-to-llvm
4242
expand-strided-metadata
4343
lower-affine
44-
convert-memref-to-llvm
44+
finalize-memref-to-llvm
4545
reconcile-unrealized-casts)
4646
// End

0 commit comments

Comments
 (0)