[PATCH] D86376: [HIP] Improve kernel launching latency

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added reviewers: tra, rjmccall.
yaxunl requested review of this revision.

Currently clang emits emits the following code for triple chevron kernel call for HIP:

  __hipPushCallConfiguration(grids, blocks, shmem, stream);
  kernel_stub();

whereas for each kernel, clang emits a kernel_stub:

  void kernel_stub() {
    __hipPopCallConfiguration(&grids, &blocks, &shmem, &stream);
    hipLaunchKernel(kernel_stub, grids, blocks, kernel_args, shmem, stream);
  }

This is really unnecessary. in host code, a kernel function is not really a "function"
since you cannot "call" it in the generated IR, you can only launch it through kernel
launching API.

This patch simplifies the generated code for kernel launching by eliminating the
call of `__hipPushCallConfiguration` and `__hipPopCallConfiguration`. For each
triple chevron, a call of `hipLaunchKernel` is directly emitted. The kernel stub
function is still emitted as an empty function, for the sole purpose of as a shadow
symbol to map to the device symbol in device binary so that runtime can use it
to find the device symbol.

This patch does not change AST for kernel since semantically a triple chevron
is like a function call. Keep it as a function call facilitates overloading resolution
and function argument type checking.

This patch only changes kernel launching codegen for HIP for the new kernel launching
API since we are sure there is no other side effect in `__hipPushCallConfiguration`
and `__hipPopCallConfiguration`.


https://reviews.llvm.org/D86376

Files:
  clang/lib/CodeGen/CGCUDANV.cpp
  clang/test/CodeGenCUDA/kernel-call.cu
  clang/test/CodeGenCUDA/kernel-call.hip
  clang/test/lit.cfg.py

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D86376.287114.patch
Type: text/x-patch
Size: 15828 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200821/a61683dc/attachment-0001.bin>