Add chipStar (SPIR-V) support for HIP backends

Makefile

@@ -77,6 +77,9 @@ endif
 ifeq (,$(filter-out undefined default,$(origin ARFLAGS)))
   ARFLAGS = $(if $(DARWIN),cr,crD)
 endif
+# Often /opt/rocm, but sometimes present on machines that don't support HIP
+ROCM_DIR ?= ${HIP_DIR}
+HIP_ARCH ?=
 NVCC ?= $(CUDA_DIR)/bin/nvcc
 NVCC_CXX ?= $(CXX)
 HIPCC ?= $(ROCM_DIR)/bin/hipcc
@@ -88,6 +91,23 @@ ifneq ($(EMSCRIPTEN),)
   EM_LDFLAGS = -s TOTAL_MEMORY=256MB
 endif
 
+# Detect HIP platform (chipStar/SPIR-V vs AMD ROCm) only when ROCM_DIR is set
+# and hipconfig is actually present, so machines without HIP can build other backends.
+HIP_LIB_NAME = amdhip64
+ifneq ($(ROCM_DIR),)
+  ifneq ($(wildcard $(ROCM_DIR)/bin/hipconfig),)
+    HIP_CONFIG_RES := $(shell $(ROCM_DIR)/bin/hipconfig 2>/dev/null)
+    ifneq (,$(findstring __HIP_PLATFORM_SPIRV__,$(HIP_CONFIG_RES)))
+      HIP_LIB_NAME = CHIP
+    else ifneq (,$(findstring __HIP_PLATFORM_HCC__,$(HIP_CONFIG_RES)))
+      HIP_LIB_NAME = amdhip64
+    else ifneq (,$(findstring __HIP_PLATFORM_AMD__,$(HIP_CONFIG_RES)))
+      HIP_LIB_NAME = amdhip64
+    else
+      $(error "HIP platform not supported")
+    endif
+  endif
+endif
 # ASAN must be left empty if you don't want to use it
 ASAN ?=
 
@@ -555,16 +575,23 @@ ifneq ($(CUDA_LIB_DIR),)
 endif
 
 # HIP Backends
-HIP_LIB_DIR := $(wildcard $(foreach d,lib lib64,$(ROCM_DIR)/$d/libamdhip64.${SO_EXT}))
+HIP_LIB_DIR := $(wildcard $(foreach d,lib lib64,$(ROCM_DIR)/$d/lib${HIP_LIB_NAME}.${SO_EXT}))
 HIP_LIB_DIR := $(patsubst %/,%,$(dir $(firstword $(HIP_LIB_DIR))))
 HIP_BACKENDS = /gpu/hip/ref /gpu/hip/shared /gpu/hip/gen
 ifneq ($(HIP_LIB_DIR),)
-  HIPCONFIG_CPPFLAGS := $(subst =,,$(shell $(ROCM_DIR)/bin/hipconfig -C))
-  $(hip-all.c:%.c=$(OBJDIR)/%.o) $(hip-all.c:%=%.tidy): CPPFLAGS += $(HIPCONFIG_CPPFLAGS)
+  HIPCONFIG_CPPFLAGS := $(shell $(ROCM_DIR)/bin/hipconfig -C)
+  # chipStar hipconfig -C includes clang-only flags (--target=, --offload=, -nohipwrapperinc, --hip-path=);
+  # strip those out for gcc-compiled C sources, keeping -D/-I/-include flags
+  ifeq ($(HIP_LIB_NAME),CHIP)
+    HIPCONFIG_CPPFLAGS_C := $(filter-out --offload% -nohipwrapperinc --hip-path% --target%,$(HIPCONFIG_CPPFLAGS)) -I$(ROCM_DIR)/include
+  else
+    HIPCONFIG_CPPFLAGS_C := $(HIPCONFIG_CPPFLAGS)
+  endif
+  $(hip-all.c:%.c=$(OBJDIR)/%.o) $(hip-all.c:%=%.tidy): CPPFLAGS += $(HIPCONFIG_CPPFLAGS_C)
   ifneq ($(CXX), $(HIPCC))
-    $(hip-all.cpp:%.cpp=$(OBJDIR)/%.o) $(hip-all.cpp:%=%.tidy): CPPFLAGS += $(HIPCONFIG_CPPFLAGS)
+    $(hip-all.cpp:%.cpp=$(OBJDIR)/%.o) $(hip-all.cpp:%=%.tidy): CPPFLAGS += $(HIPCONFIG_CPPFLAGS_C)
   endif
-  PKG_LIBS += -L$(abspath $(HIP_LIB_DIR)) -lamdhip64 -lhipblas
+  PKG_LIBS += -L$(abspath $(HIP_LIB_DIR)) -l${HIP_LIB_NAME} -lhipblas
   LIBCEED_CONTAINS_CXX = 1
   libceed.c     += $(hip-all.c)
   libceed.cpp   += $(hip-all.cpp)
@@ -640,7 +667,12 @@ endif
 
 pkgconfig-libs-private = $(PKG_LIBS)
 ifeq ($(LIBCEED_CONTAINS_CXX),1)
-  $(libceeds) : LINK = $(CXX)
+  ifneq ($(SYCL_LIB_DIR),)
+    $(libceeds) : LINK = $(SYCLCXX)
+    $(libceeds) : CEED_LDFLAGS += $(SYCLFLAGS)
+  else
+    $(libceeds) : LINK = $(CXX)
+  endif
   ifeq ($(STATIC),1)
     $(examples) $(tests) : CEED_LDLIBS += $(LIBCXX)
     pkgconfig-libs-private += $(LIBCXX)
