diff --git a/src/hip/builtin.hip b/src/hip/builtin.hip index a1a8f68..756eb09 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,105 @@ MGGA_EXC_VXC_INC_GENERATOR_DEVICE( device_eval_exc_vxc_inc_helper_polar ) { } +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 +992,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 +1010,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 +1028,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 +1053,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 +1090,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 +1146,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"); +} + }