Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Got compilation working on Windows. #4420

Closed
wants to merge 3 commits into from

Conversation

eaplatanios
Copy link
Contributor

This PR contains a couple more changes to get compilation working on Windows. I know that it's not part of the goals of the Triton team to support Windows yet as @Jokeren mentioned earlier today, but I'm hoping that the changes included here are non-invasive enough and safe to merge. The main issues that they address are:

  1. Namespace conflicts when using using namespace triton and using namespace triton::gpu that result in failures on Windows.
  2. A missing #include <string> that also resulted in a failure when trying to compile on Windows.
  3. Handling of the __builtin_ctz intrinsic for Windows as it's not supported by MSVC.

@ThomasRaoux
Copy link
Collaborator

Why can't you use clang to compile on Windows?

@eaplatanios
Copy link
Contributor Author

eaplatanios commented Jul 30, 2024

I'm using the Bazel build configuration that's propagated via XLA. I'm trying to get XLA to compile on Windows and it pulls in triton as part of its build. I tried to use clang directly but other things fail and it seemed harder to fix than getting the default XLA build on Windows working. With the changes in this PR I got everything to compile successfully on Windows using MSVC (XLA needed some more changes that I'll put up in a PR there after I can update the triton dependency).

Comment on lines 24 to 30
// We only "import" the symbols that we need to avoid name conflicts.
using triton::AxisInfo;
using triton::DialectInferLayoutInterface;
using triton::JoinOp;
using triton::ModuleAxisInfoAnalysis;
using triton::PointerType;
using triton::SplitOp;
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What do those names conflict with? I would assume using using namespace triton; should be enough?

Copy link
Contributor Author

@eaplatanios eaplatanios Jul 30, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

These names do not conflict. It's other stuff in the triton namespace that conflict. For example, there exist both mlir::detail and mlir::triton::detail iirc as one example and those conflicts were causing the build to fail.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What would it look like to make those explicit instead? (I don't see detail being used)

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Hmm I'm not sure I understand but I'm also not very familiar with C++ so maybe I'm missing something. My understanding of the issue I ran into is the following. The code previously looked something like this:

namespace mlir {
  // At this point the `mlir` already contains a symbol named `detail` coming from the MLIR codebase via the headers.

  // The following line brings all symbols from within `mlir::triton` in the current scope. That namespace also contains
  // a symbol named `detail`. That symbol now conflicts with `mlir::detail` that is already in scope. Further down in the
  // code when the compiler tries to prove a type expression in a SFINAE bound, it fails because of this ambiguity.
  using namespace triton;
}

In this PR, instead of importing all symbols from the mlir::triton namespace, I switched to only importing the ones that are used in this file, thus avoiding the name conflict/ambiguity. I'm actually not sure why this does not raise an error with Clang, but then again I've noticed how C++ compilers tend to not be super consistent with the standards sometimes and I'm not actually sure what the correct behavior would be here. This change makes it work for both compilers.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

what I'm asking is can we make the conflicting names explicit instead? There would be conflict only when trying to use a symbol that would exist in both namespace. I'm not sure I understand where that happens and why we cannot just make those names explicit.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

is there any way to repro without MSVC? I'm willing to take a look but don't have MSVC locally :')

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am tkaing a look and noticing we already use fully qualified names in some places despite the using namespace, and a lot of these symbols have a single use so using triton::symbol is definitely too bulky. Would you mind modifying this PR to remove using triton:: and use fully qualified names where appropriate?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Oh yeah definitely. I was thinking of doing that originally but wasn't sure what you'd prefer.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Regarding reproducing without MSVC, I looked a bit into whether there's any clang flag or something we could flip to reproduce the stricter behavior but I couldn't find anything. :/

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@ptillet I replaced all references with fully qualified names and removed the using statements.

@Jokeren
Copy link
Contributor

Jokeren commented Jul 30, 2024

OOO, why the other PR seems to be pretty huge? #4045 (comment)

@eaplatanios
Copy link
Contributor Author

eaplatanios commented Jul 30, 2024

OOO, why the other PR seems to be pretty huge? #4045 (comment)

This PR is only concerned with the core C++ library and not the Python bindings. Also, I'm not updating the CMake config because I'm building triton with Bazel as part of the XLA build. For context, I'm building a Rust library that uses this and that's why I don't need the Python bindings.

I have verified that with the changes in this PR, it compiles successfully on Windows. Note that it does not make any changes to CI pipelines etc., which was the main concern in that PR. I'm happy to just respond to changes in the future that break Windows support, as needed, rather than expect the Triton team to maintain Windows compatibility of course.

@ptillet
Copy link
Collaborator

ptillet commented Aug 2, 2024

I'm happy to just respond to changes in the future that break Windows support, as needed, rather than expect the Triton team to maintain Windows compatibility of cours

That is not the point. Even reviewing PR takes time (look at this one for example) and we would only accept PRs that are neutral-to-good independently of Windows support. If you want, I can setup triton-lang/triton-windows and make you an admin if you want to maintain windows support there. But you would be responsible for CI, releases and addressing issues.

@ptillet
Copy link
Collaborator

ptillet commented Aug 2, 2024

@eaplatanios although I am neutral wrt this PR (not a fan of blanket using namespace statements) and would be ok merging it in a vacuum, I think that it sets a bad precedent in this particular case. If you are OK with the above proposal of maintaining your own fork in the triton-lang org, please let me know and I can set this up. The advantage of the fork is that you can make sure that top-of-tree is always compatible with msvc, whereas in here it would constantly break

@ptillet ptillet closed this Aug 2, 2024
@eaplatanios
Copy link
Contributor Author

Unfortunately I have to be away from work for the next month for some personal reasons but I'll get back to you after that. XLA already maintains a separate fork of triton so I'll also ask those folks if adding the changes there would be ok instead of creating yet another fork.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants