================
@@ -0,0 +1,132 @@
+// 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 %t.cir %}
+// RUN: %if cir-enabled %{ FileCheck --check-prefix=CIR-HOST
--input-file=%t.cir %s %}
+// RUN: %if cir-enabled %{ %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir
-fcuda-is-device -emit-cir -target-sdk-version=12.3 %s -o %t.cir %}
+// RUN: %if cir-enabled %{ FileCheck --check-prefix=CIR-DEVICE
--input-file=%t.cir %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 %t.ll %}
+// RUN: %if cir-enabled %{ FileCheck --check-prefix=LLVM-HOST
--input-file=%t.ll %s %}
+// RUN: %if cir-enabled %{ %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir
-fcuda-is-device -emit-llvm -target-sdk-version=12.3 %s -o %t.ll %}
+// RUN: %if cir-enabled %{ FileCheck --check-prefix=LLVM-DEVICE
--input-file=%t.ll %s %}
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x cuda -emit-llvm
-target-sdk-version=12.3 %s -o %t.ll
+// RUN: FileCheck --check-prefix=OGCG-HOST --input-file=%t.ll %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm
-target-sdk-version=12.3 %s -o %t.ll
+// RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %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) {}
+}
+
+// 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) {}
+}
----------------
Men-cotton wrote:
Add a newline at the end of the file.
https://github.com/llvm/llvm-project/pull/184444
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
