Author: Srivarshitha M Date: 2026-03-19T14:53:12-07:00 New Revision: 61b9fc1d4225af14ac660cd81bd927655e35cef9
URL: https://github.com/llvm/llvm-project/commit/61b9fc1d4225af14ac660cd81bd927655e35cef9 DIFF: https://github.com/llvm/llvm-project/commit/61b9fc1d4225af14ac660cd81bd927655e35cef9.diff LOG: [CIR] Upstream CUDA mangling test with LLVM and OGCG verification (#184444) This PR upstreams the `mangling.cu` test from the ClangIR incubator. Building on the feedback from my previous upstreaming PR, I have expanded the verification for this file to include: 1. **CIR checks:** Verifying the ClangIR generated functions. 2. **LLVM checks:** Verifying the LLVM IR generated via the ClangIR pipeline. 3. **OGCG checks:** Verifying against the original CodeGen pipeline to ensure name-mangling parity. I have also moved the `Inputs/cuda.h` mock header to the upstream `Inputs` directory to support this and future CUDA tests. If this multi-stage verification approach looks correct to the maintainers, I plan to follow up by upstreaming the other currently passing CUDA test, `simple-nvptx-triple.cu`, using the same standard. Verified locally with `llvm-lit`. Partially addresses #156747. Added: clang/test/CodeGenCUDA/mangling.cu Modified: Removed: ################################################################################ diff --git a/clang/test/CodeGenCUDA/mangling.cu b/clang/test/CodeGenCUDA/mangling.cu new file mode 100644 index 0000000000000..dffd676fe8bfb --- /dev/null +++ b/clang/test/CodeGenCUDA/mangling.cu @@ -0,0 +1,126 @@ +// REQUIRES: nvptx-registered-target + +// RUN: %if cir-enabled %{ %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -x cuda -emit-cir -target-sdk-version=12.3 %s -o - | FileCheck --check-prefix=CIR-HOST %s %} +// RUN: %if cir-enabled %{ %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir -fcuda-is-device -emit-cir -target-sdk-version=12.3 %s -o - | FileCheck --check-prefix=CIR-DEVICE %s %} + +// RUN: %if cir-enabled %{ %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -x cuda -emit-llvm -target-sdk-version=12.3 %s -o - | FileCheck --check-prefix=LLVM-HOST %s %} +// RUN: %if cir-enabled %{ %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir -fcuda-is-device -emit-llvm -target-sdk-version=12.3 %s -o - | FileCheck --check-prefix=LLVM-DEVICE %s %} + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x cuda -emit-llvm -target-sdk-version=12.3 %s -o - | FileCheck --check-prefix=OGCG-HOST %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -target-sdk-version=12.3 %s -o - | FileCheck --check-prefix=OGCG-DEVICE %s + +#include "Inputs/cuda.h" + +namespace ns { + +// CIR-HOST: cir.func {{.*}} @_ZN2ns36__device_stub__cpp_global_function_1EiPif +// CIR-DEVICE: cir.func {{.*}} @_ZN2ns21cpp_global_function_1EiPif +// LLVM-HOST: define {{.*}} @_ZN2ns36__device_stub__cpp_global_function_1EiPif +// LLVM-DEVICE: define {{.*}} @_ZN2ns21cpp_global_function_1EiPif +// OGCG-HOST: define {{.*}} @_ZN2ns36__device_stub__cpp_global_function_1EiPif +// OGCG-DEVICE: define {{.*}} @_ZN2ns21cpp_global_function_1EiPif + __global__ void cpp_global_function_1(int a, int* b, float c) {} + +// CIR-HOST: cir.func {{.*}} @_ZN2ns36__device_stub__cpp_global_function_2EiPif +// CIR-DEVICE: cir.func {{.*}} @_ZN2ns21cpp_global_function_2EiPif +// LLVM-HOST: define {{.*}} @_ZN2ns36__device_stub__cpp_global_function_2EiPif +// LLVM-DEVICE: define {{.*}} @_ZN2ns21cpp_global_function_2EiPif +// OGCG-HOST: define {{.*}} @_ZN2ns36__device_stub__cpp_global_function_2EiPif +// OGCG-DEVICE: define {{.*}} @_ZN2ns21cpp_global_function_2EiPif + __global__ void cpp_global_function_2(int a, int* b, float c) {} + +// CIR-HOST: cir.func {{.*}} @_ZN2ns19cpp_host_function_1EiPif +// LLVM-HOST: define {{.*}} @_ZN2ns19cpp_host_function_1EiPif +// OGCG-HOST: define {{.*}} @_ZN2ns19cpp_host_function_1EiPif + __host__ void cpp_host_function_1(int a, int* b, float c) {} + +// CIR-HOST: cir.func {{.*}} @_ZN2ns19cpp_host_function_2EiPif +// LLVM-HOST: define {{.*}} @_ZN2ns19cpp_host_function_2EiPif +// OGCG-HOST: define {{.*}} @_ZN2ns19cpp_host_function_2EiPif + __host__ void cpp_host_function_2(int a, int* b, float c) {} + +// CIR-DEVICE: cir.func {{.*}} @_ZN2ns21cpp_device_function_1EiPif +// LLVM-DEVICE: define {{.*}} @_ZN2ns21cpp_device_function_1EiPif +// OGCG-DEVICE: define {{.*}} @_ZN2ns21cpp_device_function_1EiPif + __device__ void cpp_device_function_1(int a, int* b, float c) {} + +// CIR-DEVICE: cir.func {{.*}} @_ZN2ns21cpp_device_function_2EiPif +// LLVM-DEVICE: define {{.*}} @_ZN2ns21cpp_device_function_2EiPif +// OGCG-DEVICE: define {{.*}} @_ZN2ns21cpp_device_function_2EiPif + __device__ void cpp_device_function_2(int a, int* b, float c) {} +} // namespace ns + +// CIR-HOST: cir.func {{.*}} @_Z36__device_stub__cpp_global_function_1iPif +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_global_function_1iPif +// LLVM-HOST: define {{.*}} @_Z36__device_stub__cpp_global_function_1iPif +// LLVM-DEVICE: define {{.*}} @_Z21cpp_global_function_1iPif +// OGCG-HOST: define {{.*}} @_Z36__device_stub__cpp_global_function_1iPif +// OGCG-DEVICE: define {{.*}} @_Z21cpp_global_function_1iPif +__global__ void cpp_global_function_1(int a, int* b, float c) {} + +// CIR-HOST: cir.func {{.*}} @_Z36__device_stub__cpp_global_function_2iPif +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_global_function_2iPif +// LLVM-HOST: define {{.*}} @_Z36__device_stub__cpp_global_function_2iPif +// LLVM-DEVICE: define {{.*}} @_Z21cpp_global_function_2iPif +// OGCG-HOST: define {{.*}} @_Z36__device_stub__cpp_global_function_2iPif +// OGCG-DEVICE: define {{.*}} @_Z21cpp_global_function_2iPif +__global__ void cpp_global_function_2(int a, int* b, float c) {} + +// CIR-HOST: cir.func {{.*}} @_Z19cpp_host_function_1iPif +// LLVM-HOST: define {{.*}} @_Z19cpp_host_function_1iPif +// OGCG-HOST: define {{.*}} @_Z19cpp_host_function_1iPif +__host__ void cpp_host_function_1(int a, int* b, float c) {} + +// CIR-HOST: cir.func {{.*}} @_Z19cpp_host_function_2iPif +// LLVM-HOST: define {{.*}} @_Z19cpp_host_function_2iPif +// OGCG-HOST: define {{.*}} @_Z19cpp_host_function_2iPif +__host__ void cpp_host_function_2(int a, int* b, float c) {} + +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_device_function_1iPif +// LLVM-DEVICE: define {{.*}} @_Z21cpp_device_function_1iPif +// OGCG-DEVICE: define {{.*}} @_Z21cpp_device_function_1iPif +__device__ void cpp_device_function_1(int a, int* b, float c) {} + +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_device_function_2iPif +// LLVM-DEVICE: define {{.*}} @_Z21cpp_device_function_2iPif +// OGCG-DEVICE: define {{.*}} @_Z21cpp_device_function_2iPif +__device__ void cpp_device_function_2(int a, int* b, float c) {} + +extern "C" { + +// CIR-HOST: cir.func {{.*}} @__device_stub__c_global_function_1 +// CIR-DEVICE: cir.func {{.*}} @c_global_function_1 +// LLVM-HOST: define {{.*}} @__device_stub__c_global_function_1 +// LLVM-DEVICE: define {{.*}} @c_global_function_1 +// OGCG-HOST: define {{.*}} @__device_stub__c_global_function_1 +// OGCG-DEVICE: define {{.*}} @c_global_function_1 + __global__ void c_global_function_1(int a, int* b, float c) {} + +// CIR-HOST: cir.func {{.*}} @__device_stub__c_global_function_2 +// CIR-DEVICE: cir.func {{.*}} @c_global_function_2 +// LLVM-HOST: define {{.*}} @__device_stub__c_global_function_2 +// LLVM-DEVICE: define {{.*}} @c_global_function_2 +// OGCG-HOST: define {{.*}} @__device_stub__c_global_function_2 +// OGCG-DEVICE: define {{.*}} @c_global_function_2 + __global__ void c_global_function_2(int a, int* b, float c) {} + +// CIR-HOST: cir.func {{.*}} @c_host_function_1 +// LLVM-HOST: define {{.*}} @c_host_function_1 +// OGCG-HOST: define {{.*}} @c_host_function_1 + __host__ void c_host_function_1(int a, int* b, float c) {} + +// CIR-HOST: cir.func {{.*}} @c_host_function_2 +// LLVM-HOST: define {{.*}} @c_host_function_2 +// OGCG-HOST: define {{.*}} @c_host_function_2 + __host__ void c_host_function_2(int a, int* b, float c) {} + +// CIR-DEVICE: cir.func {{.*}} @c_device_function_1 +// LLVM-DEVICE: define {{.*}} @c_device_function_1 +// OGCG-DEVICE: define {{.*}} @c_device_function_1 + __device__ void c_device_function_1(int a, int* b, float c) {} + +// CIR-DEVICE: cir.func {{.*}} @c_device_function_2 +// LLVM-DEVICE: define {{.*}} @c_device_function_2 +// OGCG-DEVICE: define {{.*}} @c_device_function_2 + __device__ void c_device_function_2(int a, int* b, float c) {} +} // extern "C" _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
