hip - don't use JiT options that chipStar cannot process

Author: jeremylt

backends/hip/ceed-hip-compile.cpp

@@ -10,6 +10,7 @@
 #include <ceed.h>
 #include <ceed/backend.h>
 #include <ceed/jit-tools.h>
+#include <ceed/jit-source/hip/hip-chipstar.h>
 #include <stdarg.h>
 #include <string.h>
 #include <hip/hiprtc.h>
@@ -37,7 +38,6 @@ static int CeedCompileCore_Hip(Ceed ceed, const char *source, const bool throw_e
                                const CeedInt num_defines, va_list args) {
   size_t                 ptx_size;
   char                  *ptx;
-  const int              num_opts            = 4;
   CeedInt                num_jit_source_dirs = 0, num_jit_defines = 0;
   const char           **opts;
   int                    runtime_version;
@@ -78,13 +78,24 @@ static int CeedCompileCore_Hip(Ceed ceed, const char *source, const bool throw_e
 
   // Non-macro options
   CeedCallBackend(CeedCalloc(num_opts, &opts));
+#if CEED_HIP_USE_CHIPSTAR
+  const int num_opts = 1;
+
+  opts[0] = "-DCEED_RUNNING_JIT_PASS=1";
+#else
+  const int num_opts = 4;
+
   opts[0] = "-default-device";
-  CeedCallBackend(CeedGetData(ceed, (void **)&ceed_data));
-  CeedCallHip(ceed, hipGetDeviceProperties(&prop, ceed_data->device_id));
-  std::string arch_arg = "--gpu-architecture=" + std::string(prop.gcnArchName);
-  opts[1]              = arch_arg.c_str();
-  opts[2]              = "-munsafe-fp-atomics";
-  opts[3]              = "-DCEED_RUNNING_JIT_PASS=1";
+  {
+    CeedCallBackend(CeedGetData(ceed, (void **)&ceed_data));
+    CeedCallHip(ceed, hipGetDeviceProperties(&prop, ceed_data->device_id));
+    std::string arch_arg = "--gpu-architecture=" + std::string(prop.gcnArchName);
+
+    opts[1] = arch_arg.c_str();
+  }
+  opts[2] = "-munsafe-fp-atomics";
+  opts[3] = "-DCEED_RUNNING_JIT_PASS=1";
+#endif
   // Additional include dirs
   {
     const char **jit_source_dirs;

include/ceed/jit-source/hip/hip-chipstar.h

@@ -0,0 +1,17 @@
+// Copyright (c) 2017-2026, Lawrence Livermore National Security, LLC and other CEED contributors.
+// All Rights Reserved. See the top-level LICENSE and NOTICE files for details.
+//
+// SPDX-License-Identifier: BSD-2-Clause
+//
+// This file is part of CEED:  http://github.com/ceed
+
+/// @file
+/// Internal header for HIP chipStar detection
+
+// If we are using Chipstar, then we have to ensure all threads have the same workloads
+//   and hit __syncthreads() at the same places/number of times
+#ifdef __HIP_PLATFORM_SPIRV__
+#define CEED_HIP_USE_CHIPSTAR true
+#else
+#define CEED_HIP_USE_CHIPSTAR false
+#endif

include/ceed/jit-source/hip/hip-jit.h

@@ -13,12 +13,5 @@
 #define CeedPragmaSIMD
 #define CEED_Q_VLA 1
 
-// If we are using Chipstar, then we have to ensure all threads have the same workloads
-//   and hit __syncthreads() at the same places/number of times
-#ifdef __HIP_PLATFORM_SPIRV__
-#define CEED_HIP_USE_CHIPSTAR true
-#else
-#define CEED_HIP_USE_CHIPSTAR false
-#endif
-
+#include "hip-chipstar.h"
 #include "hip-types.h"