diff --git a/new-manual/gpu.md b/new-manual/gpu.md index 1b7e262e51..50e38dd2ef 100644 --- a/new-manual/gpu.md +++ b/new-manual/gpu.md @@ -50,6 +50,11 @@ It is also possible to pass some extra options by exporting or specifying the th - [INSPHERE](INSPHERE.md) - adjmat - [CONTACT_MATRIX](CONTACT_MATRIX.md) +- function + - [LESS_THAN](LESS_THAN.md) + - [MORE_THAN](MORE_THAN.md) + - [BETWEEN](BETWEEN.md) + - [COMBINE](COMBINE.md) - matrixtools - [MATRIX_PRODUCT](MATRIX_PRODUCT.md) - [MATRIX_VECTOR_PRODUCT](MATRIX_VECTOR_PRODUCT.md) diff --git a/regtest/targetGPU/rt-adjmat-eigvecs/plumed.dat b/regtest/targetGPU/rt-adjmat-eigvecs/plumed.dat index 0999d531c1..66c1772156 100644 --- a/regtest/targetGPU/rt-adjmat-eigvecs/plumed.dat +++ b/regtest/targetGPU/rt-adjmat-eigvecs/plumed.dat @@ -7,6 +7,6 @@ eigv4: SELECT_COMPONENTS ARG=diag.vecs-1 COMPONENTS=4 eigv5: SELECT_COMPONENTS ARG=diag.vecs-1 COMPONENTS=5 eigv6: SELECT_COMPONENTS ARG=diag.vecs-1 COMPONENTS=6 eigv7: SELECT_COMPONENTS ARG=diag.vecs-1 COMPONENTS=7 -cc: COMBINE ARG=eigv1,eigv2,eigv3,eigv4,eigv5,eigv6,eigv7 COEFFICIENTS=1,2,3,4,5,6,7 POWERS=2,2,3,4,3,2,1 PERIODIC=NO NORMALIZE +cc: COMBINE ARG=eigv1,eigv2,eigv3,eigv4,eigv5,eigv6,eigv7 COEFFICIENTS=1,2,3,4,5,6,7 POWERS=2,2,3,4,3,2,1 PERIODIC=NO NORMALIZE USEGPU PRINT ARG=cc FILE=colvar FMT=%8.4f BIASVALUE ARG=cc diff --git a/regtest/targetGPU/rt-multicolvar-filters/Makefile b/regtest/targetGPU/rt-multicolvar-filters/Makefile new file mode 100644 index 0000000000..3703b27cea --- /dev/null +++ b/regtest/targetGPU/rt-multicolvar-filters/Makefile @@ -0,0 +1 @@ +include ../../scripts/test.make diff --git a/regtest/targetGPU/rt-multicolvar-filters/colvar.reference b/regtest/targetGPU/rt-multicolvar-filters/colvar.reference new file mode 100644 index 0000000000..e613efa869 --- /dev/null +++ b/regtest/targetGPU/rt-multicolvar-filters/colvar.reference @@ -0,0 +1,6 @@ +#! FIELDS time d1_m d1_ltm d1_mtm d1_bt1m d1_bt2m + 0.000000 2.2713 1.2137 2.4776 2.0375 3.1830 + 0.050000 2.2754 1.2362 2.4781 2.0957 3.1959 + 0.100000 2.2785 1.2487 2.4795 2.1196 3.2319 + 0.150000 2.2799 1.2551 2.4799 2.1201 3.2387 + 0.200000 2.2761 1.2554 2.4753 2.1159 3.2374 diff --git a/regtest/targetGPU/rt-multicolvar-filters/config b/regtest/targetGPU/rt-multicolvar-filters/config new file mode 100644 index 0000000000..4804b7f2a8 --- /dev/null +++ b/regtest/targetGPU/rt-multicolvar-filters/config @@ -0,0 +1,4 @@ +type=driver +# this is to test a different name +arg="--plumed plumed.dat --trajectory-stride 10 --timestep 0.005 --ixyz trajectory.xyz --dump-forces forces --dump-forces-fmt=%8.4f" +extra_files="../../trajectories/trajectory.xyz" diff --git a/regtest/targetGPU/rt-multicolvar-filters/forces.reference b/regtest/targetGPU/rt-multicolvar-filters/forces.reference new file mode 100644 index 0000000000..e19d2dc778 --- /dev/null +++ b/regtest/targetGPU/rt-multicolvar-filters/forces.reference @@ -0,0 +1,550 @@ +108 + 16.6207 20.7691 22.9512 +X 5.5637 0.2820 0.9109 +X -0.3998 0.0052 -0.3526 +X -0.3800 -0.3736 -0.0148 +X -0.0311 -0.4006 -0.3504 +X -0.0126 -0.0185 -0.6260 +X -0.2126 -0.0104 -0.5891 +X -0.2699 -0.2568 -0.5036 +X -0.0057 -0.1935 0.5957 +X 0.0193 -0.0071 0.6261 +X -0.3726 -0.0078 0.3812 +X -0.2664 -0.2597 0.5040 +X -0.0004 -0.3852 0.3686 +X -0.0298 -0.6257 -0.0053 +X -0.2485 -0.5204 -0.2446 +X -0.2096 0.5902 0.0111 +X -0.0163 0.5971 -0.1888 +X -0.0000 -0.4339 -0.4518 +X -0.0921 -0.1905 -0.2917 +X -0.1002 0.2833 -0.1990 +X -0.0113 0.2526 0.2568 +X -0.0403 -0.4422 0.4418 +X -0.2804 -0.5026 0.2473 +X -0.1021 0.2888 0.1898 +X -0.0207 -0.5851 0.2228 +X -0.0427 0.6249 0.0044 +X -0.2622 0.5140 -0.2439 +X -0.3864 0.3664 0.0253 +X -0.0202 0.3540 -0.3981 +X -0.0212 0.4540 -0.4311 +X -0.0934 0.1942 -0.2888 +X -0.2745 0.2472 -0.5059 +X -0.0153 0.2015 -0.5929 +X -0.0035 0.4565 0.4289 +X -0.2645 0.5051 0.2593 +X -0.2539 0.2450 0.5176 +X -0.0077 0.3680 0.3857 +X -0.6258 0.0273 0.0062 +X 0.5863 0.0143 -0.2201 +X 0.5910 -0.2075 -0.0035 +X -0.5281 -0.2368 -0.2397 +X -0.4466 -0.0036 -0.4392 +X -0.2562 0.0027 -0.2535 +X 0.2882 -0.0989 -0.1924 +X -0.2019 -0.0970 0.2823 +X -0.4634 0.0029 0.4215 +X 0.5948 -0.0103 0.1961 +X 0.2912 -0.0983 0.1883 +X -0.5036 -0.2742 0.2521 +X -0.4453 -0.4406 0.0054 +X 0.2843 -0.1965 -0.1021 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +108 + 17.5381 21.4331 23.5417 +X 5.8140 1.8001 1.1707 +X -0.4376 0.0034 -0.3533 +X -0.4043 -0.3898 -0.0302 +X -0.0430 -0.4344 -0.3547 +X -0.0188 -0.0367 -0.6264 +X -0.2185 -0.0199 -0.5882 +X -0.2834 -0.2549 -0.4988 +X -0.0046 -0.1856 0.5997 +X 0.0403 -0.0133 0.6263 +X -0.3877 -0.0050 0.4074 +X -0.2749 -0.2646 0.4985 +X -0.0080 -0.4148 0.3797 +X -0.0491 -0.6258 0.0006 +X -0.2434 -0.5282 -0.2362 +X -0.2188 -0.5880 0.0212 +X -0.0275 0.5998 -0.1833 +X 0.0049 -0.4258 -0.4612 +X -0.1562 -0.3287 -0.5115 +X -0.1021 0.2799 -0.2030 +X -0.0169 0.2533 0.2559 +X -0.0697 -0.4436 0.4387 +X -0.2947 -0.5000 0.2393 +X -0.1086 0.2872 0.1890 +X -0.0341 -0.5795 0.2390 +X -0.0818 0.6224 0.0084 +X -0.2681 0.5176 -0.2331 +X -0.4036 0.3893 0.0436 +X -0.0457 0.3506 -0.4374 +X -0.0395 0.4631 -0.4220 +X -0.0930 0.1943 -0.2891 +X -0.2840 0.2360 -0.5077 +X -0.0265 0.2032 -0.5934 +X -0.0054 0.4692 0.4171 +X -0.2708 0.5025 0.2613 +X -0.2552 0.2420 0.5200 +X -0.0174 0.3820 0.4124 +X -0.6259 0.0455 0.0177 +X 0.5800 0.0308 -0.2383 +X 0.5914 -0.2106 -0.0050 +X -0.5402 -0.2234 -0.2288 +X -0.4497 -0.0077 -0.4380 +X -0.2565 0.0030 -0.2533 +X 0.2881 -0.1017 -0.1914 +X -0.2100 -0.0983 0.2761 +X -0.4771 -0.0008 0.4080 +X 0.5961 -0.0167 0.1961 +X 0.2934 -0.1025 0.1828 +X -0.4958 -0.2882 0.2554 +X -0.4478 -0.4399 0.0092 +X 0.4878 -0.3467 -0.1896 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +108 + 18.4475 21.4087 23.6649 +X 5.8798 3.4868 2.2014 +X -0.4574 0.0036 -0.3526 +X -0.4270 -0.3873 -0.0352 +X -0.0308 -0.4384 -0.3746 +X -0.0259 -0.0560 -0.6244 +X -0.2213 -0.0238 -0.5866 +X -0.2951 -0.2441 -0.4970 +X 0.0024 -0.1764 -0.6021 +X 0.0602 0.0045 0.6245 +X -0.3837 0.0152 0.4313 +X -0.2894 -0.2591 0.4927 +X -0.0432 -0.4308 0.3822 +X -0.0675 -0.6234 0.0216 +X -0.2462 -0.5256 -0.2384 +X -0.2326 -0.5821 0.0269 +X -0.0339 -0.6001 -0.1801 +X 0.0178 -0.4224 -0.4636 +X -0.1645 -0.3182 -0.5151 +X -0.1031 -0.2773 -0.2066 +X -0.0147 0.2588 0.2510 +X -0.0863 -0.4373 0.4415 +X -0.2934 -0.5000 0.2400 +X -0.1193 0.2828 0.1897 +X -0.0386 -0.5827 0.2293 +X -0.1078 0.6179 0.0156 +X -0.2768 0.5181 -0.2206 +X -0.4038 0.4091 0.0560 +X -0.0823 0.3552 -0.4478 +X -0.0607 0.4686 -0.4128 +X -0.1006 0.1931 -0.2877 +X -0.2894 0.2316 -0.5062 +X -0.0306 0.2076 -0.5913 +X -0.0074 0.4726 0.4126 +X -0.2727 0.4993 0.2645 +X -0.2567 0.2509 0.5146 +X -0.0225 0.3910 0.4244 +X -0.6240 0.0538 0.0371 +X 0.5775 0.0452 -0.2410 +X 0.5963 -0.1950 -0.0094 +X -0.5488 -0.2072 -0.2225 +X -0.4481 -0.0124 -0.4390 +X -0.2556 0.0033 -0.2547 +X 0.2879 -0.1048 -0.1907 +X -0.2168 -0.0983 0.2711 +X -0.4836 -0.0052 0.3997 +X 0.5962 -0.0103 0.1952 +X 0.5134 -0.1815 0.3116 +X -0.4956 -0.2870 0.2562 +X -0.4543 -0.4326 0.0098 +X 0.4805 -0.3499 -0.2008 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +108 + 18.8502 21.1693 23.6288 +X 6.7734 3.8768 1.5673 +X -0.4686 -0.0020 -0.3512 +X -0.4490 -0.3751 -0.0254 +X -0.0154 -0.4401 -0.3860 +X -0.0315 -0.0774 -0.6221 +X -0.2247 -0.0289 -0.5853 +X -0.3101 -0.2362 -0.4919 +X 0.0084 -0.1747 -0.6028 +X 0.0804 0.0234 0.6220 +X -0.3669 0.0419 0.4545 +X -0.3116 -0.2388 0.4897 +X -0.0824 -0.4289 0.3900 +X -0.0916 -0.6196 0.0409 +X -0.2546 -0.5194 -0.2436 +X -0.2417 -0.5786 0.0263 +X -0.0389 -0.6000 -0.1799 +X 0.0290 -0.4266 -0.4595 +X -0.1780 -0.3076 -0.5173 +X -0.1048 -0.2759 -0.2078 +X -0.0115 -0.2608 0.2493 +X -0.0863 -0.4333 0.4458 +X -0.2879 -0.4987 0.2496 +X -0.1254 0.2775 0.1938 +X -0.0413 -0.5918 0.2050 +X -0.1145 0.6163 0.0313 +X -0.2956 0.5122 -0.2104 +X -0.3996 0.4225 0.0687 +X -0.1224 0.3651 -0.4412 +X -0.0827 0.4658 -0.4124 +X -0.1078 0.1921 -0.2859 +X -0.2920 0.2293 -0.5060 +X -0.0290 0.2134 -0.5895 +X -0.0158 0.4693 0.4164 +X -0.2758 0.4927 0.2740 +X -0.2603 0.2584 0.5093 +X -0.0268 0.3949 0.4315 +X -0.6227 0.0513 0.0599 +X -0.5810 0.0547 -0.2311 +X 0.6010 -0.1804 -0.0131 +X -0.5569 -0.1910 -0.2176 +X -0.4443 -0.0164 -0.4430 +X 0.2563 0.0028 0.2542 +X 0.2853 -0.1065 -0.1937 +X -0.2224 -0.0967 0.2673 +X -0.4840 -0.0146 0.3993 +X 0.5946 -0.0004 0.2009 +X 0.5143 -0.1871 0.3073 +X -0.5031 -0.2767 0.2534 +X -0.4622 -0.4244 0.0113 +X 0.4785 -0.3521 -0.2025 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +108 + 18.9617 20.8113 23.5628 +X 7.3957 3.7290 1.2140 +X -0.4768 -0.0014 -0.3387 +X -0.4593 -0.3621 -0.0033 +X 0.0002 -0.4559 -0.3664 +X -0.0310 -0.0813 -0.6194 +X -0.2283 -0.0373 -0.5812 +X -0.3207 -0.2416 -0.4796 +X 0.0111 -0.1804 -0.5988 +X 0.0876 0.0414 0.6180 +X -0.3347 0.0613 0.4757 +X -0.3263 -0.2173 0.4874 +X -0.1116 -0.4064 0.4055 +X -0.1180 -0.6123 0.0494 +X -0.2690 -0.5077 -0.2473 +X -0.2466 -0.5743 0.0250 +X -0.0451 -0.5983 -0.1768 +X 0.0370 -0.4319 -0.4509 +X -0.1924 -0.2992 -0.5145 +X -0.1055 -0.2759 -0.2055 +X -0.0072 -0.2545 0.2543 +X -0.0716 -0.4260 0.4524 +X -0.2818 -0.4879 0.2716 +X -0.1232 0.2714 0.2016 +X -0.0354 -0.5944 0.1917 +X -0.0982 0.6159 0.0474 +X -0.3025 0.5109 -0.1968 +X -0.3928 0.4226 0.0960 +X -0.1500 0.3821 -0.4166 +X -0.0993 0.4600 -0.4121 +X -0.1134 0.1929 -0.2818 +X -0.2913 0.2268 -0.5050 +X -0.0268 0.2232 -0.5837 +X -0.0329 0.4585 0.4243 +X -0.2733 0.4817 0.2907 +X -0.2620 0.2592 0.5054 +X -0.0281 0.3923 0.4329 +X -0.6175 0.0438 0.0899 +X -0.5824 0.0640 -0.2191 +X 0.6013 -0.1721 -0.0108 +X -0.5603 -0.1799 -0.2121 +X -0.4446 -0.0114 -0.4399 +X -0.2559 0.0022 0.2530 +X 0.2837 -0.1053 -0.1947 +X -0.2238 -0.0976 0.2643 +X -0.4757 -0.0230 0.4055 +X 0.5867 0.0156 0.2164 +X 0.5072 -0.1884 0.3140 +X -0.5088 -0.2649 0.2495 +X -0.4659 -0.4170 0.0169 +X 0.4798 -0.3492 -0.1976 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 +X 0.0000 0.0000 0.0000 diff --git a/regtest/targetGPU/rt-multicolvar-filters/plumed.dat b/regtest/targetGPU/rt-multicolvar-filters/plumed.dat new file mode 100644 index 0000000000..0377f989d9 --- /dev/null +++ b/regtest/targetGPU/rt-multicolvar-filters/plumed.dat @@ -0,0 +1,99 @@ +d1: DISTANCE ... + ATOMS1=1,2 + ATOMS2=1,3 + ATOMS3=1,4 + ATOMS4=1,5 + ATOMS5=1,6 + ATOMS6=1,7 + ATOMS7=1,8 + ATOMS8=1,9 + ATOMS9=1,10 + ATOMS10=1,11 + ATOMS11=1,12 + ATOMS12=1,13 + ATOMS13=1,14 + ATOMS14=1,15 + ATOMS15=1,16 + ATOMS16=1,17 + ATOMS17=1,18 + ATOMS18=1,19 + ATOMS19=1,20 + ATOMS20=1,21 + ATOMS21=1,22 + ATOMS22=1,23 + ATOMS23=1,24 + ATOMS24=1,25 + ATOMS25=1,26 + ATOMS26=1,27 + ATOMS27=1,28 + ATOMS28=1,29 + ATOMS29=1,30 + ATOMS30=1,31 + ATOMS31=1,32 + ATOMS32=1,33 + ATOMS33=1,34 + ATOMS34=1,35 + ATOMS35=1,36 + ATOMS36=1,37 + ATOMS37=1,38 + ATOMS38=1,39 + ATOMS39=1,40 + ATOMS40=1,41 + ATOMS41=1,42 + ATOMS42=1,43 + ATOMS43=1,44 + ATOMS44=1,45 + ATOMS45=1,46 + ATOMS46=1,47 + ATOMS47=1,48 + ATOMS48=1,49 + ATOMS49=1,50 +... + +d1_m: MEAN ARG=d1 PERIODIC=NO + +d1_lt: LESS_THAN ARG=d1 SWITCH={GAUSSIAN D_0=1.5 R_0=0.00001} USEGPU +d1_lt_prod: MATHEVAL ARG=d1_lt,d1 FUNC=x*y PERIODIC=NO +d1_lt_numer: SUM ARG=d1_lt_prod PERIODIC=NO +d1_lt_denom: SUM ARG=d1_lt PERIODIC=NO +d1_ltm: MATHEVAL ARG=d1_lt_numer,d1_lt_denom FUNC=x/y PERIODIC=NO + +d1_mt: MORE_THAN ARG=d1 SWITCH={GAUSSIAN D_0=1.5 R_0=0.00001} USEGPU +d1_mt_prod: MATHEVAL ARG=d1_mt,d1 FUNC=x*y PERIODIC=NO +d1_mt_numer: SUM ARG=d1_mt_prod PERIODIC=NO +d1_mt_denom: SUM ARG=d1_mt PERIODIC=NO +d1_mtm: MATHEVAL ARG=d1_mt_numer,d1_mt_denom FUNC=x/y PERIODIC=NO + +d1_bt1: BETWEEN ARG=d1 LOWER=0 UPPER=3.0 SMEAR=0.0001 USEGPU +d1_bt1_prod: MATHEVAL ARG=d1_bt1,d1 FUNC=x*y PERIODIC=NO +d1_bt1_numer: SUM ARG=d1_bt1_prod PERIODIC=NO +d1_bt1_denom: SUM ARG=d1_bt1 PERIODIC=NO +d1_bt1m: MATHEVAL ARG=d1_bt1_numer,d1_bt1_denom FUNC=x/y PERIODIC=NO + +d1_bt2: BETWEEN ARG=d1 LOWER=3.0 UPPER=10.0 SMEAR=0.0001 USEGPU +d1_bt2_prod: MATHEVAL ARG=d1_bt2,d1 FUNC=x*y PERIODIC=NO +d1_bt2_numer: SUM ARG=d1_bt2_prod PERIODIC=NO +d1_bt2_denom: SUM ARG=d1_bt2 PERIODIC=NO +d1_bt2m: MATHEVAL ARG=d1_bt2_numer,d1_bt2_denom FUNC=x/y PERIODIC=NO + +PRINT ARG=d1_m,d1_ltm,d1_mtm,d1_bt1m,d1_bt2m FILE=colvar FMT=%8.4f + +d1_lt1: LESS_THAN ARG=d1 SWITCH={GAUSSIAN D_0=1.5 R_0=0.1} USEGPU +d1_lt1_prod: MATHEVAL ARG=d1_lt,d1 FUNC=x*y PERIODIC=NO +d1_lt1_numer: SUM ARG=d1_lt_prod PERIODIC=NO +d1_lt1_denom: SUM ARG=d1_lt PERIODIC=NO +d1_lt1m: MATHEVAL ARG=d1_lt_numer,d1_lt_denom FUNC=x/y PERIODIC=NO + +d1_mt1: MORE_THAN ARG=d1 SWITCH={GAUSSIAN D_0=1.5 R_0=0.1} USEGPU +d1_mt1_prod: MATHEVAL ARG=d1_mt,d1 FUNC=x*y PERIODIC=NO +d1_mt1_numer: SUM ARG=d1_mt_prod PERIODIC=NO +d1_mt1_denom: SUM ARG=d1_mt PERIODIC=NO +d1_mt1m: MATHEVAL ARG=d1_mt_numer,d1_mt_denom FUNC=x/y PERIODIC=NO + +d1_bt11: BETWEEN ARG=d1 LOWER=0 UPPER=3.0 SMEAR=0.1 USEGPU +d1_bt11_prod: MATHEVAL ARG=d1_bt1,d1 FUNC=x*y PERIODIC=NO +d1_bt11_numer: SUM ARG=d1_bt1_prod PERIODIC=NO +d1_bt11_denom: SUM ARG=d1_bt1 PERIODIC=NO +d1_bt11m: MATHEVAL ARG=d1_bt1_numer,d1_bt1_denom FUNC=x/y PERIODIC=NO + +RESTRAINT ARG=d1_lt1m,d1_mt1m,d1_bt11m AT=1.0,1.0,1.0 KAPPA=10,10,10 diff --git a/src/contour/ContourFindingBase.h b/src/contour/ContourFindingBase.h index 46fa6a8820..aea7156ad2 100644 --- a/src/contour/ContourFindingBase.h +++ b/src/contour/ContourFindingBase.h @@ -59,8 +59,14 @@ void ContourFindingBase::findContour( const std::vector& direction, std: inline double ContourFindingBase::getDifferenceFromContour( const std::vector& x, std::vector& der ) const { std::vector vals(1), derivs( x.size() ); - function::FunctionOutput funcout( 1, vals.data(), x.size(), derivs.data() ); - gridtools::EvaluateGridFunction::calc( function, false, View(x.data(),x.size()), funcout ); + auto funcout = function::FunctionOutput::create( 1, + vals.data(), + x.size(), + derivs.data() ); + gridtools::EvaluateGridFunction::calc( function, + false, + View(x.data(),x.size()), + funcout ); for(unsigned i=0; i A; - VectorGeneric<25> B; + static constexpr std::array A{-4.41534164647933937950E-18, + 3.33079451882223809783E-17, + -2.43127984654795469359E-16, + 1.71539128555513303061E-15, + -1.16853328779934516808E-14, + 7.67618549860493561688E-14, + -4.85644678311192946090E-13, + 2.95505266312963983461E-12, + -1.72682629144155570723E-11, + 9.67580903537323691224E-11, + -5.18979560163526290666E-10, + 2.65982372468238665035E-9, + -1.30002500998624804212E-8, + 6.04699502254191894932E-8, + -2.67079385394061173391E-7, + 1.11738753912010371815E-6, + -4.41673835845875056359E-6, + 1.64484480707288970893E-5, + -5.75419501008210370398E-5, + 1.88502885095841655729E-4, + -5.76375574538582365885E-4, + 1.63947561694133579842E-3, + -4.32430999505057594430E-3, + 1.05464603945949983183E-2, + -2.37374148058994688156E-2, + 4.93052842396707084878E-2, + -9.49010970480476444210E-2, + 1.71620901522208775349E-1, + -3.04682672343198398683E-1, + 6.76795274409476084995E-1}; + static constexpr std::arrayB {-7.23318048787475395456E-18, + -4.83050448594418207126E-18, + 4.46562142029675999901E-17, + 3.46122286769746109310E-17, + -2.82762398051658348494E-16, + -3.42548561967721913462E-16, + 1.77256013305652638360E-15, + 3.81168066935262242075E-15, + -9.55484669882830764870E-15, + -4.15056934728722208663E-14, + 1.54008621752140982691E-14, + 3.85277838274214270114E-13, + 7.18012445138366623367E-13, + -1.79417853150680611778E-12, + -1.32158118404477131188E-11, + -3.14991652796324136454E-11, + 1.18891471078464383424E-11, + 4.94060238822496958910E-10, + 3.39623202570838634515E-9, + 2.26666899049817806459E-8, + 2.04891858946906374183E-7, + 2.89137052083475648297E-6, + 6.88975834691682398426E-5, + 3.36911647825569408990E-3, + 8.04490411014108831608E-1}; public: - unsigned order; + unsigned order{0}; static void registerKeywords( Keywords& keys ); explicit Bessel(); static void read( Bessel& func, ActionWithArguments* action, FunctionOptions& options ); - template - static double chbevl(double x, const VectorGeneric& array); - static void calc( const Bessel& func, bool noderiv, const View& args, FunctionOutput& funcout ); + template + static double chbevl(double x, const std::array& array); + static void calc( const Bessel& func, bool noderiv, View args, FunctionOutput& funcout ); }; typedef FunctionShortcut BesselShortcut; @@ -84,63 +137,7 @@ void Bessel::registerKeywords(Keywords& keys) { keys.setValueDescription("scalar/vector","the value of the bessel function"); } -Bessel::Bessel(): - A(-4.41534164647933937950E-18, - 3.33079451882223809783E-17, - -2.43127984654795469359E-16, - 1.71539128555513303061E-15, - -1.16853328779934516808E-14, - 7.67618549860493561688E-14, - -4.85644678311192946090E-13, - 2.95505266312963983461E-12, - -1.72682629144155570723E-11, - 9.67580903537323691224E-11, - -5.18979560163526290666E-10, - 2.65982372468238665035E-9, - -1.30002500998624804212E-8, - 6.04699502254191894932E-8, - -2.67079385394061173391E-7, - 1.11738753912010371815E-6, - -4.41673835845875056359E-6, - 1.64484480707288970893E-5, - -5.75419501008210370398E-5, - 1.88502885095841655729E-4, - -5.76375574538582365885E-4, - 1.63947561694133579842E-3, - -4.32430999505057594430E-3, - 1.05464603945949983183E-2, - -2.37374148058994688156E-2, - 4.93052842396707084878E-2, - -9.49010970480476444210E-2, - 1.71620901522208775349E-1, - -3.04682672343198398683E-1, - 6.76795274409476084995E-1), - B(-7.23318048787475395456E-18, - -4.83050448594418207126E-18, - 4.46562142029675999901E-17, - 3.46122286769746109310E-17, - -2.82762398051658348494E-16, - -3.42548561967721913462E-16, - 1.77256013305652638360E-15, - 3.81168066935262242075E-15, - -9.55484669882830764870E-15, - -4.15056934728722208663E-14, - 1.54008621752140982691E-14, - 3.85277838274214270114E-13, - 7.18012445138366623367E-13, - -1.79417853150680611778E-12, - -1.32158118404477131188E-11, - -3.14991652796324136454E-11, - 1.18891471078464383424E-11, - 4.94060238822496958910E-10, - 3.39623202570838634515E-9, - 2.26666899049817806459E-8, - 2.04891858946906374183E-7, - 2.89137052083475648297E-6, - 6.88975834691682398426E-5, - 3.36911647825569408990E-3, - 8.04490411014108831608E-1), - order(0) { +Bessel::Bessel() { } void Bessel::read( Bessel& func, ActionWithArguments* action, FunctionOptions& options ) { @@ -160,8 +157,8 @@ void Bessel::read( Bessel& func, ActionWithArguments* action, FunctionOptions& o } } -template -double Bessel::chbevl(double x, const VectorGeneric& array) { +template +double Bessel::chbevl(const double x, const std::array& array) { double b0, b1, b2; b0 = array[0]; @@ -177,21 +174,25 @@ double Bessel::chbevl(double x, const VectorGeneric& array) { } -void Bessel::calc( const Bessel& func, bool noderiv, const View& args, FunctionOutput& funcout ) { +void Bessel::calc( const Bessel& func, + bool noderiv, + const View args, + FunctionOutput& funcout ) { +// NDEBUG seems to be ignored by nvcc... plumed_dbg_assert( args.size()==1 ); if( !noderiv ) { plumed_merror("derivatives have not been implemented for this function"); - } else if( func.order==0 ) { - double x = fabs(args[0]); - if (x <= 8.0) { - double y = (x / 2.0) - 2.0; - funcout.values[0] = chbevl(y, func.A) ; - return; - } - funcout.values[0] = chbevl(32.0 / x - 2.0, func.B) / sqrt(x) ; - } else { - plumed_error(); } + //I trust the parser into avoiding geting here with order!=0 + double x = fabs(args[0]); + if (x <= 8.0) { + double y = (x / 2.0) - 2.0; + funcout.values[0] = chbevl(y, Bessel::A) ; + return; + } + funcout.values[0] = chbevl(32.0 / x - 2.0, Bessel::B) / sqrt(x) ; + ///TODO: higher order bessel functions + } } diff --git a/src/function/Between.cpp b/src/function/Between.cpp index e0f97feb2a..1e1162dbf3 100644 --- a/src/function/Between.cpp +++ b/src/function/Between.cpp @@ -19,6 +19,9 @@ You should have received a copy of the GNU Lesser General Public License along with plumed. If not, see . +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ */ +#ifdef __PLUMED_HAS_OPENACC +#define __PLUMED_USE_OPENACC 1 +#endif //__PLUMED_HAS_OPENACC #include "FunctionSetup.h" #include "FunctionShortcut.h" #include "FunctionOfScalar.h" @@ -140,14 +143,25 @@ tells you how many of these distances are between 0.1 and 0.2 nm. */ //+ENDPLUMEDOC -class Between { -public: - bool isPeriodic; - double min=0, max=0; +struct Between { HistogramBead hist{HistogramBead::KernelType::gaussian,0.0,1.0,0.5}; static void registerKeywords( Keywords& keys ); - static void read( Between& func, ActionWithArguments* action, FunctionOptions& options ); - static void calc( const Between& func, bool noderiv, const View& args, FunctionOutput& funcout ); + static void read( Between& func, ActionWithArguments* action, + FunctionOptions& options ); + static void calc( const Between& func, + bool noderiv, + const View args, + FunctionOutput& funcout ); +#ifdef __PLUMED_USE_OPENACC + void toACCDevice() const { +#pragma acc enter data copyin(this[0:1]) + hist.toACCDevice(); + } + void removeFromACCDevice() const { + hist.removeFromACCDevice(); +#pragma acc exit data delete(order, this[0:1]) + } +#endif // __PLUMED_USE_OPENACC }; typedef FunctionShortcut BetweenShortcut; @@ -182,8 +196,8 @@ void Between::read( Between& func, ActionWithArguments* action, FunctionOptions& } std::string str_min, str_max; - func.isPeriodic = action->getPntrToArgument(0)->isPeriodic(); - if( func.isPeriodic ) { + bool isPeriodic = action->getPntrToArgument(0)->isPeriodic(); + if( isPeriodic ) { action->getPntrToArgument(0)->getDomain( str_min, str_max ); } std::string hinput; @@ -202,17 +216,23 @@ void Between::read( Between& func, ActionWithArguments* action, FunctionOptions& } action->log.printf(" %s \n", func.hist.description().c_str() ); - if( !func.isPeriodic ) { + if( !isPeriodic ) { func.hist.isNotPeriodic(); } else { - Tools::convert( str_min, func.min ); - Tools::convert( str_max, func.max ); - func.hist.isPeriodic( func.min, func.max ); + double min; + double max; + Tools::convert( str_min, min ); + Tools::convert( str_max, max ); + func.hist.isPeriodic( min, max ); } } -void Between::calc( const Between& func, bool noderiv, const View& args, FunctionOutput& funcout ) { - plumed_dbg_assert( args.size()==1 ); +void Between::calc( const Between& func, + bool noderiv, + const View args, + FunctionOutput& funcout ) { + // the presence of NDEBUG seems to be ignored by nvcc... + // plumed_dbg_assert( args.size()==1 ); double deriv; funcout.values[0] = func.hist.calculate( args[0], deriv ); if( !noderiv ) { diff --git a/src/function/Combine.cpp b/src/function/Combine.cpp index 7d95bc6864..21830abb5b 100644 --- a/src/function/Combine.cpp +++ b/src/function/Combine.cpp @@ -19,12 +19,16 @@ You should have received a copy of the GNU Lesser General Public License along with plumed. If not, see . +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ */ +#ifdef __PLUMED_HAS_OPENACC +#define __PLUMED_USE_OPENACC 1 +#endif //__PLUMED_HAS_OPENACC #include "FunctionShortcut.h" #include "FunctionOfScalar.h" #include "FunctionOfVector.h" #include "FunctionOfMatrix.h" #include "core/ActionRegister.h" +#include namespace PLMD { namespace function { @@ -151,16 +155,62 @@ the [CUSTOM](CUSTOM.md) action to transform all the components of the input vect //+ENDPLUMEDOC class Combine { + struct component { + double coefficient{0.0}; + double parameter{0.0}; + double power{0.0}; + double max_minus_min{0.0}; + double inv_max_minus_min{0.0}; + bool periodic{false}; + }; + std::vector components{}; public: - std::vector coefficients; - std::vector parameters; - std::vector powers; - std::vector periodic; - std::vector max_minus_min; - std::vector inv_max_minus_min; + std::size_t ncomponents{0}; + component * cmps=nullptr; + void update() { + cmps = components.data(); + ncomponents = components.size(); + } + Combine() = default; + ~Combine() = default; + Combine(const Combine&x): + components(x.components) { + update(); + } + Combine(Combine&&x): + components(std::move(x.components)) { + update(); + } + Combine &operator=(const Combine&x) { + if (this!=&x) { + components=x.components; + update(); + } + return *this; + } + Combine &operator=(Combine&&x) { + if (this!=&x) { + components=std::move(x.components); + update(); + } + return *this; + } static void registerKeywords(Keywords& keys); - static void read( Combine& func, ActionWithArguments* action, FunctionOptions& options ); - static void calc( const Combine& func, bool noderiv, const View& args, FunctionOutput& funcout ); + static void read( Combine& func, + ActionWithArguments* action, + FunctionOptions& options ); + static void calc( const Combine& func, + bool noderiv, + const View& args, + FunctionOutput& funcout ); +#ifdef __PLUMED_USE_OPENACC + void toACCDevice() const { +#pragma acc enter data copyin(this[0:1], ncomponents, cmps[0:ncomponents]) + } + void removeFromACCDevice() const { +#pragma acc exit data delete(cmps[0:ncomponents], ncomponents, this[0:1]) + } +#endif // __PLUMED_USE_OPENACC }; typedef FunctionShortcut CombineShortcut; @@ -178,89 +228,94 @@ void Combine::registerKeywords(Keywords& keys) { keys.add("compulsory","PARAMETERS","0.0","the parameters of the arguments in your function"); keys.add("compulsory","POWERS","1.0","the powers to which you are raising each of the arguments in your function"); keys.addFlag("NORMALIZE",false,"normalize all the coefficients so that in total they are equal to one"); - keys.setValueDescription("scalar/vector/matrix","a linear compbination"); + keys.setValueDescription("scalar/vector/matrix","a linear combination"); } -void Combine::read( Combine& func, ActionWithArguments* action, FunctionOptions& options ) { +void Combine::read( Combine& func, ActionWithArguments* action, + FunctionOptions& options ) { unsigned nargs = action->getNumberOfArguments(); ActionWithVector* av=dynamic_cast(action); if(av && av->getNumberOfMasks()>0) { nargs = nargs - av->getNumberOfMasks(); } - func.coefficients.resize( nargs ); - func.parameters.resize( nargs ); - func.powers.resize( nargs ); - action->parseVector("COEFFICIENTS",func.coefficients); - if(func.coefficients.size()!=nargs) { + std::vector coefficients(nargs); + action->parseVector("COEFFICIENTS",coefficients); + + if(coefficients.size()!=nargs) { action->error("Size of COEFFICIENTS array should be the same as number for arguments"); } - action->parseVector("PARAMETERS",func.parameters); - if(func.parameters.size()!=nargs) { + std::vector parameters(nargs); + action->parseVector("PARAMETERS",parameters); + if(parameters.size()!=nargs) { action->error("Size of PARAMETERS array should be the same as number for arguments"); } - action->parseVector("POWERS",func.powers); - if(func.powers.size()!=nargs) { + std::vector powers(nargs); + action->parseVector("POWERS",powers); + if(powers.size()!=nargs) { action->error("Size of POWERS array should be the same as number for arguments"); } bool normalize; action->parseFlag("NORMALIZE",normalize); if(normalize) { - double n=0.0; - for(unsigned i=0; ilog.printf(" with coefficients:"); - for(unsigned i=0; ilog.printf(" %f",func.coefficients[i]); + func.components.resize( nargs ); + func.update(); + for(unsigned i=0; ilog.printf(" %f",func.components[i].coefficient); } action->log.printf("\n with parameters:"); - for(unsigned i=0; ilog.printf(" %f",func.parameters[i]); + for(unsigned i=0; ilog.printf(" %f",func.components[i].parameter); } action->log.printf("\n and powers:"); - for(unsigned i=0; ilog.printf(" %f",func.powers[i]); + for(unsigned i=0; ilog.printf(" %f",func.components[i].power); } action->log.printf("\n"); // Store periodicity stuff - func.periodic.resize( nargs, false ); - func.max_minus_min.resize( nargs, 0 ); - func.inv_max_minus_min.resize( nargs, 0 ); for(unsigned i=0; igetPntrToArgument(i))->isPeriodic() ) { - func.periodic[i] = true; + func.components[i].periodic = true; std::string min, max; (action->getPntrToArgument(i))->getDomain( min, max ); double dmin, dmax; Tools::convert( min, dmin ); Tools::convert( max, dmax ); - func.max_minus_min[i] = dmax - dmin; - func.inv_max_minus_min[i] = 1 / func.max_minus_min[i]; + func.components[i].max_minus_min = dmax - dmin; + func.components[i].inv_max_minus_min = 1.0 / func.components[i].max_minus_min ; } } + } -void Combine::calc( const Combine& func, bool noderiv, const View& args, FunctionOutput& funcout ) { +void Combine::calc( const Combine& func, + bool noderiv, + const View& args, + FunctionOutput& funcout ) { funcout.values[0]=0.0; - for(unsigned i=0; i::FunctionOfMatrix(const ActionOptions&ao): } } // Setup the values - FunctionData myfunc; + // Get the function data from the parallel task manager, to avoid copies + auto & myfunc = taskmanager.getActionInput(); myfunc.argstart = argstart; myfunc.nscalars = nscalars; FunctionData::setup( myfunc.f, keywords.getOutputComponents(), shape, false, this ); @@ -178,9 +179,8 @@ FunctionOfMatrix::FunctionOfMatrix(const ActionOptions&ao): for(unsigned i=0; isetSymmetric( symmetric ); } - taskmanager.setupParallelTaskManager( getNumberOfFunctionArguments() - argstart, nscalars ); - // Pass the function to the parallel task manager - taskmanager.setActionInput( myfunc ); + taskmanager.setupParallelTaskManager( getNumberOfFunctionArguments() - argstart, + nscalars ); } template @@ -326,8 +326,15 @@ void FunctionOfMatrix::performTask( std::size_t task_index, const FunctionData& actiondata, ParallelActionsInput& input, ParallelActionsOutput& output ) { - FunctionOutput funcout( input.ncomponents, output.values.data(), input.nderivatives_per_scalar, output.derivatives.data() ); - T::calc( actiondata.f, input.noderiv, View( input.inputdata + task_index*input.nderivatives_per_scalar, input.nderivatives_per_scalar ), funcout ); + auto funcout = FunctionOutput::create( input.ncomponents, + output.values.data(), + input.nderivatives_per_scalar, + output.derivatives.data() ); + T::calc( actiondata.f, + input.noderiv, + View( input.inputdata + task_index*input.nderivatives_per_scalar, + input.nderivatives_per_scalar ), + funcout ); } template diff --git a/src/function/FunctionOfScalar.h b/src/function/FunctionOfScalar.h index b8c3a9adcd..b3a94fba65 100644 --- a/src/function/FunctionOfScalar.h +++ b/src/function/FunctionOfScalar.h @@ -103,8 +103,14 @@ void FunctionOfScalar::calculate() { args[i-argstart]=getPntrToArgument(i)->get(); } std::vector vals( getNumberOfComponents() ), deriv( getNumberOfComponents()*args.size() ); - FunctionOutput funcout( getNumberOfComponents(), vals.data(), args.size(), deriv.data() ); - T::calc( myfunc, doNotCalculateDerivatives(), View(args.data(),args.size()), funcout ); + auto funcout = FunctionOutput::create( getNumberOfComponents(), + vals.data(), + args.size(), + deriv.data() ); + T::calc( myfunc, + doNotCalculateDerivatives(), + View(args.data(),args.size()), + funcout ); for(unsigned i=0; iset(vals[i]); } diff --git a/src/function/FunctionOfVector.h b/src/function/FunctionOfVector.h index 616921b85b..703755d1d7 100644 --- a/src/function/FunctionOfVector.h +++ b/src/function/FunctionOfVector.h @@ -152,15 +152,14 @@ FunctionOfVector::FunctionOfVector(const ActionOptions&ao): error("input mask has wrong size"); } - // Setup the function and the values values - FunctionData myfunc; + // Setup the function and the values + // Get the function data from the parallel task manager, to avoid copies + auto & myfunc = taskmanager.getActionInput(); myfunc.argstart = argstart; myfunc.nscalars = nscalars; FunctionData::setup( myfunc.f, keywords.getOutputComponents(), shape, false, this ); // Setup the parallel task manager taskmanager.setupParallelTaskManager( nargs-argstart, nscalars ); - // Pass the function to the parallel task manager - taskmanager.setActionInput( myfunc ); } template @@ -251,8 +250,15 @@ void FunctionOfVector::performTask( std::size_t task_index, const FunctionData& actiondata, ParallelActionsInput& input, ParallelActionsOutput& output ) { - FunctionOutput funcout( input.ncomponents, output.values.data(), input.nderivatives_per_scalar, output.derivatives.data() ); - T::calc( actiondata.f, input.noderiv, View( input.inputdata + task_index*input.nderivatives_per_scalar, input.nderivatives_per_scalar ), funcout ); + auto funcout = FunctionOutput::create( input.ncomponents, + output.values.data(), + input.nderivatives_per_scalar, + output.derivatives.data() ); + T::calc( actiondata.f, + input.noderiv, + View( input.inputdata + task_index*input.nderivatives_per_scalar, + input.nderivatives_per_scalar ), + funcout ); } template diff --git a/src/function/FunctionSetup.h b/src/function/FunctionSetup.h index d478898c50..4f942b9051 100644 --- a/src/function/FunctionSetup.h +++ b/src/function/FunctionSetup.h @@ -29,8 +29,7 @@ namespace PLMD { namespace function { -class FunctionOptions { -public: +struct FunctionOptions { /// Is the derivative zero if the value is zero bool derivativeZeroIfValueIsZero = false; /// Are multiple components registered with a single name as in SphericalHarmonic and Moments @@ -38,20 +37,37 @@ class FunctionOptions { }; template -class FunctionData { -public: +struct FunctionData { /// Set equal to one if we are doing EvaluateGridFunction unsigned argstart = 0; // Number of scalars that appear in input unsigned nscalars = 0; T f; // This is for setting up the functions - static void setup( T& myfunc, const std::vector& components, const std::vector& shape, bool hasderiv, ActionWithValue* action ); + static void setup( T& myfunc, + const std::vector& components, + const std::vector& shape, + bool hasderiv, + ActionWithValue* action ); +#ifdef __PLUMED_USE_OPENACC + void toACCDevice() const { +#pragma acc enter data copyin(this[0:1], argstart, nscalars) + f.toACCDevice(); + } + void removeFromACCDevice() const { + f.removeFromACCDevice(); +#pragma acc exit data delete(nscalars, argstart, this[0:1]) + } +#endif //__PLUMED_USE_OPENACC }; template -void FunctionData::setup( T& myfunc, const std::vector& components, const std::vector& shape, bool hasderiv, ActionWithValue* action ) { - ActionWithArguments* aarg = dynamic_cast( action ); +void FunctionData::setup( T& myfunc, + const std::vector& components, + const std::vector& shape, + bool hasderiv, + ActionWithValue* action ) { + ActionWithArguments* aarg = action->castToActionWithArguments(); plumed_assert( aarg ); FunctionOptions options; T::read( myfunc, aarg, options ); @@ -145,20 +161,20 @@ void FunctionData::setup( T& myfunc, const std::vector& componen } } -class FunctionOutput { -public: - unsigned nvals; - View values; - unsigned nder; - View2D derivs; - FunctionOutput( unsigned nv, double* v, unsigned na, double* d ): - nvals(nv), - values(v,nv), - nder(na), - derivs(d,nv,na) { +struct FunctionOutput { + View values; + View2D derivs; + static FunctionOutput create( unsigned nvals, + double* vals, + unsigned ndev, + double* devs ) { + return FunctionOutput{ + View(vals,nvals), + View2D(devs,nvals,ndev) + }; } }; -} -} +} // namespace function +} // namespace PLMD #endif diff --git a/src/function/FunctionWithSingleArgument.h b/src/function/FunctionWithSingleArgument.h index 08d9102595..cc8d460314 100644 --- a/src/function/FunctionWithSingleArgument.h +++ b/src/function/FunctionWithSingleArgument.h @@ -107,8 +107,14 @@ void FunctionWithSingleArgument::calculate() { } } std::vector vals( getNumberOfComponents() ), deriv( getNumberOfComponents()*args.size() ); - FunctionOutput funcout( getNumberOfComponents(), vals.data(), args.size(), deriv.data() ); - T::calc( f, doNotCalculateDerivatives(), View(args.data(), args.size()), funcout ); + auto funcout = FunctionOutput::create( getNumberOfComponents(), + vals.data(), + args.size(), + deriv.data() ); + T::calc( f, + doNotCalculateDerivatives(), + View(args.data(), args.size()), + funcout ); for(unsigned i=0; iset( vals[i] ); diff --git a/src/function/LessThan.cpp b/src/function/LessThan.cpp index 478098af4c..cac9c5b03a 100644 --- a/src/function/LessThan.cpp +++ b/src/function/LessThan.cpp @@ -19,6 +19,9 @@ You should have received a copy of the GNU Lesser General Public License along with plumed. If not, see . +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ */ +#ifdef __PLUMED_HAS_OPENACC +#define __PLUMED_USE_OPENACC 1 +#endif //__PLUMED_HAS_OPENACC #include "FunctionSetup.h" #include "FunctionShortcut.h" #include "FunctionOfScalar.h" @@ -402,19 +405,29 @@ was described in the previous section is not performed. class LessThan { public: bool squared; - std::string sfinput; +#ifdef __PLUMED_USE_OPENACC + SwitchingFunctionAccelerable switchingFunction; +#else SwitchingFunction switchingFunction; +#endif //__PLUMED_USE_OPENACC static void registerKeywords( Keywords& keys ); - static void read( LessThan& func, ActionWithArguments* action, FunctionOptions& options ); - static void calc( const LessThan& func, bool noderiv, const View& args, FunctionOutput& funcout ); - LessThan& operator=(const LessThan& m ) { - squared = m.squared; - sfinput = m.sfinput; - std::string errors; - switchingFunction.set( sfinput, errors ); - plumed_assert( errors.length()==0 ); - return *this; + static void read( LessThan& func, + ActionWithArguments* action, + FunctionOptions& options ); + static void calc( const LessThan& func, + bool noderiv, + const View& args, + FunctionOutput& funcout ); +#ifdef __PLUMED_USE_OPENACC + void toACCDevice() const { +#pragma acc enter data copyin(this[0:1],squared) + switchingFunction.toACCDevice(); + } + void removeFromACCDevice() const { + switchingFunction.removeFromACCDevice(); +#pragma acc exit data delete(squared,this[0:1]) } +#endif //__PLUMED_USE_OPENACC }; typedef FunctionShortcut LessThanShortcut; @@ -452,9 +465,10 @@ void LessThan::read( LessThan& func, ActionWithArguments* action, FunctionOption std::string errors; - action->parse("SWITCH",func.sfinput); - if(func.sfinput.length()>0) { - func.switchingFunction.set(func.sfinput,errors); + std::string sfinput; + action->parse("SWITCH",sfinput); + if(sfinput.length()>0) { + func.switchingFunction.set(sfinput,errors); if( errors.length()!=0 ) { action->error("problem reading SWITCH keyword : " + errors ); } @@ -471,12 +485,6 @@ void LessThan::read( LessThan& func, ActionWithArguments* action, FunctionOption action->parse("NN",nn); action->parse("MM",mm); func.switchingFunction.set(nn,mm,r0,d0); - std::string str_nn, str_mm, str_d0, str_r0; - Tools::convert( nn, str_nn ); - Tools::convert( mm, str_mm ); - Tools::convert( d0, str_d0 ); - Tools::convert( r0, str_r0 ); - func.sfinput = "RATIONAL R_0=" + str_r0 + " D_0=" + str_d0 + " NN=" + str_nn + " MM=" + str_mm; } action->log<<" using switching function with cutoff "<parseFlag("SQUARED",func.squared); @@ -485,7 +493,10 @@ void LessThan::read( LessThan& func, ActionWithArguments* action, FunctionOption } } -void LessThan::calc( const LessThan& func, bool noderiv, const View& args, FunctionOutput& funcout ) { +void LessThan::calc( const LessThan& func, + bool noderiv, + const View& args, + FunctionOutput& funcout ) { double d; if( func.squared ) { funcout.values[0] = func.switchingFunction.calculateSqr( args[0], d ); diff --git a/src/function/MoreThan.cpp b/src/function/MoreThan.cpp index d7d34fec92..bf07d37f5b 100644 --- a/src/function/MoreThan.cpp +++ b/src/function/MoreThan.cpp @@ -19,6 +19,9 @@ You should have received a copy of the GNU Lesser General Public License along with plumed. If not, see . +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ */ +#ifdef __PLUMED_HAS_OPENACC +#define __PLUMED_USE_OPENACC 1 +#endif //__PLUMED_HAS_OPENACC #include "FunctionSetup.h" #include "FunctionShortcut.h" #include "FunctionOfScalar.h" @@ -107,19 +110,29 @@ tells you how many of these distances are greater than 0.2 nm. class MoreThan { public: bool squared; - std::string sfinput; +#ifdef __PLUMED_USE_OPENACC + SwitchingFunctionAccelerable switchingFunction; +#else SwitchingFunction switchingFunction; +#endif //__PLUMED_USE_OPENACC static void registerKeywords( Keywords& keys ); - static void read( MoreThan& func, ActionWithArguments* action, FunctionOptions& options ); - static void calc( const MoreThan& func, bool noderiv, const View& args, FunctionOutput& funcout ); - MoreThan& operator=(const MoreThan& m ) { - squared = m.squared; - sfinput = m.sfinput; - std::string errors; - switchingFunction.set( sfinput, errors ); - plumed_assert( errors.length()==0 ); - return *this; + static void read( MoreThan& func, + ActionWithArguments* action, + FunctionOptions& options ); + static void calc( const MoreThan& func, + bool noderiv, + const View& args, + FunctionOutput& funcout ); +#ifdef __PLUMED_USE_OPENACC + void toACCDevice() const { +#pragma acc enter data copyin(this[0:1],squared) + switchingFunction.toACCDevice(); + } + void removeFromACCDevice() const { + switchingFunction.removeFromACCDevice(); +#pragma acc exit data delete(squared,this[0:1]) } +#endif //__PLUMED_USE_OPENACC }; typedef FunctionShortcut MoreThanShortcut; @@ -157,9 +170,10 @@ void MoreThan::read( MoreThan& func, ActionWithArguments* action, FunctionOption std::string errors; - action->parse("SWITCH",func.sfinput); - if(func.sfinput.length()>0) { - func.switchingFunction.set(func.sfinput,errors); + std::string sfinput; + action->parse("SWITCH", sfinput); + if(sfinput.length()>0) { + func.switchingFunction.set(sfinput, errors); if( errors.length()!=0 ) { action->error("problem reading SWITCH keyword : " + errors ); } @@ -176,12 +190,6 @@ void MoreThan::read( MoreThan& func, ActionWithArguments* action, FunctionOption action->parse("NN",nn); action->parse("MM",mm); func.switchingFunction.set(nn,mm,r0,d0); - std::string str_nn, str_mm, str_d0, str_r0; - Tools::convert( nn, str_nn ); - Tools::convert( mm, str_mm ); - Tools::convert( d0, str_d0 ); - Tools::convert( r0, str_r0 ); - func.sfinput = "RATIONAL R_0=" + str_r0 + " D_0=" + str_d0 + " NN=" + str_nn + " MM=" + str_mm; } action->log<<" using switching function with cutoff "<parseFlag("SQUARED",func.squared); @@ -190,8 +198,9 @@ void MoreThan::read( MoreThan& func, ActionWithArguments* action, FunctionOption } } -void MoreThan::calc( const MoreThan& func, bool noderiv, const View& args, FunctionOutput& funcout ) { - plumed_dbg_assert( args.size()==1 ); +void MoreThan::calc( const MoreThan& func, bool noderiv, const View& args, FunctionOutput& funcout ) { + // the presence of NDEBUG seems to be ignored by nvcc... + // plumed_dbg_assert( args.size()==1 ); double d; if( func.squared ) { funcout.values[0] = 1.0 - func.switchingFunction.calculateSqr( args[0], d ); diff --git a/src/gridtools/FunctionOfGrid.h b/src/gridtools/FunctionOfGrid.h index 1a6b82c46a..c030b96b11 100644 --- a/src/gridtools/FunctionOfGrid.h +++ b/src/gridtools/FunctionOfGrid.h @@ -178,8 +178,14 @@ void FunctionOfGrid::performTask( const unsigned& current, MultiValue& myvals } // Calculate the function and its derivatives std::vector vals(getNumberOfComponents()), deriv( getNumberOfComponents()*args.size() ); - function::FunctionOutput funcout( getNumberOfComponents(), vals.data(), args.size(), deriv.data() ); - T::calc( myfunc, false, View(args.data(), args.size()), funcout ); + auto funcout = function::FunctionOutput::create( getNumberOfComponents(), + vals.data(), + args.size(), + deriv.data() ); + T::calc( myfunc, + false, + View(args.data(), args.size()), + funcout ); unsigned np = myvals.getTaskIndex(); // And set the values and derivatives myvals.addValue( 0, vals[0] ); diff --git a/src/gridtools/InterpolateGrid.cpp b/src/gridtools/InterpolateGrid.cpp index 6e8164e4ec..d4db835d82 100644 --- a/src/gridtools/InterpolateGrid.cpp +++ b/src/gridtools/InterpolateGrid.cpp @@ -193,8 +193,14 @@ void InterpolateGrid::performTask( const unsigned& current, MultiValue& myvals ) std::vector pos( output_grid.getDimension() ); output_grid.getGridPointCoordinates( current, pos ); std::vector val(1), der( output_grid.getDimension() ); - function::FunctionOutput funcout( 1, val.data(), output_grid.getDimension(), der.data() ); - EvaluateGridFunction::calc( input_grid, false, View(pos.data(),pos.size()), funcout ); + auto funcout = function::FunctionOutput::create( 1, + val.data(), + output_grid.getDimension(), + der.data() ); + EvaluateGridFunction::calc( input_grid, + false, + View(pos.data(),pos.size()), + funcout ); myvals.setValue( 0, val[0] ); for(unsigned i=0; i