Skip to content

[rocPRIM] Add gfx1150 support #3388

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.

Sign up for GitHub

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

Merged
amd-mtrifuno merged 2 commits into from
Jan 9, 2026
Merged

[rocPRIM] Add gfx1150 support #3388

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
0514b10
Add gfx1150 support
amd-mtrifuno Dec 15, 2025
ec640e2
Merge branch 'develop' into users/mtrifuno/rocprim_gfx1150
amd-mtrifuno Jan 8, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion projects/rocprim/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -162,7 +162,7 @@ else()
)
else()
rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
TARGETS "gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201"
TARGETS "gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201"
)
endif()
set(GPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for" FORCE)
Expand Down
2 changes: 1 addition & 1 deletion projects/rocprim/docs/install/rocPRIM-build-install-linux.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,7 @@ Use the appropriate CMake directive:
* ``BUILD_BENCHMARK``: Set to ``ON`` to build benchmarking tests. ``OFF`` by default.
* ``BENCHMARK_CONFIG_TUNING``: Set to ``ON`` to find the best kernel configuration parameters for benchmarking. Turning this on might increase compilation time significantly. ``OFF`` by default.
* ``BENCHMARK_USE_AMDSMI``: Set to ``ON`` to let benchmarks use AMD SMI to output more GPU statistics. ``OFF`` by default.
* ``AMDGPU_TARGETS``: Set this to build the library, examples, tests, examples, and benchmarks for specific architecture targets. When not set, the examples, tests, and benchmarks are built for gfx803, gfx900:xnack-, gfx906:xnack-, gfx908:xnack-, gfx90a:xnack-, gfx90a:xnack+, gfx942;gfx950, gfx1030, gfx1100, gfx1101, gfx1102, gfx1151, gfx1152, gfx1153, gfx1200, and gfx1201 architectures. The list of targets must be separated by a semicolon (``;``).
* ``AMDGPU_TARGETS``: Set this to build the library, examples, tests, examples, and benchmarks for specific architecture targets. When not set, the examples, tests, and benchmarks are built for gfx803, gfx900:xnack-, gfx906:xnack-, gfx908:xnack-, gfx90a:xnack-, gfx90a:xnack+, gfx942;gfx950, gfx1030, gfx1100, gfx1101, gfx1102, gfx1150, gfx1151, gfx1152, gfx1153, gfx1200, and gfx1201 architectures. The list of targets must be separated by a semicolon (``;``).
* ``AMDGPU_TEST_TARGETS``: Set this to build tests for a subset of the architectures specified by ``AMDGPU_TARGETS``. When set, copies of the same test will be generated for each of the architectures listed. These tests can be run using ``ctest -R "TARGET_ARCHITECTURE"``. The list of targets must be separated by a semicolon (``;``).
* ``USE_SYSTEM_LIB``: Set to ``ON`` to use the installed ``ROCm`` libraries when building the tests. Off by default. For this option to take effect, ``BUILD_TEST`` must be ``ON``.
* ``ONLY_INSTALL``: Set to ``ON`` to ignore any example, test, or benchmark build instructions. ``OFF`` by default.
Expand Down
2 changes: 1 addition & 1 deletion projects/rocprim/rmake.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,7 @@ def parse_args():
Checks build arguments
""")

default_gpus = 'gfx906:xnack-,gfx1030,gfx1100,gfx1101,gfx1102,gfx1151,gfx1200,gfx1201'
default_gpus = 'gfx906:xnack-,gfx1030,gfx1100,gfx1101,gfx1102,gfx1150,gfx1151,gfx1200,gfx1201'
Comment thread
amd-mtrifuno marked this conversation as resolved.

parser.add_argument('-g', '--debug', required=False, default=False, action='store_true',
help='Generate Debug build (default: False)')
Expand Down
1 change: 1 addition & 0 deletions projects/rocprim/rocprim/include/rocprim/config.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -169,6 +169,7 @@
__builtin_amdgcn_processor_is("gfx1100") || __builtin_amdgcn_processor_is("gfx1101") \
|| __builtin_amdgcn_processor_is("gfx1102") \
|| __builtin_amdgcn_processor_is("gfx1103") \
|| __builtin_amdgcn_processor_is("gfx1150") \
|| __builtin_amdgcn_processor_is("gfx1152") \
|| __builtin_amdgcn_processor_is("gfx1153") \
|| __builtin_amdgcn_processor_is("gfx11-generic")
Expand Down
4 changes: 4 additions & 0 deletions projects/rocprim/rocprim/include/rocprim/device/config_types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -173,6 +173,7 @@ enum class target_arch : unsigned int
gfx1030 = 1030,
gfx1100 = 1100,
gfx1102 = 1102,
gfx1150 = 1150,
gfx1152 = 1152,
gfx1153 = 1153,
gfx1200 = 1200,
Expand Down Expand Up @@ -224,6 +225,7 @@ constexpr auto target_arch_descriptors = std::array{
X(gfx1030),
X(gfx1100),
X(gfx1102),
X(gfx1150),
X(gfx1152),
X(gfx1153),
X(gfx1200),
Expand Down Expand Up @@ -271,6 +273,7 @@ constexpr arch::wavefront::target arch_wavefront_size(const target_arch target_a
case target_arch::gfx1030: return arch::wavefront::target::size32;
case target_arch::gfx1100: return arch::wavefront::target::size32;
case target_arch::gfx1102: return arch::wavefront::target::size32;
case target_arch::gfx1150: return arch::wavefront::target::size32;
case target_arch::gfx1152: return arch::wavefront::target::size32;
case target_arch::gfx1153: return arch::wavefront::target::size32;
case target_arch::gfx1200: return arch::wavefront::target::size32;
Expand Down Expand Up @@ -505,6 +508,7 @@ auto dispatch_target_arch([[maybe_unused]] const target_arch target_arch)
X(gfx1030);
X(gfx1100);
X(gfx1102);
X(gfx1150);
X(gfx1152);
X(gfx1153);
X(gfx1200);
Expand Down