# from llvm-project/libcxx/test/configs/cmake-bridge.cfg.in import os, platform, site, sys IS_WINDOWS = platform.system() == 'Windows' config.target_triple = 'x86_64-pc-windows-msvc' if IS_WINDOWS else 'x86_64-unknown-linux-gnu' config.std = 'c++17' config.enable_experimental = False # Auto-detect HIPTHREADS_SOURCE_DIR from this file's location (test/lit.cfg -> hipThreads root) _THIS_DIR = os.path.dirname(os.path.abspath(__file__)) HIPTHREADS_SOURCE_DIR = os.environ.get('HIPTHREADS_SOURCE_DIR', os.path.dirname(_THIS_DIR)) HIPTHREADS_BUILD_DIR = os.environ.get('HIPTHREADS_BUILD_DIR', os.path.join(HIPTHREADS_SOURCE_DIR, 'build')) ROCM_PATH = os.environ.get('ROCM_PATH', '/opt/rocm') if IS_WINDOWS: ROCM_PATH = os.environ.get('ROCM_PATH', os.environ.get('HIP_PATH')) if ROCM_PATH is None: lit_config.fatal('Set HIP_PATH or ROCM_PATH env var to your ROCm install root') else: ROCM_PATH = os.environ.get('ROCM_PATH', '/opt/rocm') # Prefer hipcc on both platforms — clang++ doesn't compile .cpp as HIP # (it uses -x c++ instead of -x hip, missing __host__/__device__ keywords) _hipcc = 'hipcc.exe' if IS_WINDOWS else 'hipcc' compiler_path = os.path.join(ROCM_PATH, 'bin', _hipcc) if not os.path.exists(compiler_path): compiler_path = _hipcc # fall back to PATH config.compiler = compiler_path site.addsitedir(os.path.join(HIPTHREADS_SOURCE_DIR, 'test/utils')) import libcxx.test.format # Basic configuration of the test suite config.name = 'hipThreads' config.test_source_root = os.path.join(HIPTHREADS_SOURCE_DIR, 'test') config.test_format = libcxx.test.format.CxxStandardLibraryTest() config.recursiveExpansionLimit = 10 config.test_exec_root = os.path.join(HIPTHREADS_BUILD_DIR, 'test') # Add substitutions for bootstrapping the test suite configuration import shlex config.substitutions.append(('%{cxx}', shlex.quote(compiler_path))) config.substitutions.append(('%{executor}', f'{sys.executable} {os.path.join(HIPTHREADS_SOURCE_DIR, "test/utils/run.py")}')) config.substitutions.append(('%{test-tools}', os.path.join(HIPTHREADS_SOURCE_DIR, 'test/utils'))) config.substitutions.append(('%{benchmark_flags}', f'-I {os.path.join(HIPTHREADS_BUILD_DIR, "test/benchmarks/google-benchmark/include")} ' f'-L {os.path.join(HIPTHREADS_BUILD_DIR, "test/benchmarks/google-benchmark/lib")} ' f'-L {os.path.join(HIPTHREADS_BUILD_DIR, "test/benchmarks/google-benchmark/lib64")} ' '-l benchmark')) # On Windows, hipcc has no GPU auto-detection; without --offload-arch it defaults # to gfx906, causing linker errors against hipthreads built for a different arch. _hip_arch = os.environ.get('HIP_ARCHITECTURES', '') _arch_flag = f' --offload-arch={_hip_arch}' if _hip_arch else '' # CMake builds hipthreads with the dynamic CRT (/MD), but hipcc defaults to static # CRT (/MT) when invoked directly. This mismatch causes lld-link /failifmismatch # errors, so we force the dynamic CRT to match. _win_runtime = ' -D_DLL -D_MT -Xclang --dependent-lib=msvcrt' if IS_WINDOWS else '' config.substitutions.append(('%{flags}', f'-Wno-unused-command-line-argument -v -fgpu-rdc{_arch_flag}{_win_runtime}' )) rocm_include = os.path.join(ROCM_PATH, 'include') # Determine libhipcxx include path # In CI (venv setup): ROCM_PATH contains site-packages, libhipcxx is at venv base include # Locally: libhipcxx is at ROCM_PATH/include/libhipcxx if 'site-packages' in ROCM_PATH: # Extract venv base path from ROCM_PATH # e.g., /path/to/.venv/lib/python3.12/site-packages/_rocm_sdk_devel -> /path/to/.venv venv_base = ROCM_PATH.split('/lib/python')[0] if '/lib/python' in ROCM_PATH else ROCM_PATH venv_include = os.path.join(venv_base, 'include') # Need both .venv/include (for cuda/) and .venv/include/libhipcxx (for hip/ namespace) libhipcxx_include = f'-isystem "{venv_include}" -isystem "{os.path.join(venv_include, "libhipcxx")}" ' else: # Local installation: libhipcxx is at rocm_include/libhipcxx libhipcxx_include = f'-isystem "{os.path.join(rocm_include, "libhipcxx")}" ' # _CRT_SECURE_NO_WARNINGS: MSVC headers deprecate fopen/sscanf/etc, which becomes # errors under -Werror. # NOMINMAX: prevent min/max macros from breaking std::max() calls in # hipthreads and libhipcxx headers. _win_extra = '-D_CRT_SECURE_NO_WARNINGS -DNOMINMAX ' if IS_WINDOWS else '' config.substitutions.append(('%{compile_flags}', f'{_win_extra}-D_LIBCUDACXX_ALLOW_UNSUPPORTED_ARCHITECTURE -isystem "{rocm_include}" ' f'{libhipcxx_include}' f'-I {os.path.join(HIPTHREADS_SOURCE_DIR, "inc")} ' f'-I {os.path.join(HIPTHREADS_SOURCE_DIR, "test/support")} -flax-vector-conversions=all' )) # Link flags — Windows doesn't need -latomic (GCC-specific). # TheRock embeds /DEFAULTLIB:c++ in amdhip64.lib but doesn't ship libc++ on Windows, # and clang's HIP driver auto-injects -lc++. Suppress both; MSVC's stdlib pulls in # automatically via #pragma comment(lib, "msvcprt") from MSVC headers. if IS_WINDOWS: config.substitutions.append(('%{link_flags}', f'-lhipthreads -L "{os.path.join(HIPTHREADS_BUILD_DIR, "lib")}" ' f'-nostdlib++ -Xlinker /NODEFAULTLIB:c++.lib' )) else: config.substitutions.append(('%{link_flags}', f'-latomic -lhipthreads -L "{os.path.join(HIPTHREADS_BUILD_DIR, "lib")}"' )) # Executor + library path environment # Windows: cmd.exe for mkdir; PATH is also the DLL search path, so prepend # ROCM_PATH/bin so test binaries can find amdhip64_7.dll etc. # Linux: mkdir -p; LD_LIBRARY_PATH forwarded to spawned test binaries via --env # # run.py builds a clean environment for the test process, so shell env vars aren't # inherited. Everything the library reads at runtime must be forwarded via --env # (its nargs="*" means a second --env flag would overwrite the first, so all vars # go under one --env). HIPTHREADS_VCORES_PER_WGP is read by thread.cxx via getenv # and applies on both platforms; it's appended only when set, so when unset # thread.cxx sees getenv == nullptr and applies its default (16). vcores_per_wgp = os.environ.get('HIPTHREADS_VCORES_PER_WGP') vcores_env_arg = f' HIPTHREADS_VCORES_PER_WGP={vcores_per_wgp}' if vcores_per_wgp else '' if IS_WINDOWS: rocm_bin = os.path.join(ROCM_PATH, 'bin') path = rocm_bin + os.pathsep + os.environ.get('PATH', '') config.environment['PATH'] = path config.substitutions.append(('%{exec}', f'cmd /c "if not exist %t.d mkdir %t.d" && %{{executor}} --execdir %t.d ' f'--env PATH="{path}"{vcores_env_arg} -- ' )) else: ld_library_path = os.environ.get('LD_LIBRARY_PATH', os.path.join(ROCM_PATH, 'lib')) config.environment['LD_LIBRARY_PATH'] = ld_library_path config.substitutions.append(('%{exec}', f'mkdir -p %t.d && %{{executor}} --execdir %t.d ' f'--env LD_LIBRARY_PATH={ld_library_path}{vcores_env_arg} -- ' )) import libcxx.test.params, libcxx.test.config libcxx.test.config.configure( libcxx.test.params.DEFAULT_PARAMETERS, libcxx.test.features.DEFAULT_FEATURES, config, lit_config ) # Override flags added by configure() that conflict with HIP compilation new_substitutions = [] for (key, value) in config.substitutions: if key == '%{compile_flags}': # If 'none' is present, replace it; otherwise append 'all' to the end if '-flax-vector-conversions=none' in value: value = value.replace('-flax-vector-conversions=none', '-flax-vector-conversions=all') else: value = value + ' -flax-vector-conversions=all' # Ensure libhipcxx "#warning" doesn't become an error under -Werror if '-Wno-error=#warnings' not in value: value = value + ' -Wno-error=#warnings' if key == '%{flags}': # Put this late so it can override a later-added -Werror if '-Wno-error=#warnings' not in value: value = value + ' -Wno-error=#warnings' new_substitutions.append((key, value)) config.substitutions = new_substitutions # Run tests with: pip install lit && lit -a -v -j 1 test/