diff options
| author | Christopher Bate <cbate@nvidia.com> | 2022-06-17 13:02:47 -0600 |
|---|---|---|
| committer | Christopher Bate <cbate@nvidia.com> | 2022-06-17 16:26:40 -0600 |
| commit | 829c84ec5b8b08d97952b5f371291e95ec5fee43 (patch) | |
| tree | 70351cad49da041fae061f625bcf58787cddc45f /mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp | |
| parent | 91688716ba49942051dccdf7b9c4f81a7ec8feaf (diff) | |
| download | llvm-829c84ec5b8b08d97952b5f371291e95ec5fee43.tar.gz | |
[mlir][nvgpu] fix MSVC warning regarding left shift
Differential Revision: https://reviews.llvm.org/D128088
Diffstat (limited to 'mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp')
| -rw-r--r-- | mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp | 4 |
1 files changed, 2 insertions, 2 deletions
diff --git a/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp b/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp index 23d8c79a2518..3d01e2ee0998 100644 --- a/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp +++ b/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp @@ -64,7 +64,7 @@ static Value permuteVectorOffset(OpBuilder &b, Location loc, int64_t M = llvm::Log2_64(memrefTy.getDimSize(tgtDim)); // Capture bits[0:(M-N)] of src by first creating a (M-N) mask. - int64_t mask = (1 << (M - N)) - 1; + int64_t mask = (1LL << (M - N)) - 1; if (permuteEveryN > 1) mask = mask << llvm::Log2_64(permuteEveryN); Value srcBits = b.create<arith::ConstantIndexOp>(loc, mask); @@ -191,7 +191,7 @@ mlir::nvgpu::optimizeSharedMemoryReadsAndWrites(Operation *parentOp, (8 * kSharedMemoryLineSizeBytes / memRefType.getElementTypeBitWidth()) / rowSize; const int64_t threadGroupSize = - 1 << (7 - llvm::Log2_64(kDefaultVectorSizeBits / 8)); + 1LL << (7 - llvm::Log2_64(kDefaultVectorSizeBits / 8)); if (rowsPerLine >= threadGroupSize) return failure(); |