11#include " Dialect/TritonIntelGPU/IR/Dialect.h"
2- #include " mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
3- #include " mlir/IR/Matchers.h"
4- #include " mlir/IR/TypeUtilities.h"
5- #include " triton/Tools/LayoutUtils.h"
6- #include " llvm/ADT/SmallVector.h"
7- #include " llvm/ADT/TypeSwitch.h"
8-
92#include " PatternTritonGPUOpToLLVM.h"
103#include " TargetInfo.h"
114#include " Utility.h"
12- #include " triton/Conversion/TritonGPUToLLVM/Utility.h"
13-
145#include " intel/include/Dialect/TritonIntelGPU/IR/Attributes.h"
156#include " intel/include/Dialect/TritonIntelGPU/Transforms/Utility.h"
167#include " intel/include/Utils/Utility.h"
8+ #include " mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
9+ #include " mlir/IR/Matchers.h"
10+ #include " mlir/IR/TypeUtilities.h"
11+ #include " triton/Conversion/TritonGPUToLLVM/Utility.h"
12+ #include " triton/Tools/LayoutUtils.h"
1713#include " triton/Tools/LinearLayout.h"
14+ #include " triton/Tools/Sys/GetEnv.hpp"
15+ #include " llvm/ADT/SmallVector.h"
16+ #include " llvm/ADT/TypeSwitch.h"
1817#include < optional>
19- #include < triton/Tools/Sys/GetEnv.hpp>
2018
2119using namespace mlir ;
2220using namespace mlir ::triton;
@@ -26,11 +24,11 @@ using namespace mlir::triton::gpu::intel;
2624#define S (v ) StringAttr::get(ctx, (v))
2725
2826#if defined(_MSC_VER) && !defined(__clang__)
29- // MSVC does not provide the GCC/Clang built-ins __builtin_clz and __builtin_ctz.
30- // The following implementations use MSVC intrinsics to provide equivalent functionality.
31- // This is only needed when compiling with MSVC (and not Clang), to ensure cross-platform compatibility.
32- // See: https://gist.github.com/pps83/3210a2f980fd02bb2ba2e5a1fc4a2ef0
33- # if defined(_MSC_VER) && !defined(__clang__)
27+ // MSVC does not provide the GCC/Clang built-ins __builtin_clz and
28+ // __builtin_ctz. The following implementations use MSVC intrinsics to provide
29+ // equivalent functionality. This is only needed when compiling with MSVC (and
30+ // not Clang), to ensure cross-platform compatibility. See:
31+ // https://gist.github.com/pps83/3210a2f980fd02bb2ba2e5a1fc4a2ef0
3432#include < intrin.h>
3533
3634static int __builtin_clz (unsigned x) {
@@ -44,7 +42,6 @@ static int __builtin_ctz(unsigned x) {
4442 _BitScanForward (&r, x);
4543 return static_cast <int >(r);
4644}
47-
4845#endif
4946
5047namespace {
@@ -381,9 +378,10 @@ struct BlockIOConversionBase : public LoadStoreConversionBase {
381378 // \param ptr The pointer value whose pitch is to be computed.
382379 // \param elemSizeInBits The size of each element in bits.
383380 // \param dim The dimension along which to compute the pitch (stride).
384- // The default value is 0, which typically refers to the first (innermost) dimension.
385- // Use the default when you want the pitch for the first dimension; specify another
386- // value if you need the pitch for a different dimension.
381+ // The default value is 0, which typically refers to the first
382+ // (innermost) dimension. Use the default when you want the pitch
383+ // for the first dimension; specify another value if you need the
384+ // pitch for a different dimension.
387385 Value getPitch (ConversionPatternRewriter &rewriter, Value ptr,
388386 unsigned elemSizeInBits, unsigned dim = 0 ) const {
389387 Location loc = ptr.getLoc ();
0 commit comments