@@ -798,7 +830,7 @@ NPROC_POOL ?= 1
 export NPROC_POOL
 
 run-% : $(OBJDIR)/%
-	@$(PYTHON) tests/junit.py --mode tap --ceed-backends $(BACKENDS) --nproc $(NPROC_TEST) --pool-size $(NPROC_POOL) --search '$(subsearch)' $(<:$(OBJDIR)/%=%)
+	@$(PYTHON) tests/junit.py --mode tap --ceed-backends $(BACKENDS) --nproc $(NPROC_TEST) --pool-size $(NPROC_POOL) --search '$(subsearch)' $(JUNIT_ARGS) $(<:$(OBJDIR)/%=%)
 
 # The test and prove targets can be controlled via pattern searches.  The
 # default is to run tests and those examples that have no external dependencies.
@@ -812,6 +844,8 @@ realsearch = $(search:%=%%)
 matched = $(foreach pattern,$(realsearch),$(filter $(OBJDIR)/$(pattern),$(tests) $(allexamples)))
 subsearch ?= .*
 JUNIT_BATCH ?= ''
+# Extra arguments forwarded to tests/junit.py (e.g. --env CHIP_LOGLEVEL=crit)
+JUNIT_ARGS ?=
 
 # Test core libCEED
 test : $(matched:$(OBJDIR)/%=run-%)
@@ -825,15 +859,15 @@ ctc-% : $(ctests);@$(foreach tst,$(ctests),$(tst) /cpu/$*;)
 # https://testanything.org/tap-specification.html
 prove : $(matched)
 	$(info Testing backends: $(BACKENDS))
-	$(PROVE) $(PROVE_OPTS) --exec '$(PYTHON) tests/junit.py' $(matched:$(OBJDIR)/%=%) :: --mode tap --ceed-backends $(BACKENDS) --nproc $(NPROC_TEST) --pool-size $(NPROC_POOL) --search '$(subsearch)'
+	$(PROVE) $(PROVE_OPTS) --exec '$(PYTHON) tests/junit.py' $(matched:$(OBJDIR)/%=%) :: --mode tap --ceed-backends $(BACKENDS) --nproc $(NPROC_TEST) --pool-size $(NPROC_POOL) --search '$(subsearch)' $(JUNIT_ARGS)
 # Run prove target in parallel
 prv : ;@$(MAKE) $(MFLAGS) V=$(V) prove
 
 prove-all :
 	+$(MAKE) prove realsearch=%
 
 junit-% : $(OBJDIR)/%
-	@printf "  %10s %s\n" TEST $(<:$(OBJDIR)/%=%); $(PYTHON) tests/junit.py --ceed-backends $(BACKENDS) --nproc $(NPROC_TEST) --pool-size $(NPROC_POOL) --search '$(subsearch)' --junit-batch $(JUNIT_BATCH) $(<:$(OBJDIR)/%=%)
+	@printf "  %10s %s\n" TEST $(<:$(OBJDIR)/%=%); $(PYTHON) tests/junit.py --ceed-backends $(BACKENDS) --nproc $(NPROC_TEST) --pool-size $(NPROC_POOL) --search '$(subsearch)' --junit-batch $(JUNIT_BATCH) $(JUNIT_ARGS) $(<:$(OBJDIR)/%=%)
 
 junit : $(matched:$(OBJDIR)/%=junit-%)
 

backends/hip-gen/ceed-hip-gen-operator-build.cpp

@@ -12,6 +12,7 @@
 #include <ceed/gen-tools.h>
 #include <ceed/jit-tools.h>
 
+#include <cstring>
 #include <iostream>
 #include <sstream>
 #include <string>

tests/junit.py

@@ -2,6 +2,26 @@
 from junit_common import *
 
 
+def parse_env_assignment(arg: str) -> Tuple[str, str]:
+    """Parse a KEY=VAL string from the --env flag.
+
+    Args:
+        arg (str): The KEY=VAL string
+
+    Returns:
+        Tuple[str, str]: (key, value) pair
+
+    Raises:
+        argparse.ArgumentTypeError: if `arg` is not in KEY=VAL form
+    """
+    if '=' not in arg:
+        raise argparse.ArgumentTypeError(f"--env expects KEY=VAL, got {arg!r}")
+    key, _, value = arg.partition('=')
+    if not key:
+        raise argparse.ArgumentTypeError(f"--env KEY must be non-empty, got {arg!r}")
+    return key, value
+
+
 def create_argparser() -> argparse.ArgumentParser:
     """Creates argument parser to read command line arguments
 
@@ -30,6 +50,11 @@ def create_argparser() -> argparse.ArgumentParser:
                         help='Search string to filter tests, using `re` package format')
     parser.add_argument('-v', '--verbose', action='store_true', default=False,
                         help='print details for all runs, not just failures')
+    parser.add_argument('--env', dest='extra_env', action='append', default=[], metavar='KEY=VAL',
+                        type=parse_env_assignment,
+                        help='Set an environment variable for test subprocesses; may be repeated. '
+                             'Useful for backend-specific runtime knobs (e.g. --env CHIP_LOGLEVEL=crit '
+                             'to silence chipStar runtime messages on stderr).')
     parser.add_argument('test', help='Test executable', nargs='?')
 
     return parser
@@ -198,6 +223,11 @@ def check_allowed_stdout(self, test: str) -> bool:
 if __name__ == '__main__':
     args = create_argparser().parse_args()
 
+    # Apply --env KEY=VAL settings to the parent environment so worker
+    # processes (which copy os.environ in init_process) inherit them.
+    for key, value in args.extra_env:
+        os.environ[key] = value
+
     result: TestSuite = run_tests(
         args.test,
         args.ceed_backends,