[llvm-branch-commits] [clang] [CIR][CUDA][NFC] Add CIR-to-LLVM lowering checks for existing registration support (PR #195002)

[via llvm-branch-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang

Author: David Rivera (RiverDave)

<details>
<summary>Changes</summary>

This was originally meant to be added by https://github.com/llvm/llvm-project/pull/187636. Lowering was/will be fixed in https://github.com/llvm/llvm-project/pull/194988 - so this merely adds coverage to the existing CUDA runtime calls.

---
Full diff: https://github.com/llvm/llvm-project/pull/195002.diff


1 Files Affected:

- (modified) clang/test/CIR/CodeGenCUDA/device-stub.cu (+23) 


``````````diff
diff --git a/clang/test/CIR/CodeGenCUDA/device-stub.cu b/clang/test/CIR/CodeGenCUDA/device-stub.cu
index 0f9d4d68d67ff..4b3b920a63b84 100644
--- a/clang/test/CIR/CodeGenCUDA/device-stub.cu
+++ b/clang/test/CIR/CodeGenCUDA/device-stub.cu
@@ -7,6 +7,10 @@
 // RUN:   -target-sdk-version=12.3 -fcuda-include-gpubinary %t -o %t.cir
 // RUN: FileCheck --input-file=%t.cir %s --check-prefix=CIR
 
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -fclangir -emit-llvm %s -x cuda \
+// RUN:   -target-sdk-version=12.3 -fcuda-include-gpubinary %t -o %t-cir.ll
+// RUN: FileCheck --input-file=%t-cir.ll %s --check-prefix=LLVM
+
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -x cuda \
 // RUN:   -target-sdk-version=12.3 -fcuda-include-gpubinary %t -o %t.ll
 // RUN: FileCheck --input-file=%t.ll %s --check-prefix=OGCG
@@ -105,6 +109,25 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
 // OGCG: load ptr, ptr @__cuda_gpubin_handle
 // OGCG: call void @__cudaUnregisterFatBinary
 
+// LLVM: constant [25 x i8] c"GPU binary would be here.", section ".nv_fatbin", align 8
+// LLVM: @__cuda_fatbin_wrapper = {{.*}}constant { i32, i32, ptr, ptr } { i32 1180844977, i32 1, ptr @{{.*}}, ptr null }, section ".nvFatBinSegment"
+// LLVM: @__cuda_gpubin_handle = internal global ptr null
+// LLVM: @llvm.global_ctors = appending global {{.*}}@__cuda_module_ctor
+
+// LLVM: define internal void @__cuda_module_dtor
+// LLVM: load ptr, ptr @__cuda_gpubin_handle
+// LLVM: call void @__cudaUnregisterFatBinary
+
+// LLVM: define internal void @__cuda_register_globals
+// LLVM: call{{.*}}@__cudaRegisterFunction(ptr %{{.*}}, ptr @{{.*}}kernelfunc{{.*}}, ptr @{{.*}}, ptr @{{.*}}, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
+// LLVM: ret void
+
+// LLVM: define internal void @__cuda_module_ctor
+// LLVM: call{{.*}}@__cudaRegisterFatBinary(ptr @__cuda_fatbin_wrapper)
+// LLVM: store ptr %{{.*}}, ptr @__cuda_gpubin_handle
+// LLVM-NEXT: call void @__cuda_register_globals
+// LLVM: call i32 @atexit(ptr @__cuda_module_dtor)
+
 // No GPU binary — no registration infrastructure at all.
 // NOGPUBIN-NOT: fatbin
 // NOGPUBIN-NOT: gpubin

``````````

</details>


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