[clang] Support '-fgpu-default-stream=per-thread' for NVIDIA CUDA

[boxu-zhang]

Summary

I'm using clang to compile CUDA code. And just found that clang doesn't support the per-thread stream option for NV CUDA. I don't know if there is another solution.

[tra]

Looking at CUDA headers, it appears that changing only compiler-generated-glue may be insufficient. A lot of other CUDA API calls need to be changed to _ptsz variant and for that we need to have CUDA_API_PER_THREAD_DEFAULT_STREAM defined.

[boxu-zhang]

Append 'CUDA_API_PER_THREAD_DEFAULT_STREAM' as a defined macro

[boxu-zhang]

In D154822#4485700, @tra wrote:

Looking at CUDA headers, it appears that changing only compiler-generated-glue may be insufficient. A lot of other CUDA API calls need to be changed to _ptsz variant and for that we need to have CUDA_API_PER_THREAD_DEFAULT_STREAM defined.

CUDA_API_PER_THREAD_DEFAULT_STREAM is defined now.

[boxu-zhang]

Another point that I don't get is why the libcxx CI failed. Anyone know the reason this message?
"
Running global pre-checkout hook
Preparing working directory
Running global post-checkout hook
Running commands
$ trap 'kill -- $$' INT TERM QUIT; libcxx/utils/ci/generate-buildkite-pipeline | buildkite-agent pipeline upload
2023-07-11 02:36:45 INFO Reading pipeline config from STDIN
2023-07-11 02:36:46 INFO Updating BUILDKITE_COMMIT to "a297905cd83911c8a03f060cb9d96bc99aae3f8c"
2023-07-11 02:36:46 FATAL Pipeline parsing of "(stdin)" failed (Expected identifier to start with a letter, got ')
🚨 Error: The command exited with status 1
user command error: exit status 1
"

[boxu-zhang]

Add component 'clang' in commit message

[boxu-zhang]

I don't have the permission to push to main branch. Can anyone push this?

[tra]

Can anyone push this?

I can help with this. How do you want your commit to be attributed? The patch currently has boxu.zhang <[e-mail]>. Do you want it to be changed to something else?

[boxu-zhang]

In D154822#4494287, @tra wrote:

Can anyone push this?

I can help with this. How do you want your commit to be attributed? The patch currently has boxu.zhang <[e-mail]>. Do you want it to be changed to something else?

Just use that, it's my real name and mailbox, thanks for your help. :)

clang/lib/CodeGen/CGCUDANV.cpp

void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
   //                              void **args, size_t sharedMem,
   //                              cudaStream_t stream);
   // hipError_t hipLaunchKernel[_spt](const void *func, dim3 gridDim,
   //                                  dim3 blockDim, void **args,
   //                                  size_t sharedMem, hipStream_t stream);
   TranslationUnitDecl *TUDecl = CGM.getContext().getTranslationUnitDecl();
   DeclContext *DC = TranslationUnitDecl::castToDeclContext(TUDecl);
   std::string KernelLaunchAPI = "LaunchKernel";
-  if (CGF.getLangOpts().HIP && CGF.getLangOpts().GPUDefaultStream ==
-                                   LangOptions::GPUDefaultStreamKind::PerThread)
+  if (CGF.getLangOpts().GPUDefaultStream ==
+      LangOptions::GPUDefaultStreamKind::PerThread) {
+    if (CGF.getLangOpts().HIP)
       KernelLaunchAPI = KernelLaunchAPI + "_spt";
+    else if (CGF.getLangOpts().CUDA)
+      KernelLaunchAPI = KernelLaunchAPI + "_ptsz";
+  }
   auto LaunchKernelName = addPrefixToName(KernelLaunchAPI);
   IdentifierInfo &cudaLaunchKernelII =
       CGM.getContext().Idents.get(LaunchKernelName);
   FunctionDecl *cudaLaunchKernelFD = nullptr;
   for (auto *Result : DC->lookup(&cudaLaunchKernelII)) {
     if (FunctionDecl *FD = dyn_cast<FunctionDecl>(Result))
       cudaLaunchKernelFD = FD;
   }

clang/lib/Frontend/InitPreprocessor.cpp

static void InitializeStandardPredefinedMacros(const TargetInfo &TI,
   // Not "standard" per se, but available even with the -undef flag.
   if (LangOpts.AsmPreprocessor)
     Builder.defineMacro("__ASSEMBLER__");
   if (LangOpts.CUDA) {
     if (LangOpts.GPURelocatableDeviceCode)
       Builder.defineMacro("__CLANG_RDC__");
     if (!LangOpts.HIP)
       Builder.defineMacro("__CUDA__");
+    if (LangOpts.GPUDefaultStream ==
+        LangOptions::GPUDefaultStreamKind::PerThread)
+      Builder.defineMacro("CUDA_API_PER_THREAD_DEFAULT_STREAM");
   }
   if (LangOpts.HIP) {
     Builder.defineMacro("__HIP__");
     Builder.defineMacro("__HIPCC__");
     Builder.defineMacro("__HIP_MEMORY_SCOPE_SINGLETHREAD", "1");
     Builder.defineMacro("__HIP_MEMORY_SCOPE_WAVEFRONT", "2");
     Builder.defineMacro("__HIP_MEMORY_SCOPE_WORKGROUP", "3");
     Builder.defineMacro("__HIP_MEMORY_SCOPE_AGENT", "4");

clang/test/CodeGenCUDA/Inputs/cuda.h

extern "C" int cudaConfigureCall(dim3 gridSize, dim3 blockSize,
                                  size_t sharedSize = 0,
                                  cudaStream_t stream = 0);
 extern "C" int __cudaPushCallConfiguration(dim3 gridSize, dim3 blockSize,
                                            size_t sharedSize = 0,
                                            cudaStream_t stream = 0);
 extern "C" cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim,
                                         dim3 blockDim, void **args,
                                         size_t sharedMem, cudaStream_t stream);
+extern "C" cudaError_t cudaLaunchKernel_ptsz(const void *func, dim3 gridDim,
+                                        dim3 blockDim, void **args,
+                                        size_t sharedMem, cudaStream_t stream);
+
 #endif
 
 extern "C" __device__ int printf(const char*, ...);

clang/test/CodeGenCUDA/kernel-call.cu

 // RUN: %clang_cc1 -target-sdk-version=8.0 -emit-llvm %s -o - \
 // RUN: | FileCheck %s --check-prefixes=CUDA-OLD,CHECK
 // RUN: %clang_cc1 -target-sdk-version=9.2  -emit-llvm %s -o - \
 // RUN: | FileCheck %s --check-prefixes=CUDA-NEW,CHECK
+// RUN: %clang_cc1 -target-sdk-version=9.2  -emit-llvm %s -o - \
+// RUN:   -fgpu-default-stream=per-thread -DCUDA_API_PER_THREAD_DEFAULT_STREAM \
+// RUN: | FileCheck %s --check-prefixes=CUDA-PTH,CHECK
 // RUN: %clang_cc1 -x hip -emit-llvm %s -o - \
 // RUN: | FileCheck %s --check-prefixes=HIP-OLD,CHECK
 // RUN: %clang_cc1 -fhip-new-launch-api -x hip -emit-llvm %s -o - \
 // RUN: | FileCheck %s --check-prefixes=HIP-NEW,LEGACY,CHECK
 // RUN: %clang_cc1 -fhip-new-launch-api -x hip -emit-llvm %s -o - \
 // RUN:   -fgpu-default-stream=legacy \
 // RUN:   | FileCheck %s --check-prefixes=HIP-NEW,LEGACY,CHECK
 // RUN: %clang_cc1 -fhip-new-launch-api -x hip -emit-llvm %s -o - \
 // RUN:   -fgpu-default-stream=per-thread -DHIP_API_PER_THREAD_DEFAULT_STREAM \
 // RUN:   | FileCheck %s --check-prefixes=HIP-NEW,PTH,CHECK
 
 #include "Inputs/cuda.h"
 
 // CHECK-LABEL: define{{.*}}g1
 // HIP-OLD: call{{.*}}hipSetupArgument
 // HIP-OLD: call{{.*}}hipLaunchByPtr
 // HIP-NEW: call{{.*}}__hipPopCallConfiguration
 // LEGACY: call{{.*}}hipLaunchKernel
 // PTH: call{{.*}}hipLaunchKernel_spt
 // CUDA-OLD: call{{.*}}cudaSetupArgument
 // CUDA-OLD: call{{.*}}cudaLaunch
 // CUDA-NEW: call{{.*}}__cudaPopCallConfiguration
 // CUDA-NEW: call{{.*}}cudaLaunchKernel
+// CUDA-PTH: call{{.*}}cudaLaunchKernel_ptsz
 __global__ void g1(int x) {}
 
 // CHECK-LABEL: define{{.*}}main
 int main(void) {
   // HIP-OLD: call{{.*}}hipConfigureCall
   // HIP-NEW: call{{.*}}__hipPushCallConfiguration
   // CUDA-OLD: call{{.*}}cudaConfigureCall
   // CUDA-NEW: call{{.*}}__cudaPushCallConfiguration
   // CHECK: icmp
   // CHECK: br
   // CHECK: call{{.*}}g1
   g1<<<1, 1>>>(42);
 }