[async] More fixes of constant folding in async mode

[k-ye]

This is a follow-up of #2034. Besides the Kernel constructor, running kernel in multiple threads could cause problems on device synchronization (at least on Metal, where metal::KernelManager isn't thread safe).

jit_evaluator_cache_mut also guards the fetch_result call. This is because we have to make sure the content of the global Program's result buffer is not written by two threads, too.

Related issue = #

[Click here for the format server]

---

[codecov]

Codecov Report

Merging #2035 (5e49f42) into master (980886e) will not change coverage.
The diff coverage is n/a.

@@           Coverage Diff           @@
##           master    #2035   +/-   ##
=======================================
  Coverage   43.54%   43.54%           
=======================================
  Files          45       45           
  Lines        6267     6267           
  Branches     1110     1110           
=======================================
  Hits         2729     2729           
  Misses       3366     3366           
  Partials      172      172           

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 980886e...5e49f42. Read the comment docs.

[xumingkuan]

LGTM!

[archibate]

LGTM! Btw, why we have two (*ker)(launch_ctx);? According to the DRY principle, shall we move them (including the mutex guard) into a separate function (like jit_launch_and_fetch_result)?