[clang] [CUDA] Add device-side kernel launch support (PR #165519)

[via cfe-commits]

================
@@ -0,0 +1,23 @@
+// RUN: %clang_cc1 -fcuda-is-device -verify=nordc %s
+// RUN: %clang_cc1 -fcuda-is-device -fgpu-rdc -verify=rdc %s
+// RUN: %clang_cc1 -x hip -fcuda-is-device -verify=hip %s
+
+// rdc-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+__global__ void g2(int x) {}
+
+// CHECK-LABEL: define{{.*}}g1
+__global__ void g1(void) {
+  // CHECK: [[CONFIG:%.*]] = call{{.*}}_Z22cudaGetParameterBuffermm(i64{{.*}}64, i64{{.*}}4)
+  // CHECK-NEXT: [[FLAG:%.*]] = icmp ne ptr [[CONFIG]], null
+  // CHECK-NEXT: br i1 [[FLAG]], label %[[THEN:.*]], label %[[ENDIF:.*]]
+  // CHECK: [[THEN]]:
+  // CHECK-NEXT: [[PPTR:%.*]] = getelementptr{{.*}}i8, ptr [[CONFIG]], i64 0
+  // CHECK-NEXT: store i32 42, ptr [[PPTR]]
+  // CHECK: = call{{.*}} i32 @_Z16cudaLaunchDevicePvS_4dim3S0_jP10cudaStream(ptr{{.*}} @_Z2g2i, ptr{{.*}} [[CONFIG]],
----------------
darkbuck wrote:

Done. clean up in 28e200495e5b39b7599846c511e61723cde2f478

https://github.com/llvm/llvm-project/pull/165519