From c4c768a34b10ad733526277e4c1ba8be39a9881c Mon Sep 17 00:00:00 2001 From: Aaron Siddhartha Mondal Date: Wed, 29 May 2024 16:10:56 +0200 Subject: [PATCH] =?UTF-8?q?=F0=9F=A5=AC=20Add=20`cuda=5Fnvptx=5Fnvcc`=20to?= =?UTF-8?q?olchain=20(#245)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit This is an initial take on a CUDA toolchain that uses `nvcc` as device compiler. Since `nvcc` doesn't properly work with `libcxx` this toolchain builds against `libstdc++` from nixpkgs. --- docs/reference/defs.md | 6 +-- docs/reference/ll.md | 6 +-- docs/reference/toolchain.md | 19 ++++---- docs/rules/defs.md | 6 +-- examples/BUILD.bazel | 12 +++++- examples/cuda_example/BUILD.bazel | 55 ++++++++++++++--------- examples/cuda_example/example.cpp | 16 +++---- ll/BUILD.bazel | 23 +++++++++- ll/args.bzl | 72 ++++++++++++++++++++----------- ll/attributes.bzl | 15 +++++++ ll/driver.bzl | 5 ++- ll/environment.bzl | 2 +- ll/toolchain.bzl | 3 ++ ll/tools.bzl | 2 + ll/transitions.bzl | 1 + modules/defaultLlEnv.nix | 18 +++++++- 16 files changed, 184 insertions(+), 77 deletions(-) diff --git a/docs/reference/defs.md b/docs/reference/defs.md index 6638795e..b26dfa52 100644 --- a/docs/reference/defs.md +++ b/docs/reference/defs.md @@ -36,7 +36,7 @@ Example: | `data` | List of labels, optional, defaults to [].

Extra files made available to compilation and linking steps.

Not appended to the default command line arguments, but available to actions. Reference these files manually for instance in the `includes`, and `compile_flags` attributes.

Use this attribute to make intermediary outputs from non-ll targets, for example from `rules_cc` or `filegroup`, available to the rule. | | `hdrs` | List of labels, optional, defaults to [].

Header files for this target.

When including a header file with a nested path, for instance `#include "some/path/myheader.h"`, add `"some/path"` to `includes` to make it visible to the rule.

Unavailable to downstream targets. | | `angled_includes` | List of strings, optional, defaults to [].

Angled include paths, relative to the target workspace.

Useful if you require include prefix stripping for dynamic paths, for instance the ones generated by `bzlmod`. Instead of `compile_flags = ["-Iexternal/mydep.someversion/include"]`, use `angled_includes = ["include"]` to add the path to the workspace automatically.

Expands paths starting with `$(GENERATED)` to the workspace location in the `GENDIR` path.

Unavailable to downstream targets. | -| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | +| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"cuda_nvptx_nvcc"` The CUDA toolchain with `nvcc` as device compiler.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | | `compile_flags` | List of strings, optional, defaults to [].

Flags for the compiler.

Pass a list of strings here. For instance `["-O3", "-std=c++20"]`.

Split flag pairs `-Xclang -somearg` into separate flags `["-Xclang", "-somearg"]`.

Unavailable to downstream targets. | | `compile_string_flags` | List of labels, optional, defaults to [].

Flags for the compiler in the form of `string_flag`s.

Splits the values of each `string_flag` along colons like so:

load("@bazel_skylib//rules:common_settings.bzl", "string_flag")

string_flag(
    name = "myflags",
    build_setting_default = "a:b:c",
)

ll_library(
    # ...
    # Equivalent to `compile_flags = ["a", "b", "c"]`
    compile_string_flags = [":myflags"],
)


Useful for externally configurable build attributes, such as generated flags from Nix environments. | | `defines` | List of strings, optional, defaults to [].

Defines for this target.

Pass a list of strings here. For instance `["MYDEFINE_1", "MYDEFINE_2"]`.

Unavailable to downstream targets. | @@ -120,7 +120,7 @@ Example: | `data` | List of labels, optional, defaults to [].

Extra files made available to compilation and linking steps.

Not appended to the default command line arguments, but available to actions. Reference these files manually for instance in the `includes`, and `compile_flags` attributes.

Use this attribute to make intermediary outputs from non-ll targets, for example from `rules_cc` or `filegroup`, available to the rule. | | `hdrs` | List of labels, optional, defaults to [].

Header files for this target.

When including a header file with a nested path, for instance `#include "some/path/myheader.h"`, add `"some/path"` to `includes` to make it visible to the rule.

Unavailable to downstream targets. | | `angled_includes` | List of strings, optional, defaults to [].

Angled include paths, relative to the target workspace.

Useful if you require include prefix stripping for dynamic paths, for instance the ones generated by `bzlmod`. Instead of `compile_flags = ["-Iexternal/mydep.someversion/include"]`, use `angled_includes = ["include"]` to add the path to the workspace automatically.

Expands paths starting with `$(GENERATED)` to the workspace location in the `GENDIR` path.

Unavailable to downstream targets. | -| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | +| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"cuda_nvptx_nvcc"` The CUDA toolchain with `nvcc` as device compiler.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | | `compile_flags` | List of strings, optional, defaults to [].

Flags for the compiler.

Pass a list of strings here. For instance `["-O3", "-std=c++20"]`.

Split flag pairs `-Xclang -somearg` into separate flags `["-Xclang", "-somearg"]`.

Unavailable to downstream targets. | | `compile_string_flags` | List of labels, optional, defaults to [].

Flags for the compiler in the form of `string_flag`s.

Splits the values of each `string_flag` along colons like so:

load("@bazel_skylib//rules:common_settings.bzl", "string_flag")

string_flag(
    name = "myflags",
    build_setting_default = "a:b:c",
)

ll_library(
    # ...
    # Equivalent to `compile_flags = ["a", "b", "c"]`
    compile_string_flags = [":myflags"],
)


Useful for externally configurable build attributes, such as generated flags from Nix environments. | | `defines` | List of strings, optional, defaults to [].

Defines for this target.

Pass a list of strings here. For instance `["MYDEFINE_1", "MYDEFINE_2"]`.

Unavailable to downstream targets. | @@ -177,7 +177,7 @@ Example: | `data` | List of labels, optional, defaults to [].

Extra files made available to compilation and linking steps.

Not appended to the default command line arguments, but available to actions. Reference these files manually for instance in the `includes`, and `compile_flags` attributes.

Use this attribute to make intermediary outputs from non-ll targets, for example from `rules_cc` or `filegroup`, available to the rule. | | `hdrs` | List of labels, optional, defaults to [].

Header files for this target.

When including a header file with a nested path, for instance `#include "some/path/myheader.h"`, add `"some/path"` to `includes` to make it visible to the rule.

Unavailable to downstream targets. | | `angled_includes` | List of strings, optional, defaults to [].

Angled include paths, relative to the target workspace.

Useful if you require include prefix stripping for dynamic paths, for instance the ones generated by `bzlmod`. Instead of `compile_flags = ["-Iexternal/mydep.someversion/include"]`, use `angled_includes = ["include"]` to add the path to the workspace automatically.

Expands paths starting with `$(GENERATED)` to the workspace location in the `GENDIR` path.

Unavailable to downstream targets. | -| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | +| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"cuda_nvptx_nvcc"` The CUDA toolchain with `nvcc` as device compiler.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | | `compile_flags` | List of strings, optional, defaults to [].

Flags for the compiler.

Pass a list of strings here. For instance `["-O3", "-std=c++20"]`.

Split flag pairs `-Xclang -somearg` into separate flags `["-Xclang", "-somearg"]`.

Unavailable to downstream targets. | | `compile_string_flags` | List of labels, optional, defaults to [].

Flags for the compiler in the form of `string_flag`s.

Splits the values of each `string_flag` along colons like so:

load("@bazel_skylib//rules:common_settings.bzl", "string_flag")

string_flag(
    name = "myflags",
    build_setting_default = "a:b:c",
)

ll_library(
    # ...
    # Equivalent to `compile_flags = ["a", "b", "c"]`
    compile_string_flags = [":myflags"],
)


Useful for externally configurable build attributes, such as generated flags from Nix environments. | | `defines` | List of strings, optional, defaults to [].

Defines for this target.

Pass a list of strings here. For instance `["MYDEFINE_1", "MYDEFINE_2"]`.

Unavailable to downstream targets. | diff --git a/docs/reference/ll.md b/docs/reference/ll.md index 9658858b..a1593b92 100644 --- a/docs/reference/ll.md +++ b/docs/reference/ll.md @@ -32,7 +32,7 @@ Example: | `data` | List of labels, optional, defaults to [].

Extra files made available to compilation and linking steps.

Not appended to the default command line arguments, but available to actions. Reference these files manually for instance in the `includes`, and `compile_flags` attributes.

Use this attribute to make intermediary outputs from non-ll targets, for example from `rules_cc` or `filegroup`, available to the rule. | | `hdrs` | List of labels, optional, defaults to [].

Header files for this target.

When including a header file with a nested path, for instance `#include "some/path/myheader.h"`, add `"some/path"` to `includes` to make it visible to the rule.

Unavailable to downstream targets. | | `angled_includes` | List of strings, optional, defaults to [].

Angled include paths, relative to the target workspace.

Useful if you require include prefix stripping for dynamic paths, for instance the ones generated by `bzlmod`. Instead of `compile_flags = ["-Iexternal/mydep.someversion/include"]`, use `angled_includes = ["include"]` to add the path to the workspace automatically.

Expands paths starting with `$(GENERATED)` to the workspace location in the `GENDIR` path.

Unavailable to downstream targets. | -| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | +| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"cuda_nvptx_nvcc"` The CUDA toolchain with `nvcc` as device compiler.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | | `compile_flags` | List of strings, optional, defaults to [].

Flags for the compiler.

Pass a list of strings here. For instance `["-O3", "-std=c++20"]`.

Split flag pairs `-Xclang -somearg` into separate flags `["-Xclang", "-somearg"]`.

Unavailable to downstream targets. | | `compile_string_flags` | List of labels, optional, defaults to [].

Flags for the compiler in the form of `string_flag`s.

Splits the values of each `string_flag` along colons like so:

load("@bazel_skylib//rules:common_settings.bzl", "string_flag")

string_flag(
    name = "myflags",
    build_setting_default = "a:b:c",
)

ll_library(
    # ...
    # Equivalent to `compile_flags = ["a", "b", "c"]`
    compile_string_flags = [":myflags"],
)


Useful for externally configurable build attributes, such as generated flags from Nix environments. | | `defines` | List of strings, optional, defaults to [].

Defines for this target.

Pass a list of strings here. For instance `["MYDEFINE_1", "MYDEFINE_2"]`.

Unavailable to downstream targets. | @@ -80,7 +80,7 @@ Example: | `data` | List of labels, optional, defaults to [].

Extra files made available to compilation and linking steps.

Not appended to the default command line arguments, but available to actions. Reference these files manually for instance in the `includes`, and `compile_flags` attributes.

Use this attribute to make intermediary outputs from non-ll targets, for example from `rules_cc` or `filegroup`, available to the rule. | | `hdrs` | List of labels, optional, defaults to [].

Header files for this target.

When including a header file with a nested path, for instance `#include "some/path/myheader.h"`, add `"some/path"` to `includes` to make it visible to the rule.

Unavailable to downstream targets. | | `angled_includes` | List of strings, optional, defaults to [].

Angled include paths, relative to the target workspace.

Useful if you require include prefix stripping for dynamic paths, for instance the ones generated by `bzlmod`. Instead of `compile_flags = ["-Iexternal/mydep.someversion/include"]`, use `angled_includes = ["include"]` to add the path to the workspace automatically.

Expands paths starting with `$(GENERATED)` to the workspace location in the `GENDIR` path.

Unavailable to downstream targets. | -| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | +| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"cuda_nvptx_nvcc"` The CUDA toolchain with `nvcc` as device compiler.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | | `compile_flags` | List of strings, optional, defaults to [].

Flags for the compiler.

Pass a list of strings here. For instance `["-O3", "-std=c++20"]`.

Split flag pairs `-Xclang -somearg` into separate flags `["-Xclang", "-somearg"]`.

Unavailable to downstream targets. | | `compile_string_flags` | List of labels, optional, defaults to [].

Flags for the compiler in the form of `string_flag`s.

Splits the values of each `string_flag` along colons like so:

load("@bazel_skylib//rules:common_settings.bzl", "string_flag")

string_flag(
    name = "myflags",
    build_setting_default = "a:b:c",
)

ll_library(
    # ...
    # Equivalent to `compile_flags = ["a", "b", "c"]`
    compile_string_flags = [":myflags"],
)


Useful for externally configurable build attributes, such as generated flags from Nix environments. | | `defines` | List of strings, optional, defaults to [].

Defines for this target.

Pass a list of strings here. For instance `["MYDEFINE_1", "MYDEFINE_2"]`.

Unavailable to downstream targets. | @@ -137,7 +137,7 @@ Example: | `data` | List of labels, optional, defaults to [].

Extra files made available to compilation and linking steps.

Not appended to the default command line arguments, but available to actions. Reference these files manually for instance in the `includes`, and `compile_flags` attributes.

Use this attribute to make intermediary outputs from non-ll targets, for example from `rules_cc` or `filegroup`, available to the rule. | | `hdrs` | List of labels, optional, defaults to [].

Header files for this target.

When including a header file with a nested path, for instance `#include "some/path/myheader.h"`, add `"some/path"` to `includes` to make it visible to the rule.

Unavailable to downstream targets. | | `angled_includes` | List of strings, optional, defaults to [].

Angled include paths, relative to the target workspace.

Useful if you require include prefix stripping for dynamic paths, for instance the ones generated by `bzlmod`. Instead of `compile_flags = ["-Iexternal/mydep.someversion/include"]`, use `angled_includes = ["include"]` to add the path to the workspace automatically.

Expands paths starting with `$(GENERATED)` to the workspace location in the `GENDIR` path.

Unavailable to downstream targets. | -| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | +| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"cuda_nvptx_nvcc"` The CUDA toolchain with `nvcc` as device compiler.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | | `compile_flags` | List of strings, optional, defaults to [].

Flags for the compiler.

Pass a list of strings here. For instance `["-O3", "-std=c++20"]`.

Split flag pairs `-Xclang -somearg` into separate flags `["-Xclang", "-somearg"]`.

Unavailable to downstream targets. | | `compile_string_flags` | List of labels, optional, defaults to [].

Flags for the compiler in the form of `string_flag`s.

Splits the values of each `string_flag` along colons like so:

load("@bazel_skylib//rules:common_settings.bzl", "string_flag")

string_flag(
    name = "myflags",
    build_setting_default = "a:b:c",
)

ll_library(
    # ...
    # Equivalent to `compile_flags = ["a", "b", "c"]`
    compile_string_flags = [":myflags"],
)


Useful for externally configurable build attributes, such as generated flags from Nix environments. | | `defines` | List of strings, optional, defaults to [].

Defines for this target.

Pass a list of strings here. For instance `["MYDEFINE_1", "MYDEFINE_2"]`.

Unavailable to downstream targets. | diff --git a/docs/reference/toolchain.md b/docs/reference/toolchain.md index 023b129c..4817a3a2 100644 --- a/docs/reference/toolchain.md +++ b/docs/reference/toolchain.md @@ -6,14 +6,14 @@ This file declares the `ll_toolchain` rule. ## `ll_toolchain` -
ll_toolchain(name, LL_AMD_INCLUDES, LL_AMD_LIBRARIES, LL_CFLAGS, LL_CUDA_DRIVER, LL_CUDA_TOOLKIT,
-             LL_DYNAMIC_LINKER, LL_LDFLAGS, address_sanitizer, archiver, bitcode_linker,
-             builtin_includes, c_driver, clang_tidy, clang_tidy_runner, compiler_runtime, cov,
-             cpp_abihdrs, cpp_abilib, cpp_driver, cpp_stdhdrs, cpp_stdlib, hip_libraries, hip_runtime,
-             leak_sanitizer, linker, linker_wrapper, llvm_project_deps, machine_code_tool,
-             memory_sanitizer, objcopy, offload_bundler, offload_packager, opt, profdata, profile,
-             rocm_device_libs, symbolizer, thread_sanitizer, undefined_behavior_sanitizer,
-             unwind_library)
+
ll_toolchain(name, LL_AMD_INCLUDES, LL_AMD_LIBRARIES, LL_CFLAGS, LL_CUDA_DRIVER, LL_CUDA_NVCC,
+             LL_CUDA_NVCC_CFLAGS, LL_CUDA_NVCC_LDFLAGS, LL_CUDA_TOOLKIT, LL_DYNAMIC_LINKER,
+             LL_LDFLAGS, address_sanitizer, archiver, bitcode_linker, builtin_includes, c_driver,
+             clang_tidy, clang_tidy_runner, compiler_runtime, cov, cpp_abihdrs, cpp_abilib,
+             cpp_driver, cpp_stdhdrs, cpp_stdlib, hip_libraries, hip_runtime, leak_sanitizer, linker,
+             linker_wrapper, llvm_project_deps, machine_code_tool, memory_sanitizer, objcopy,
+             offload_bundler, offload_packager, opt, profdata, profile, rocm_device_libs, symbolizer,
+             thread_sanitizer, undefined_behavior_sanitizer, unwind_library)
`attributes` @@ -24,6 +24,9 @@ This file declares the `ll_toolchain` rule. | `LL_AMD_LIBRARIES` | Label, optional, defaults to None.

Link search paths for dependencies making use of AMD toolchains.

Affects the `hip_amdgpu` toolchain. | | `LL_CFLAGS` | Label, optional, defaults to None.

Arbitrary flags added to all compile actions. | | `LL_CUDA_DRIVER` | Label, optional, defaults to None.

The path to the CUDA driver.

Affects the `cuda_nvptx` and `hip_nvptx` toolchains. | +| `LL_CUDA_NVCC` | Label, optional, defaults to None.

The path to the `nvcc` compiler.

Affects the `cuda_nvptx_nvcc` toolchain. Other `*_nvptx` toolchains use the `c_driver` or `cpp_driver` as device compiler. | +| `LL_CUDA_NVCC_CFLAGS` | Label, optional, defaults to None.

Compile flags used by the `cuda_nvptx_nvcc` toolchain. | +| `LL_CUDA_NVCC_LDFLAGS` | Label, optional, defaults to None.

Link flags used by the `cuda_nvptx_nvcc` toolchain. | | `LL_CUDA_TOOLKIT` | Label, optional, defaults to None.

The path to the CUDA toolkit.

Affects the `cuda_nvptx` and `hip_nvptx` toolchains. | | `LL_DYNAMIC_LINKER` | Label, optional, defaults to None.

The linker from the glibc we compile and link against. | | `LL_LDFLAGS` | Label, optional, defaults to None.

Arbitrary flags added to all link actions. | diff --git a/docs/rules/defs.md b/docs/rules/defs.md index 6638795e..b26dfa52 100644 --- a/docs/rules/defs.md +++ b/docs/rules/defs.md @@ -36,7 +36,7 @@ Example: | `data` | List of labels, optional, defaults to [].

Extra files made available to compilation and linking steps.

Not appended to the default command line arguments, but available to actions. Reference these files manually for instance in the `includes`, and `compile_flags` attributes.

Use this attribute to make intermediary outputs from non-ll targets, for example from `rules_cc` or `filegroup`, available to the rule. | | `hdrs` | List of labels, optional, defaults to [].

Header files for this target.

When including a header file with a nested path, for instance `#include "some/path/myheader.h"`, add `"some/path"` to `includes` to make it visible to the rule.

Unavailable to downstream targets. | | `angled_includes` | List of strings, optional, defaults to [].

Angled include paths, relative to the target workspace.

Useful if you require include prefix stripping for dynamic paths, for instance the ones generated by `bzlmod`. Instead of `compile_flags = ["-Iexternal/mydep.someversion/include"]`, use `angled_includes = ["include"]` to add the path to the workspace automatically.

Expands paths starting with `$(GENERATED)` to the workspace location in the `GENDIR` path.

Unavailable to downstream targets. | -| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | +| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"cuda_nvptx_nvcc"` The CUDA toolchain with `nvcc` as device compiler.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | | `compile_flags` | List of strings, optional, defaults to [].

Flags for the compiler.

Pass a list of strings here. For instance `["-O3", "-std=c++20"]`.

Split flag pairs `-Xclang -somearg` into separate flags `["-Xclang", "-somearg"]`.

Unavailable to downstream targets. | | `compile_string_flags` | List of labels, optional, defaults to [].

Flags for the compiler in the form of `string_flag`s.

Splits the values of each `string_flag` along colons like so:

load("@bazel_skylib//rules:common_settings.bzl", "string_flag")

string_flag(
    name = "myflags",
    build_setting_default = "a:b:c",
)

ll_library(
    # ...
    # Equivalent to `compile_flags = ["a", "b", "c"]`
    compile_string_flags = [":myflags"],
)


Useful for externally configurable build attributes, such as generated flags from Nix environments. | | `defines` | List of strings, optional, defaults to [].

Defines for this target.

Pass a list of strings here. For instance `["MYDEFINE_1", "MYDEFINE_2"]`.

Unavailable to downstream targets. | @@ -120,7 +120,7 @@ Example: | `data` | List of labels, optional, defaults to [].

Extra files made available to compilation and linking steps.

Not appended to the default command line arguments, but available to actions. Reference these files manually for instance in the `includes`, and `compile_flags` attributes.

Use this attribute to make intermediary outputs from non-ll targets, for example from `rules_cc` or `filegroup`, available to the rule. | | `hdrs` | List of labels, optional, defaults to [].

Header files for this target.

When including a header file with a nested path, for instance `#include "some/path/myheader.h"`, add `"some/path"` to `includes` to make it visible to the rule.

Unavailable to downstream targets. | | `angled_includes` | List of strings, optional, defaults to [].

Angled include paths, relative to the target workspace.

Useful if you require include prefix stripping for dynamic paths, for instance the ones generated by `bzlmod`. Instead of `compile_flags = ["-Iexternal/mydep.someversion/include"]`, use `angled_includes = ["include"]` to add the path to the workspace automatically.

Expands paths starting with `$(GENERATED)` to the workspace location in the `GENDIR` path.

Unavailable to downstream targets. | -| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | +| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"cuda_nvptx_nvcc"` The CUDA toolchain with `nvcc` as device compiler.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | | `compile_flags` | List of strings, optional, defaults to [].

Flags for the compiler.

Pass a list of strings here. For instance `["-O3", "-std=c++20"]`.

Split flag pairs `-Xclang -somearg` into separate flags `["-Xclang", "-somearg"]`.

Unavailable to downstream targets. | | `compile_string_flags` | List of labels, optional, defaults to [].

Flags for the compiler in the form of `string_flag`s.

Splits the values of each `string_flag` along colons like so:

load("@bazel_skylib//rules:common_settings.bzl", "string_flag")

string_flag(
    name = "myflags",
    build_setting_default = "a:b:c",
)

ll_library(
    # ...
    # Equivalent to `compile_flags = ["a", "b", "c"]`
    compile_string_flags = [":myflags"],
)


Useful for externally configurable build attributes, such as generated flags from Nix environments. | | `defines` | List of strings, optional, defaults to [].

Defines for this target.

Pass a list of strings here. For instance `["MYDEFINE_1", "MYDEFINE_2"]`.

Unavailable to downstream targets. | @@ -177,7 +177,7 @@ Example: | `data` | List of labels, optional, defaults to [].

Extra files made available to compilation and linking steps.

Not appended to the default command line arguments, but available to actions. Reference these files manually for instance in the `includes`, and `compile_flags` attributes.

Use this attribute to make intermediary outputs from non-ll targets, for example from `rules_cc` or `filegroup`, available to the rule. | | `hdrs` | List of labels, optional, defaults to [].

Header files for this target.

When including a header file with a nested path, for instance `#include "some/path/myheader.h"`, add `"some/path"` to `includes` to make it visible to the rule.

Unavailable to downstream targets. | | `angled_includes` | List of strings, optional, defaults to [].

Angled include paths, relative to the target workspace.

Useful if you require include prefix stripping for dynamic paths, for instance the ones generated by `bzlmod`. Instead of `compile_flags = ["-Iexternal/mydep.someversion/include"]`, use `angled_includes = ["include"]` to add the path to the workspace automatically.

Expands paths starting with `$(GENERATED)` to the workspace location in the `GENDIR` path.

Unavailable to downstream targets. | -| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | +| `compilation_mode` | String, optional, defaults to "cpp".

Enables compilation of heterogeneous single source files.

Prefer this attribute over adding SYCL/HIP/CUDA flags manually in the `compile_flags` and `link_flags`.

See [CUDA and HIP](../guides/cuda_and_hip.md).

`"cpp"` The default C++ toolchain.

`"cuda_nvptx"` The CUDA toolchain.

`"cuda_nvptx_nvcc"` The CUDA toolchain with `nvcc` as device compiler.

`"hip_nvptx"` The HIP toolchain.

`"bootstrap"` The bootstrap toolchain used by internal dependencies of the `ll_toolchain`. | | `compile_flags` | List of strings, optional, defaults to [].

Flags for the compiler.

Pass a list of strings here. For instance `["-O3", "-std=c++20"]`.

Split flag pairs `-Xclang -somearg` into separate flags `["-Xclang", "-somearg"]`.

Unavailable to downstream targets. | | `compile_string_flags` | List of labels, optional, defaults to [].

Flags for the compiler in the form of `string_flag`s.

Splits the values of each `string_flag` along colons like so:

load("@bazel_skylib//rules:common_settings.bzl", "string_flag")

string_flag(
    name = "myflags",
    build_setting_default = "a:b:c",
)

ll_library(
    # ...
    # Equivalent to `compile_flags = ["a", "b", "c"]`
    compile_string_flags = [":myflags"],
)


Useful for externally configurable build attributes, such as generated flags from Nix environments. | | `defines` | List of strings, optional, defaults to [].

Defines for this target.

Pass a list of strings here. For instance `["MYDEFINE_1", "MYDEFINE_2"]`.

Unavailable to downstream targets. | diff --git a/examples/BUILD.bazel b/examples/BUILD.bazel index 975707be..3d7080d6 100644 --- a/examples/BUILD.bazel +++ b/examples/BUILD.bazel @@ -12,7 +12,7 @@ ll_compilation_database( config = ":clang_tidy_config", targets = [ # "//clang_tidy_example", # This one runs separately. - "//cuda_example", + "//cuda_example:nvptx", "//format_example", "//frontend_action_example", "//external_dependency_example", @@ -35,6 +35,7 @@ test_suite( ":amdgpu", ":cpp", ":nvptx", + ":nvptx_nvcc", ], ) @@ -79,13 +80,20 @@ test_suite( name = "nvptx", tags = ["nvptx"], tests = [ - "//cuda_example:test", + "//cuda_example:nvptx_test", "//hip_example:nvptx_test", "//hip_modules_example:nvptx_test", "//hip_rdc_example:nvptx_test", ], ) +test_suite( + name = "nvptx_nvcc", + tags = ["nvptx_nvcc"], + tests = [ + "//cuda_example:nvptx_nvcc_test", + ], +) # AMDGPU tests. test_suite( diff --git a/examples/cuda_example/BUILD.bazel b/examples/cuda_example/BUILD.bazel index d0dfc9f6..e52178c1 100644 --- a/examples/cuda_example/BUILD.bazel +++ b/examples/cuda_example/BUILD.bazel @@ -1,22 +1,37 @@ -load("@rules_ll//ll:defs.bzl", "OFFLOAD_ALL_NVPTX", "ll_binary", "ll_test") - -ll_binary( - name = "cuda_example", - srcs = ["example.cpp"], - compilation_mode = "cuda_nvptx", - compile_flags = OFFLOAD_ALL_NVPTX + [ - "--std=c++20", - ], - visibility = ["@//:__pkg__"], +load( + "@rules_ll//ll:defs.bzl", + "OFFLOAD_ALL_NVPTX", + "ll_binary", + "ll_test", ) -ll_test( - name = "test", - srcs = ["example.cpp"], - compilation_mode = "cuda_nvptx", - compile_flags = OFFLOAD_ALL_NVPTX + [ - "--std=c++20", - ], - tags = ["nvptx"], - visibility = ["@//:__pkg__"], -) +[ + ( + ll_binary( + name = name, + srcs = ["example.cpp"], + compilation_mode = "cuda_{}".format(name), + compile_flags = flags, + visibility = ["@//:__pkg__"], + ), + ll_test( + name = "{}_test".format(name), + size = "small", + srcs = ["example.cpp"], + compilation_mode = "cuda_{}".format(name), + compile_flags = flags, + tags = [name], + visibility = ["@//:__pkg__"], + ), + ) + for (name, flags) in [ + ( + "nvptx", + OFFLOAD_ALL_NVPTX + ["-std=c++20"], + ), + ( + "nvptx_nvcc", + ["--gpu-architecture=sm_50"] + ["-std=c++17"], + ), + ] +] diff --git a/examples/cuda_example/example.cpp b/examples/cuda_example/example.cpp index 272cff28..72c9603d 100644 --- a/examples/cuda_example/example.cpp +++ b/examples/cuda_example/example.cpp @@ -1,5 +1,6 @@ #include #include +#include #include #include "cuda_runtime.h" @@ -9,8 +10,6 @@ constexpr float kInputB = 2.0F; constexpr float kExpectedOutput = 3.0F; constexpr int kDimension = 1 << 20; constexpr auto kThreadsPerBlockX = 128; -constexpr auto kThreadsPerBlockY = 1; -constexpr auto kThreadsPerBlockZ = 1; template constexpr void cuda_assert(const T value) { @@ -30,7 +29,8 @@ void print_device_info() { int count = 0; const cudaError_t err = cudaGetDeviceCount(&count); if (err == cudaErrorInvalidDevice) { - std::cout << "FAIL: invalid device" << '\n'; + std::cerr << "FAIL: invalid device" << '\n'; + return; } std::cout << "Number of devices is " << count << '\n'; @@ -38,11 +38,7 @@ void print_device_info() { cuda_assert(cudaGetDeviceProperties(&device_prop, 0)); std::cout << "System major: " << device_prop.major << '\n'; std::cout << "System minor: " << device_prop.minor << '\n'; - std::cout << "Device name : "; - for (auto character : device_prop.name) { - std::cout << character; - } - std::cout << '\n'; + std::cout << "Device name : " << device_prop.name << '\n'; } auto count_errors(const float *result) -> int { @@ -92,8 +88,8 @@ auto main() -> int { cuda_assert(cudaMemcpy(device_input_b, host_input_b, kDimension * sizeof(float), cudaMemcpyHostToDevice)); - const dim3 grid_dim = dim3(kDimension / kThreadsPerBlockX); - const dim3 block_dim = dim3(kThreadsPerBlockX); + const dim3 grid_dim(kDimension / kThreadsPerBlockX); + const dim3 block_dim(kThreadsPerBlockX); // This is not pretty, but it is close to the HIP implementation. // NOLINTBEGIN cppcoreguidelines-pro-type-reinterpret-cast diff --git a/ll/BUILD.bazel b/ll/BUILD.bazel index a8244a83..005bbaab 100644 --- a/ll/BUILD.bazel +++ b/ll/BUILD.bazel @@ -95,6 +95,9 @@ string_flag( "LL_AMD_LIBRARIES", "LL_CUDA_TOOLKIT", "LL_CUDA_DRIVER", + "LL_CUDA_NVCC", + "LL_CUDA_NVCC_CFLAGS", + "LL_CUDA_NVCC_LDFLAGS", # Unset values default to an empty string. "LL_UNSET", @@ -122,11 +125,25 @@ ll_toolchain( LL_CUDA_DRIVER = select({ ":hip_nvptx": ":LL_CUDA_DRIVER", ":cuda_nvptx": ":LL_CUDA_DRIVER", - "//conditions:default": "LL_UNSET", + ":cuda_nvptx_nvcc": ":LL_CUDA_DRIVER", + "//conditions:default": ":LL_UNSET", + }), + LL_CUDA_NVCC = select({ + ":cuda_nvptx_nvcc": ":LL_CUDA_NVCC", + "//conditions:default": ":LL_UNSET", + }), + LL_CUDA_NVCC_CFLAGS = select({ + ":cuda_nvptx_nvcc": ":LL_CUDA_NVCC_CFLAGS", + "//conditions:default": ":LL_UNSET", + }), + LL_CUDA_NVCC_LDFLAGS = select({ + ":cuda_nvptx_nvcc": ":LL_CUDA_NVCC_LDFLAGS", + "//conditions:default": ":LL_UNSET", }), LL_CUDA_TOOLKIT = select({ ":hip_nvptx": ":LL_CUDA_TOOLKIT", ":cuda_nvptx": ":LL_CUDA_TOOLKIT", + ":cuda_nvptx_nvcc": ":LL_CUDA_TOOLKIT", "//conditions:default": ":LL_UNSET", }), LL_DYNAMIC_LINKER = ":LL_DYNAMIC_LINKER", @@ -148,18 +165,22 @@ ll_toolchain( }), cpp_abihdrs = select({ ":bootstrap": None, + ":cuda_nvptx_nvcc": None, "//conditions:default": "@llvm-project//libcxxabi:headers", }), cpp_abilib = select({ ":bootstrap": None, + ":cuda_nvptx_nvcc": None, "//conditions:default": "@llvm-project//libcxxabi", }), cpp_stdhdrs = select({ ":bootstrap": None, + ":cuda_nvptx_nvcc": None, "//conditions:default": "@llvm-project//libcxx:headers", }), cpp_stdlib = select({ ":bootstrap": None, + ":cuda_nvptx_nvcc": None, "//conditions:default": "@llvm-project//libcxx", }), exec_compatible_with = [ diff --git a/ll/args.bzl b/ll/args.bzl index 9a168d88..f8ecc874 100644 --- a/ll/args.bzl +++ b/ll/args.bzl @@ -120,6 +120,11 @@ def compile_object_args( args = ctx.actions.args() + if ctx.attr.compilation_mode == "cuda_nvptx_nvcc": + args.add("--forward-unknown-to-host-compiler") + args.add("--compiler-bindir") + args.add(toolchain.cpp_driver) + args.add("-fcolor-diagnostics") # Reproducibility. @@ -227,7 +232,24 @@ def compile_object_args( args.add("-xcuda") if toolchain.LL_CUDA_TOOLKIT != "": args.add(toolchain.LL_CUDA_TOOLKIT, format = "--cuda-path=%s") - if ctx.attr.compilation_mode in ["hip_nvptx", "hip_amdgpu"]: + + if ctx.attr.compilation_mode == "cuda_nvptx_nvcc": + args.add("--x") + args.add("cu") + + # Upstream Clang is always "unsupported". + args.add("--allow-unsupported-compiler") + + # Ensure clang doesn't build device code. + args.add("--cuda-host-only") + + if toolchain.LL_CUDA_TOOLKIT != "": + args.add(toolchain.LL_CUDA_TOOLKIT, format = "-I%s/include") + + if ctx.attr.compilation_mode in [ + "hip_nvptx", + "hip_amdgpu", + ]: args.add_all( [ Label("@hip").workspace_root, @@ -275,6 +297,14 @@ def compile_object_args( # 3. Search directories specified via -isystem for quoted and angled # includes. This is not exposed via target attributes. + + if ctx.attr.compilation_mode == "cuda_nvptx_nvcc": + if toolchain.LL_CUDA_NVCC_CFLAGS != "": + args.add_all( + toolchain.LL_CUDA_NVCC_CFLAGS.split(":"), + omit_if_empty = True, + ) + llvm_workspace_root = Label("@llvm-project").workspace_root args.add_all( [ @@ -292,6 +322,7 @@ def compile_object_args( # become system includes. format_each = "-isystem%s", ) + if toolchain.LL_CFLAGS != "": args.add_all(toolchain.LL_CFLAGS.split(":")) @@ -482,6 +513,7 @@ def link_executable_args(ctx, in_files, out_file, mode): if ctx.attr.compilation_mode in [ "cuda_nvptx", + "cuda_nvptx_nvcc", "hip_nvptx", ]: # Both the CUDA driver and the CUDA toolkit contain `libcuda.so`. @@ -530,29 +562,21 @@ def link_executable_args(ctx, in_files, out_file, mode): fail("Invalid linking mode") # Add archives and objects. - # if ctx.attr.depends_on_llvm: - link_files = [ - file - for file in in_files.to_list() - if file.extension in ["a", "o"] - ] - # else: - # link_files = [ - # file - # for file in in_files.to_list() - # if file.extension == "o" - # ] + [ - # file - # for file in ctx.files.deps - # if file.extension == "a" - # ] - # link_files = [ - # file - # for file in in_files.to_list() - # if file.extension in ["a", "o"] - # ] - - args.add_all(link_files) + args.add_all( + [ + file + for file in in_files.to_list() + if file.extension in ["a", "o"] + ], + omit_if_empty = True, + ) + + if ctx.attr.compilation_mode == "cuda_nvptx_nvcc": + if toolchain.LL_CUDA_NVCC_LDFLAGS != "": + args.add_all( + toolchain.LL_CUDA_NVCC_LDFLAGS.split(":"), + omit_if_empty = True, + ) # Link shared libraries in a way that is accessible via `bazel run` and # via manual execution, as long as the relative paths to the shared diff --git a/ll/attributes.bzl b/ll/attributes.bzl index 539040f3..ee502a4b 100644 --- a/ll/attributes.bzl +++ b/ll/attributes.bzl @@ -26,6 +26,8 @@ DEFAULT_ATTRS = { `"cuda_nvptx"` The CUDA toolchain. + `"cuda_nvptx_nvcc"` The CUDA toolchain with `nvcc` as device compiler. + `"hip_nvptx"` The HIP toolchain. `"bootstrap"` The bootstrap toolchain used by internal dependencies of @@ -589,6 +591,19 @@ LL_TOOLCHAIN_ATTRS = { Affects the `cuda_nvptx` and `hip_nvptx` toolchains. """, ), + "LL_CUDA_NVCC": attr.label( + doc = """The path to the `nvcc` compiler. + + Affects the `cuda_nvptx_nvcc` toolchain. Other `*_nvptx` toolchains use + the `c_driver` or `cpp_driver` as device compiler. + """, + ), + "LL_CUDA_NVCC_CFLAGS": attr.label( + doc = "Compile flags used by the `cuda_nvptx_nvcc` toolchain.", + ), + "LL_CUDA_NVCC_LDFLAGS": attr.label( + doc = "Link flags used by the `cuda_nvptx_nvcc` toolchain.", + ), } LL_LIBRARY_ATTRS = dicts.add(DEFAULT_ATTRS, LIBRARY_ATTRS) diff --git a/ll/driver.bzl b/ll/driver.bzl index 314f3ff4..4cef59f8 100644 --- a/ll/driver.bzl +++ b/ll/driver.bzl @@ -16,7 +16,10 @@ def compiler_driver(ctx, in_file): toolchain = ctx.toolchains["//ll:toolchain_type"] driver = None - if in_file.extension in ["c", "S"]: + + if ctx.attr.compilation_mode == "cuda_nvptx_nvcc": + driver = toolchain.LL_CUDA_NVCC + elif in_file.extension in ["c", "S"]: driver = toolchain.c_driver elif in_file.extension in [ "cc", diff --git a/ll/environment.bzl b/ll/environment.bzl index a940253f..c3fdeac9 100644 --- a/ll/environment.bzl +++ b/ll/environment.bzl @@ -20,7 +20,7 @@ def compile_object_environment(ctx): config = ctx.attr.toolchain_configuration[BuildSettingInfo].value toolchain = ctx.toolchains["//ll:toolchain_type"] - if config in ["cpp"]: + if config in ["cpp", "cuda_nvptx_nvcc"]: return { "LINK": toolchain.bitcode_linker.path, "LLD": toolchain.linker.path, diff --git a/ll/toolchain.bzl b/ll/toolchain.bzl index 461afa98..120365db 100644 --- a/ll/toolchain.bzl +++ b/ll/toolchain.bzl @@ -71,6 +71,9 @@ def _ll_toolchain_impl(ctx): LL_AMD_LIBRARIES = ctx.attr.LL_AMD_LIBRARIES[BuildSettingInfo].value, LL_CUDA_TOOLKIT = ctx.attr.LL_CUDA_TOOLKIT[BuildSettingInfo].value, LL_CUDA_DRIVER = ctx.attr.LL_CUDA_DRIVER[BuildSettingInfo].value, + LL_CUDA_NVCC = ctx.attr.LL_CUDA_NVCC[BuildSettingInfo].value, + LL_CUDA_NVCC_CFLAGS = ctx.attr.LL_CUDA_NVCC_CFLAGS[BuildSettingInfo].value, + LL_CUDA_NVCC_LDFLAGS = ctx.attr.LL_CUDA_NVCC_LDFLAGS[BuildSettingInfo].value, ), ] diff --git a/ll/tools.bzl b/ll/tools.bzl index 9be61664..ebd1176a 100644 --- a/ll/tools.bzl +++ b/ll/tools.bzl @@ -16,6 +16,8 @@ def compile_object_tools(ctx): return [ toolchain.bitcode_linker, + toolchain.c_driver, + toolchain.cpp_driver, toolchain.linker, toolchain.linker_executable, toolchain.linker_wrapper, diff --git a/ll/transitions.bzl b/ll/transitions.bzl index eb2f9cf0..dc3d8f82 100644 --- a/ll/transitions.bzl +++ b/ll/transitions.bzl @@ -8,6 +8,7 @@ COMPILATION_MODES = [ "bootstrap", "cpp", "cuda_nvptx", + "cuda_nvptx_nvcc", "hip_amdgpu", "hip_nvptx", ] diff --git a/modules/defaultLlEnv.nix b/modules/defaultLlEnv.nix index e52dcd29..eb67d4bb 100644 --- a/modules/defaultLlEnv.nix +++ b/modules/defaultLlEnv.nix @@ -6,7 +6,7 @@ # variants don't work. , llvmPackages ? pkgs.llvmPackages_17 , nvidia_x11 ? pkgs.linuxKernel.packages.linux_latest_libre.nvidia_x11 -, cudatoolkit ? pkgs.cudaPackages.cudatoolkit +, cudaPackages ? pkgs.cudaPackages , ... }: @@ -23,6 +23,9 @@ let "-rpath=${pkgs.glibc}/lib" "-L${pkgs.libxcrypt}/lib" ]; + + cudatoolkit = cudaPackages.cudatoolkit; + nvcc = cudaPackages.cuda_nvcc; in [ @@ -75,4 +78,17 @@ in # Flags for CUDA dependencies. "LL_CUDA_TOOLKIT=${lib.strings.optionalString pkgs.config.cudaSupport "${cudatoolkit}"}" "LL_CUDA_DRIVER=${lib.strings.optionalString pkgs.config.cudaSupport "${nvidia_x11}"}" + "LL_CUDA_NVCC=${lib.strings.optionalString pkgs.config.cudaSupport "${nvcc}"}/bin/nvcc" + + # TODO(aaronmondal): At the moment `nvcc` doesn't properly work with `libc++`. + # Attempt to get it to work and remove this or integrate it + # more elegantly. + "LL_CUDA_NVCC_CFLAGS=${lib.strings.optionalString pkgs.config.cudaSupport (lib.concatStringsSep ":" [ + "-isystem${pkgs.gcc.cc}/include/c++/13.2.0" + "-isystem${pkgs.gcc.cc}/include/c++/13.2.0/x86_64-unknown-linux-gnu" + ])}" + "LL_CUDA_NVCC_LDFLAGS=${lib.strings.optionalString pkgs.config.cudaSupport (lib.concatStringsSep ":" [ + "-L${pkgs.gcc.cc}/lib" + "-lstdc++" + ])}" ]