From 2c83ba8995ab739a7e6111cfe260301ffbbe4b53 Mon Sep 17 00:00:00 2001 From: Ajay Panyala Date: Wed, 26 Aug 2026 13:02:15 -0400 Subject: [PATCH 1/2] HIP: implement missing Fxc/vxc_fxc device stubs and complete builtin functional coverage The HIP backend (unlike CUDA and SYCL) never implemented Fxc/vxc_fxc (second XC derivative) device evaluation, and separately, builtin.hip's explicit template instantiation list only covered 28 of the 102 functional types CUDA's builtin.cu instantiates. Both left vtable symbols unresolved at link time for any executable that actually links (as opposed to just archiving) against libexchcxx built for HIP -- e.g. ordinary SVWN5/M06-2X calculations, or anything touching Fxc contraction. - hip/builtin.hip: add generic (KernelType-templated) stub definitions for the LDA/GGA/MGGA device_eval_{fxc,vxc_fxc}[_inc]_helper_* function templates, throwing std::runtime_error at call time. Wire them into the {LDA,GGA,MGGA}_GENERATE_DEVICE_HELPERS macros so every instantiated functional type gets them. Also complete the type list itself: append the 74 functional types (mostly GGA/meta-GGA -- PBE, M06 family, TPSS, SCAN-adjacent, etc.) present in cuda/builtin.cu but missing here entirely, so their ordinary EXC/VXC evaluation -- which the shared kernel_traits math already supports -- actually gets compiled for HIP. - hip/libxc_device.hip: add the missing LibxcKernelImpl eval_fxc_device_/eval_vxc_fxc_device_ (and _inc_ variants) overrides using the existing UNUSED_DEVICE_(INC_)INTERFACE_GENERATOR macros -- the same idiom already used here for eval_*_inc_device_ and in builtin_kernel.cxx for disabled kernels. - hip/xc_functional_device.hip: stub XCFunctional::eval_fxc_device/ eval_vxc_fxc_device (LDA/GGA/MGGA) the same way; these aggregate over each constituent kernel's (now-stubbed) eval_*_fxc_device. --- src/hip/builtin.hip | 217 ++++++++++++++++++++++++++++++- src/hip/libxc_device.hip | 33 ++++- src/hip/xc_functional_device.hip | 26 ++++ 3 files changed, 270 insertions(+), 6 deletions(-) diff --git a/src/hip/builtin.hip b/src/hip/builtin.hip index a1a8f68..bdaf8ca 100644 --- a/src/hip/builtin.hip +++ b/src/hip/builtin.hip @@ -53,6 +53,7 @@ #include #include #include +#include namespace ExchCXX { namespace detail { @@ -884,6 +885,116 @@ MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ) { } +// ExchCXX's HIP backend does not implement Fxc/vxc_fxc device evaluation +// (unlike CUDA and SYCL). Rather than leaving BuiltinKernelImpl's +// vtable (eval_fxc_device/eval_vxc_fxc_device, declared in kernel_type.hpp) +// unresolved at link time for every functional, provide a generic stub that +// throws at runtime; ordinary exc/vxc evaluation is unaffected. The +// underlying math (kernel_traits::eval_fxc_*) already exists and is +// shared with the host/CUDA/SYCL backends, so a real port -- mirroring +// cuda/builtin.cu's device_eval_*_helper_* kernels -- remains possible as a +// follow-up; that port must account for HIP-specific numerics differences +// already present in this file (e.g. the fmax(rho,0.) clamps above) rather +// than being a verbatim copy. +template +LDA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_unpolar ) { + throw std::runtime_error("LDA Fxc device evaluation NYI for HIP backend"); +} +template +LDA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_polar ) { + throw std::runtime_error("LDA Fxc device evaluation NYI for HIP backend"); +} +template +LDA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_unpolar ) { + throw std::runtime_error("LDA Fxc device evaluation NYI for HIP backend"); +} +template +LDA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_polar ) { + throw std::runtime_error("LDA Fxc device evaluation NYI for HIP backend"); +} +template +LDA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_unpolar ) { + throw std::runtime_error("LDA Fxc device evaluation NYI for HIP backend"); +} +template +LDA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_polar ) { + throw std::runtime_error("LDA Fxc device evaluation NYI for HIP backend"); +} +template +LDA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_unpolar ) { + throw std::runtime_error("LDA Fxc device evaluation NYI for HIP backend"); +} +template +LDA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_polar ) { + throw std::runtime_error("LDA Fxc device evaluation NYI for HIP backend"); +} + +template +GGA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_unpolar ) { + throw std::runtime_error("GGA Fxc device evaluation NYI for HIP backend"); +} +template +GGA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_polar ) { + throw std::runtime_error("GGA Fxc device evaluation NYI for HIP backend"); +} +template +GGA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_unpolar ) { + throw std::runtime_error("GGA Fxc device evaluation NYI for HIP backend"); +} +template +GGA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_polar ) { + throw std::runtime_error("GGA Fxc device evaluation NYI for HIP backend"); +} +template +GGA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_unpolar ) { + throw std::runtime_error("GGA Fxc device evaluation NYI for HIP backend"); +} +template +GGA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_polar ) { + throw std::runtime_error("GGA Fxc device evaluation NYI for HIP backend"); +} +template +GGA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_unpolar ) { + throw std::runtime_error("GGA Fxc device evaluation NYI for HIP backend"); +} +template +GGA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_polar ) { + throw std::runtime_error("GGA Fxc device evaluation NYI for HIP backend"); +} + +template +MGGA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_unpolar ) { + throw std::runtime_error("MGGA Fxc device evaluation NYI for HIP backend"); +} +template +MGGA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_polar ) { + throw std::runtime_error("MGGA Fxc device evaluation NYI for HIP backend"); +} +template +MGGA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_unpolar ) { + throw std::runtime_error("MGGA Fxc device evaluation NYI for HIP backend"); +} +template +MGGA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_polar ) { + throw std::runtime_error("MGGA Fxc device evaluation NYI for HIP backend"); +} +template +MGGA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_unpolar ) { + throw std::runtime_error("MGGA Fxc device evaluation NYI for HIP backend"); +} +template +MGGA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_polar ) { + throw std::runtime_error("MGGA Fxc device evaluation NYI for HIP backend"); +} +template +MGGA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_unpolar ) { + throw std::runtime_error("MGGA Fxc device evaluation NYI for HIP backend"); +} +template +MGGA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_polar ) { + throw std::runtime_error("MGGA Fxc device evaluation NYI for HIP backend"); +} + #define LDA_GENERATE_DEVICE_HELPERS(KERN) \ template LDA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_unpolar ); \ template LDA_EXC_VXC_GENERATOR_DEVICE( device_eval_exc_vxc_helper_unpolar ); \ @@ -892,7 +1003,15 @@ MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ) { template LDA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_polar ); \ template LDA_EXC_VXC_GENERATOR_DEVICE( device_eval_exc_vxc_helper_polar ); \ template LDA_EXC_INC_GENERATOR_DEVICE( device_eval_exc_inc_helper_polar ); \ - template LDA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ); + template LDA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ); \ + template LDA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_unpolar ); \ + template LDA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_polar ); \ + template LDA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_unpolar ); \ + template LDA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_polar ); \ + template LDA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_unpolar ); \ + template LDA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_polar ); \ + template LDA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_unpolar ); \ + template LDA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_polar ); #define GGA_GENERATE_DEVICE_HELPERS(KERN) \ template GGA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_unpolar ); \ @@ -902,7 +1021,15 @@ MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ) { template GGA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_polar ); \ template GGA_EXC_VXC_GENERATOR_DEVICE( device_eval_exc_vxc_helper_polar ); \ template GGA_EXC_INC_GENERATOR_DEVICE( device_eval_exc_inc_helper_polar ); \ - template GGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ); + template GGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ); \ + template GGA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_unpolar ); \ + template GGA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_polar ); \ + template GGA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_unpolar ); \ + template GGA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_polar ); \ + template GGA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_unpolar ); \ + template GGA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_polar ); \ + template GGA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_unpolar ); \ + template GGA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_polar ); #define MGGA_GENERATE_DEVICE_HELPERS(KERN) \ template MGGA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_unpolar ); \ @@ -912,11 +1039,20 @@ MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ) { template MGGA_EXC_GENERATOR_DEVICE( device_eval_exc_helper_polar ); \ template MGGA_EXC_VXC_GENERATOR_DEVICE( device_eval_exc_vxc_helper_polar ); \ template MGGA_EXC_INC_GENERATOR_DEVICE( device_eval_exc_inc_helper_polar ); \ - template MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ); + template MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ); \ + template MGGA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_unpolar ); \ + template MGGA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_polar ); \ + template MGGA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_unpolar ); \ + template MGGA_VXC_FXC_GENERATOR_DEVICE( device_eval_vxc_fxc_helper_polar ); \ + template MGGA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_unpolar ); \ + template MGGA_FXC_INC_GENERATOR_DEVICE( device_eval_fxc_inc_helper_polar ); \ + template MGGA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_unpolar ); \ + template MGGA_VXC_FXC_INC_GENERATOR_DEVICE( device_eval_vxc_fxc_inc_helper_polar ); LDA_GENERATE_DEVICE_HELPERS( BuiltinSlaterExchange ); LDA_GENERATE_DEVICE_HELPERS( BuiltinVWN3 ); LDA_GENERATE_DEVICE_HELPERS( BuiltinVWN_RPA ); +LDA_GENERATE_DEVICE_HELPERS( BuiltinVWN ); LDA_GENERATE_DEVICE_HELPERS( BuiltinPW91_LDA ); LDA_GENERATE_DEVICE_HELPERS( BuiltinPW91_LDA_MOD ); LDA_GENERATE_DEVICE_HELPERS( BuiltinPW91_LDA_RPA ); @@ -928,6 +1064,36 @@ GGA_GENERATE_DEVICE_HELPERS( BuiltinLYP ); GGA_GENERATE_DEVICE_HELPERS( BuiltinPBE_X ); GGA_GENERATE_DEVICE_HELPERS( BuiltinRevPBE_X ); GGA_GENERATE_DEVICE_HELPERS( BuiltinPBE_C ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinB97_D ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinITYH_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinITYH_X_033 ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinITYH_X_015 ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinP86_C ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinP86VWN_FT_C ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinPW91_C ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinPBE_SOL_C ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinBMK_C ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinN12_C ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinN12_SX_C ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinSOGGA11_X_C ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinPW91_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinMPW91_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinOPTX_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinRPBE_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinSOGGA11_X_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinPW86_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinWB97_XC ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinWB97X_XC ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinWB97X_V_XC ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinWB97X_D_XC ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinWB97X_D3_XC ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinHJS_PBE_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinLCwPBE_wPBEh_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinLRCwPBE_HJS_PBE_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinLRCwPBEh_HJS_PBE_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinWPBEh_X_default0 ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinHSE03_wPBEh_X ); +GGA_GENERATE_DEVICE_HELPERS( BuiltinHSE06_wPBEh_X ); MGGA_GENERATE_DEVICE_HELPERS( BuiltinSCAN_X ); @@ -935,7 +1101,49 @@ MGGA_GENERATE_DEVICE_HELPERS( BuiltinSCAN_C ); MGGA_GENERATE_DEVICE_HELPERS( BuiltinR2SCAN_X ); MGGA_GENERATE_DEVICE_HELPERS( BuiltinR2SCAN_C ); MGGA_GENERATE_DEVICE_HELPERS( BuiltinFT98_X ); - +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM062X_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM062X_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinPKZB_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinPKZB_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinTPSS_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinRevTPSS_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM06_L_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM06_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM06_HF_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinRevM06_L_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM06_SX_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM06_L_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM06_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM06_HF_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinRevM06_L_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM06_SX_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM05_2X_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM05_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM08_HX_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM08_SO_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinCF22D_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM11_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinMN12_L_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinMN12_SX_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinMN15_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinMN15_L_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinTPSS_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinRevTPSS_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinRSCAN_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinBC95_C ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinMBEEF_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinRSCAN_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinBMK_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM08_HX_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM08_SO_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinMN12_L_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinMN15_L_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinMN15_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinCF22D_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinMN12_SX_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM11_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM05_X ); +MGGA_GENERATE_DEVICE_HELPERS( BuiltinM05_2X_X ); MGGA_GENERATE_DEVICE_HELPERS( BuiltinPC07_K ); MGGA_GENERATE_DEVICE_HELPERS( BuiltinPC07OPT_K ); @@ -949,6 +1157,7 @@ LDA_GENERATE_DEVICE_HELPERS( BuiltinEPC17_2 ) LDA_GENERATE_DEVICE_HELPERS( BuiltinEPC18_1 ) LDA_GENERATE_DEVICE_HELPERS( BuiltinEPC18_2 ) + } } diff --git a/src/hip/libxc_device.hip b/src/hip/libxc_device.hip index a133c1b..4eabf29 100644 --- a/src/hip/libxc_device.hip +++ b/src/hip/libxc_device.hip @@ -306,10 +306,39 @@ UNUSED_DEVICE_INC_INTERFACE_GENERATOR( GGA, EXC, LibxcKernelImpl::eval_exc_inc_device_, const ) UNUSED_DEVICE_INC_INTERFACE_GENERATOR( GGA, EXC_VXC, LibxcKernelImpl::eval_exc_vxc_inc_device_, const ) -UNUSED_DEVICE_INC_INTERFACE_GENERATOR( MGGA, EXC, +UNUSED_DEVICE_INC_INTERFACE_GENERATOR( MGGA, EXC, LibxcKernelImpl::eval_exc_inc_device_, const ) -UNUSED_DEVICE_INC_INTERFACE_GENERATOR( MGGA, EXC_VXC, +UNUSED_DEVICE_INC_INTERFACE_GENERATOR( MGGA, EXC_VXC, LibxcKernelImpl::eval_exc_vxc_inc_device_, const ) +// ExchCXX's HIP backend does not implement Fxc/vxc_fxc device evaluation +// for Libxc-backed kernels (unlike CUDA and SYCL); disable rather than +// leave the LibxcKernelImpl vtable unresolved at link time. +UNUSED_DEVICE_INTERFACE_GENERATOR( LDA, FXC, + LibxcKernelImpl::eval_fxc_device_, const ) +UNUSED_DEVICE_INTERFACE_GENERATOR( LDA, VXC_FXC, + LibxcKernelImpl::eval_vxc_fxc_device_, const ) +UNUSED_DEVICE_INTERFACE_GENERATOR( GGA, FXC, + LibxcKernelImpl::eval_fxc_device_, const ) +UNUSED_DEVICE_INTERFACE_GENERATOR( GGA, VXC_FXC, + LibxcKernelImpl::eval_vxc_fxc_device_, const ) +UNUSED_DEVICE_INTERFACE_GENERATOR( MGGA, FXC, + LibxcKernelImpl::eval_fxc_device_, const ) +UNUSED_DEVICE_INTERFACE_GENERATOR( MGGA, VXC_FXC, + LibxcKernelImpl::eval_vxc_fxc_device_, const ) + +UNUSED_DEVICE_INC_INTERFACE_GENERATOR( LDA, FXC, + LibxcKernelImpl::eval_fxc_inc_device_, const ) +UNUSED_DEVICE_INC_INTERFACE_GENERATOR( GGA, FXC, + LibxcKernelImpl::eval_fxc_inc_device_, const ) +UNUSED_DEVICE_INC_INTERFACE_GENERATOR( MGGA, FXC, + LibxcKernelImpl::eval_fxc_inc_device_, const ) +UNUSED_DEVICE_INC_INTERFACE_GENERATOR( LDA, VXC_FXC, + LibxcKernelImpl::eval_vxc_fxc_inc_device_, const ) +UNUSED_DEVICE_INC_INTERFACE_GENERATOR( GGA, VXC_FXC, + LibxcKernelImpl::eval_vxc_fxc_inc_device_, const ) +UNUSED_DEVICE_INC_INTERFACE_GENERATOR( MGGA, VXC_FXC, + LibxcKernelImpl::eval_vxc_fxc_inc_device_, const ) + } } diff --git a/src/hip/xc_functional_device.hip b/src/hip/xc_functional_device.hip index bb9a81a..df69d56 100644 --- a/src/hip/xc_functional_device.hip +++ b/src/hip/xc_functional_device.hip @@ -46,6 +46,7 @@ #include #include #include +#include __global__ void scal_kernel( const int N, const double fact, const double* X_device, double* Y_device ) { @@ -488,4 +489,29 @@ MGGA_EXC_VXC_GENERATOR_DEVICE( XCFunctional::eval_exc_vxc_device ) const { if( vtau_scr ) hipFree( vtau_scr ); } +// ExchCXX's HIP backend does not implement Fxc/vxc_fxc device evaluation +// (unlike CUDA and SYCL); XCFunctional::eval_fxc_device/eval_vxc_fxc_device +// forward to each constituent kernel's eval_fxc_device/eval_vxc_fxc_device, +// which are themselves stubbed (see hip/builtin.hip, hip/libxc_device.hip), +// so stub these aggregators the same way rather than leave them unresolved +// at link time. +LDA_FXC_GENERATOR_DEVICE( XCFunctional::eval_fxc_device ) const { + throw std::runtime_error("LDA Fxc device evaluation NYI for HIP backend"); +} +LDA_VXC_FXC_GENERATOR_DEVICE( XCFunctional::eval_vxc_fxc_device ) const { + throw std::runtime_error("LDA Fxc device evaluation NYI for HIP backend"); +} +GGA_FXC_GENERATOR_DEVICE( XCFunctional::eval_fxc_device ) const { + throw std::runtime_error("GGA Fxc device evaluation NYI for HIP backend"); +} +GGA_VXC_FXC_GENERATOR_DEVICE( XCFunctional::eval_vxc_fxc_device ) const { + throw std::runtime_error("GGA Fxc device evaluation NYI for HIP backend"); +} +MGGA_FXC_GENERATOR_DEVICE( XCFunctional::eval_fxc_device ) const { + throw std::runtime_error("MGGA Fxc device evaluation NYI for HIP backend"); +} +MGGA_VXC_FXC_GENERATOR_DEVICE( XCFunctional::eval_vxc_fxc_device ) const { + throw std::runtime_error("MGGA Fxc device evaluation NYI for HIP backend"); +} + } From 31ab417d5bd21741958c689ff216e7f08174edf7 Mon Sep 17 00:00:00 2001 From: Ajay Panyala Date: Wed, 26 Aug 2026 13:43:17 -0400 Subject: [PATCH 2/2] remove some comments --- src/hip/builtin.hip | 11 ----------- 1 file changed, 11 deletions(-) diff --git a/src/hip/builtin.hip b/src/hip/builtin.hip index bdaf8ca..756eb09 100644 --- a/src/hip/builtin.hip +++ b/src/hip/builtin.hip @@ -885,17 +885,6 @@ MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ) { } -// ExchCXX's HIP backend does not implement Fxc/vxc_fxc device evaluation -// (unlike CUDA and SYCL). Rather than leaving BuiltinKernelImpl's -// vtable (eval_fxc_device/eval_vxc_fxc_device, declared in kernel_type.hpp) -// unresolved at link time for every functional, provide a generic stub that -// throws at runtime; ordinary exc/vxc evaluation is unaffected. The -// underlying math (kernel_traits::eval_fxc_*) already exists and is -// shared with the host/CUDA/SYCL backends, so a real port -- mirroring -// cuda/builtin.cu's device_eval_*_helper_* kernels -- remains possible as a -// follow-up; that port must account for HIP-specific numerics differences -// already present in this file (e.g. the fmax(rho,0.) clamps above) rather -// than being a verbatim copy. template LDA_FXC_GENERATOR_DEVICE( device_eval_fxc_helper_unpolar ) { throw std::runtime_error("LDA Fxc device evaluation NYI for HIP backend");