https://github.com/16srivarshitha updated https://github.com/llvm/llvm-project/pull/184444
>From 224da4ccb6c136719d1b3fc39896e8b53b33d03e Mon Sep 17 00:00:00 2001 From: 16srivarshitha <[email protected]> Date: Wed, 4 Mar 2026 03:37:44 +0530 Subject: [PATCH 1/2] [CIR] Upstream CUDA mangling test with LLVM/OGCG checks --- clang/test/CIR/CodeGen/CUDA/mangling.cu | 81 +++++++++++++++++++++++++ 1 file changed, 81 insertions(+) create mode 100644 clang/test/CIR/CodeGen/CUDA/mangling.cu diff --git a/clang/test/CIR/CodeGen/CUDA/mangling.cu b/clang/test/CIR/CodeGen/CUDA/mangling.cu new file mode 100644 index 0000000000000..bad62892cf318 --- /dev/null +++ b/clang/test/CIR/CodeGen/CUDA/mangling.cu @@ -0,0 +1,81 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -x cuda -emit-cir -target-sdk-version=12.3 %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir -fcuda-is-device -emit-cir -target-sdk-version=12.3 %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -x cuda -emit-llvm -target-sdk-version=12.3 %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM-HOST --input-file=%t.ll %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir -fcuda-is-device -emit-llvm -target-sdk-version=12.3 %s -o %t.ll +// RUN: 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 { + __global__ void cpp_global_function_1(int a, int* b, float c) {} + __global__ void cpp_global_function_2(int a, int* b, float c) {} + __host__ void cpp_host_function_1(int a, int* b, float c) {} + __host__ void cpp_host_function_2(int a, int* b, float c) {} + __device__ void cpp_device_function_1(int a, int* b, float c) {} + __device__ void cpp_device_function_2(int a, int* b, float c) {} +} + +__global__ void cpp_global_function_1(int a, int* b, float c) {} +__global__ void cpp_global_function_2(int a, int* b, float c) {} +__host__ void cpp_host_function_1(int a, int* b, float c) {} +__host__ void cpp_host_function_2(int a, int* b, float c) {} +__device__ void cpp_device_function_1(int a, int* b, float c) {} +__device__ void cpp_device_function_2(int a, int* b, float c) {} + +extern "C" { + __global__ void c_global_function_1(int a, int* b, float c) {} + __global__ void c_global_function_2(int a, int* b, float c) {} + __host__ void c_host_function_1(int a, int* b, float c) {} + __host__ void c_host_function_2(int a, int* b, float c) {} + __device__ void c_device_function_1(int a, int* b, float c) {} + __device__ void c_device_function_2(int a, int* b, float c) {} +} + +// 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 + +// CIR-HOST: cir.func {{.*}} @_ZN2ns36__device_stub__cpp_global_function_2EiPif +// CIR-DEVICE: cir.func {{.*}} @_ZN2ns21cpp_global_function_2EiPif + +// CIR-HOST: cir.func {{.*}} @_ZN2ns19cpp_host_function_1EiPif +// CIR-HOST: cir.func {{.*}} @_ZN2ns19cpp_host_function_2EiPif + +// CIR-DEVICE: cir.func {{.*}} @_ZN2ns21cpp_device_function_1EiPif +// CIR-DEVICE: cir.func {{.*}} @_ZN2ns21cpp_device_function_2EiPif + +// CIR-HOST: cir.func {{.*}} @_Z36__device_stub__cpp_global_function_1iPif +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_global_function_1iPif + +// CIR-HOST: cir.func {{.*}} @_Z36__device_stub__cpp_global_function_2iPif +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_global_function_2iPif + +// CIR-HOST: cir.func {{.*}} @_Z19cpp_host_function_1iPif +// CIR-HOST: cir.func {{.*}} @_Z19cpp_host_function_2iPif + +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_device_function_1iPif +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_device_function_2iPif + +// CIR-HOST: cir.func {{.*}} @__device_stub__c_global_function_1 +// CIR-DEVICE: cir.func {{.*}} @c_global_function_1 + +// CIR-HOST: cir.func {{.*}} @__device_stub__c_global_function_2 +// CIR-DEVICE: cir.func {{.*}} @c_global_function_2 + +// CIR-HOST: cir.func {{.*}} @c_host_function_1 +// CIR-HOST: cir.func {{.*}} @c_host_function_2 + +// CIR-DEVICE: cir.func {{.*}} @c_device_function_1 +// CIR-DEVICE: cir.func {{.*}} @c_device_function_2 >From 52193d543649fba85adcf90c88acf8000555edbe Mon Sep 17 00:00:00 2001 From: 16srivarshitha <[email protected]> Date: Thu, 5 Mar 2026 00:31:25 +0530 Subject: [PATCH 2/2] [CIR] Move CUDA test to standard CodeGenCUDA directory --- clang/test/CodeGenCUDA/Inputs/cuda.h | 183 +-------------------------- clang/test/CodeGenCUDA/mangling.cu | 81 ++++++++++++ 2 files changed, 83 insertions(+), 181 deletions(-) create mode 100644 clang/test/CodeGenCUDA/mangling.cu diff --git a/clang/test/CodeGenCUDA/Inputs/cuda.h b/clang/test/CodeGenCUDA/Inputs/cuda.h index 421fa4dd7dbae..204bf2972088d 100644 --- a/clang/test/CodeGenCUDA/Inputs/cuda.h +++ b/clang/test/CodeGenCUDA/Inputs/cuda.h @@ -1,5 +1,5 @@ /* Minimal declarations for CUDA support. Testing purposes only. */ - +/* From test/CodeGenCUDA/Inputs/cuda.h. */ #include <stddef.h> #if __HIP__ || __CUDA__ @@ -13,8 +13,6 @@ #endif #define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__))) #define __grid_constant__ __attribute__((grid_constant)) -#define __cluster_dims__(...) __attribute__((cluster_dims(__VA_ARGS__))) -#define __no_cluster__ __attribute__((no_cluster)) #else #define __constant__ #define __device__ @@ -24,8 +22,6 @@ #define __managed__ #define __launch_bounds__(...) #define __grid_constant__ -#define __cluster_dims__(...) -#define __no_cluster__ #endif struct dim3 { @@ -72,182 +68,7 @@ extern "C" cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim, extern "C" cudaError_t cudaLaunchKernel_ptsz(const void *func, dim3 gridDim, dim3 blockDim, void **args, size_t sharedMem, cudaStream_t stream); -extern "C" __device__ cudaError_t cudaLaunchDevice(void *func, - void *parameterBuffer, - dim3 gridDim, dim3 blockDim, - unsigned int sharedMem, - cudaStream_t stream); -extern "C" __device__ void *cudaGetParameterBuffer(size_t alignment, - size_t size); + #endif extern "C" __device__ int printf(const char*, ...); - -struct char1 { - char x; - __host__ __device__ char1(char x = 0) : x(x) {} -}; -struct char2 { - char x, y; - __host__ __device__ char2(char x = 0, char y = 0) : x(x), y(y) {} -}; -struct char4 { - char x, y, z, w; - __host__ __device__ char4(char x = 0, char y = 0, char z = 0, char w = 0) : x(x), y(y), z(z), w(w) {} -}; - -struct uchar1 { - unsigned char x; - __host__ __device__ uchar1(unsigned char x = 0) : x(x) {} -}; -struct uchar2 { - unsigned char x, y; - __host__ __device__ uchar2(unsigned char x = 0, unsigned char y = 0) : x(x), y(y) {} -}; -struct uchar4 { - unsigned char x, y, z, w; - __host__ __device__ uchar4(unsigned char x = 0, unsigned char y = 0, unsigned char z = 0, unsigned char w = 0) : x(x), y(y), z(z), w(w) {} -}; - -struct short1 { - short x; - __host__ __device__ short1(short x = 0) : x(x) {} -}; -struct short2 { - short x, y; - __host__ __device__ short2(short x = 0, short y = 0) : x(x), y(y) {} -}; -struct short4 { - short x, y, z, w; - __host__ __device__ short4(short x = 0, short y = 0, short z = 0, short w = 0) : x(x), y(y), z(z), w(w) {} -}; - -struct ushort1 { - unsigned short x; - __host__ __device__ ushort1(unsigned short x = 0) : x(x) {} -}; -struct ushort2 { - unsigned short x, y; - __host__ __device__ ushort2(unsigned short x = 0, unsigned short y = 0) : x(x), y(y) {} -}; -struct ushort4 { - unsigned short x, y, z, w; - __host__ __device__ ushort4(unsigned short x = 0, unsigned short y = 0, unsigned short z = 0, unsigned short w = 0) : x(x), y(y), z(z), w(w) {} -}; - -struct int1 { - int x; - __host__ __device__ int1(int x = 0) : x(x) {} -}; -struct int2 { - int x, y; - __host__ __device__ int2(int x = 0, int y = 0) : x(x), y(y) {} -}; -struct int4 { - int x, y, z, w; - __host__ __device__ int4(int x = 0, int y = 0, int z = 0, int w = 0) : x(x), y(y), z(z), w(w) {} -}; - -struct uint1 { - unsigned x; - __host__ __device__ uint1(unsigned x = 0) : x(x) {} -}; -struct uint2 { - unsigned x, y; - __host__ __device__ uint2(unsigned x = 0, unsigned y = 0) : x(x), y(y) {} -}; -struct uint3 { - unsigned x, y, z; - __host__ __device__ uint3(unsigned x = 0, unsigned y = 0, unsigned z = 0) : x(x), y(y), z(z) {} -}; -struct uint4 { - unsigned x, y, z, w; - __host__ __device__ uint4(unsigned x = 0, unsigned y = 0, unsigned z = 0, unsigned w = 0) : x(x), y(y), z(z), w(w) {} -}; - -struct longlong1 { - long long x; - __host__ __device__ longlong1(long long x = 0) : x(x) {} -}; -struct longlong2 { - long long x, y; - __host__ __device__ longlong2(long long x = 0, long long y = 0) : x(x), y(y) {} -}; -struct longlong4 { - long long x, y, z, w; - __host__ __device__ longlong4(long long x = 0, long long y = 0, long long z = 0, long long w = 0) : x(x), y(y), z(z), w(w) {} -}; - -struct ulonglong1 { - unsigned long long x; - __host__ __device__ ulonglong1(unsigned long long x = 0) : x(x) {} -}; -struct ulonglong2 { - unsigned long long x, y; - __host__ __device__ ulonglong2(unsigned long long x = 0, unsigned long long y = 0) : x(x), y(y) {} -}; -struct ulonglong4 { - unsigned long long x, y, z, w; - __host__ __device__ ulonglong4(unsigned long long x = 0, unsigned long long y = 0, unsigned long long z = 0, unsigned long long w = 0) : x(x), y(y), z(z), w(w) {} -}; - -struct float1 { - float x; - __host__ __device__ float1(float x = 0) : x(x) {} -}; -struct float2 { - float x, y; - __host__ __device__ float2(float x = 0, float y = 0) : x(x), y(y) {} -}; -struct float4 { - float x, y, z, w; - __host__ __device__ float4(float x = 0, float y = 0, float z = 0, float w = 0) : x(x), y(y), z(z), w(w) {} -}; - -struct double1 { - double x; - __host__ __device__ double1(double x = 0) : x(x) {} -}; -struct double2 { - double x, y; - __host__ __device__ double2(double x = 0, double y = 0) : x(x), y(y) {} -}; -struct double4 { - double x, y, z, w; - __host__ __device__ double4(double x = 0, double y = 0, double z = 0, double w = 0) : x(x), y(y), z(z), w(w) {} -}; - -typedef unsigned long long cudaTextureObject_t; -typedef unsigned long long cudaSurfaceObject_t; - -enum cudaTextureReadMode { - cudaReadModeNormalizedFloat, - cudaReadModeElementType -}; - -enum cudaSurfaceBoundaryMode { - cudaBoundaryModeZero, - cudaBoundaryModeClamp, - cudaBoundaryModeTrap -}; - -enum { - cudaTextureType1D, - cudaTextureType2D, - cudaTextureType3D, - cudaTextureTypeCubemap, - cudaTextureType1DLayered, - cudaTextureType2DLayered, - cudaTextureTypeCubemapLayered -}; - -struct textureReference { }; -template <class T, int texType = cudaTextureType1D, - enum cudaTextureReadMode mode = cudaReadModeElementType> -struct __attribute__((device_builtin_texture_type)) texture - : public textureReference {}; - -struct surfaceReference { int desc; }; - -template <typename T, int dim = 1> -struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {}; diff --git a/clang/test/CodeGenCUDA/mangling.cu b/clang/test/CodeGenCUDA/mangling.cu new file mode 100644 index 0000000000000..437ae07f03725 --- /dev/null +++ b/clang/test/CodeGenCUDA/mangling.cu @@ -0,0 +1,81 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -x cuda -emit-cir -target-sdk-version=12.3 %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir -fcuda-is-device -emit-cir -target-sdk-version=12.3 %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -x cuda -emit-llvm -target-sdk-version=12.3 %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM-HOST --input-file=%t.ll %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir -fcuda-is-device -emit-llvm -target-sdk-version=12.3 %s -o %t.ll +// RUN: 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 { + __global__ void cpp_global_function_1(int a, int* b, float c) {} + __global__ void cpp_global_function_2(int a, int* b, float c) {} + __host__ void cpp_host_function_1(int a, int* b, float c) {} + __host__ void cpp_host_function_2(int a, int* b, float c) {} + __device__ void cpp_device_function_1(int a, int* b, float c) {} + __device__ void cpp_device_function_2(int a, int* b, float c) {} +} + +__global__ void cpp_global_function_1(int a, int* b, float c) {} +__global__ void cpp_global_function_2(int a, int* b, float c) {} +__host__ void cpp_host_function_1(int a, int* b, float c) {} +__host__ void cpp_host_function_2(int a, int* b, float c) {} +__device__ void cpp_device_function_1(int a, int* b, float c) {} +__device__ void cpp_device_function_2(int a, int* b, float c) {} + +extern "C" { + __global__ void c_global_function_1(int a, int* b, float c) {} + __global__ void c_global_function_2(int a, int* b, float c) {} + __host__ void c_host_function_1(int a, int* b, float c) {} + __host__ void c_host_function_2(int a, int* b, float c) {} + __device__ void c_device_function_1(int a, int* b, float c) {} + __device__ void c_device_function_2(int a, int* b, float c) {} +} + +// 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 + +// CIR-HOST: cir.func {{.*}} @_ZN2ns36__device_stub__cpp_global_function_2EiPif +// CIR-DEVICE: cir.func {{.*}} @_ZN2ns21cpp_global_function_2EiPif + +// CIR-HOST: cir.func {{.*}} @_ZN2ns19cpp_host_function_1EiPif +// CIR-HOST: cir.func {{.*}} @_ZN2ns19cpp_host_function_2EiPif + +// CIR-DEVICE: cir.func {{.*}} @_ZN2ns21cpp_device_function_1EiPif +// CIR-DEVICE: cir.func {{.*}} @_ZN2ns21cpp_device_function_2EiPif + +// CIR-HOST: cir.func {{.*}} @_Z36__device_stub__cpp_global_function_1iPif +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_global_function_1iPif + +// CIR-HOST: cir.func {{.*}} @_Z36__device_stub__cpp_global_function_2iPif +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_global_function_2iPif + +// CIR-HOST: cir.func {{.*}} @_Z19cpp_host_function_1iPif +// CIR-HOST: cir.func {{.*}} @_Z19cpp_host_function_2iPif + +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_device_function_1iPif +// CIR-DEVICE: cir.func {{.*}} @_Z21cpp_device_function_2iPif + +// CIR-HOST: cir.func {{.*}} @__device_stub__c_global_function_1 +// CIR-DEVICE: cir.func {{.*}} @c_global_function_1 + +// CIR-HOST: cir.func {{.*}} @__device_stub__c_global_function_2 +// CIR-DEVICE: cir.func {{.*}} @c_global_function_2 + +// CIR-HOST: cir.func {{.*}} @c_host_function_1 +// CIR-HOST: cir.func {{.*}} @c_host_function_2 + +// CIR-DEVICE: cir.func {{.*}} @c_device_function_1 +// CIR-DEVICE: cir.func {{.*}} @c_device_function_2 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
