Author: Austin Schuh Date: 2025-04-04T15:52:44-07:00 New Revision: bdff739c7e79933aa7b3233da1bdadceeb7e03e6
URL: https://github.com/llvm/llvm-project/commit/bdff739c7e79933aa7b3233da1bdadceeb7e03e6 DIFF: https://github.com/llvm/llvm-project/commit/bdff739c7e79933aa7b3233da1bdadceeb7e03e6.diff LOG: cuda clang: Clean up test dependency for CUDA surfaces (#134459) https://github.com/llvm/llvm-project/pull/132883 added support for cuda surfaces but reached into clang/test/Headers/ from clang/test/CodeGen/ to grab the minimal cuda.h. Duplicate that file instead based on comments in the review, to fix remote test runs. Signed-off-by: Austin Schuh <austin.li...@gmail.com> Added: clang/test/CodeGen/include/cuda.h Modified: clang/test/CodeGen/nvptx-surface.cu clang/test/Headers/Inputs/include/cuda.h Removed: ################################################################################ diff --git a/clang/test/CodeGen/include/cuda.h b/clang/test/CodeGen/include/cuda.h new file mode 100644 index 0000000000000..58202442e1f8c --- /dev/null +++ b/clang/test/CodeGen/include/cuda.h @@ -0,0 +1,194 @@ +/* Minimal declarations for CUDA support. Testing purposes only. + * This should stay in sync with clang/test/Headers/Inputs/include/cuda.h + */ +#pragma once + +// Make this file work with nvcc, for testing compatibility. + +#ifndef __NVCC__ +#define __constant__ __attribute__((constant)) +#define __device__ __attribute__((device)) +#define __global__ __attribute__((global)) +#define __host__ __attribute__((host)) +#define __shared__ __attribute__((shared)) +#define __managed__ __attribute__((managed)) +#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__))) + +struct dim3 { + unsigned x, y, z; + __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {} +}; + +// Host- and device-side placement new overloads. +void *operator new(__SIZE_TYPE__, void *p) { return p; } +void *operator new[](__SIZE_TYPE__, void *p) { return p; } +__device__ void *operator new(__SIZE_TYPE__, void *p) { return p; } +__device__ void *operator new[](__SIZE_TYPE__, void *p) { return p; } + +#define CUDA_VERSION 10100 + +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 {}; + +#endif // !__NVCC__ diff --git a/clang/test/CodeGen/nvptx-surface.cu b/clang/test/CodeGen/nvptx-surface.cu index 7c42e5d118153..cf1fe76893a17 100644 --- a/clang/test/CodeGen/nvptx-surface.cu +++ b/clang/test/CodeGen/nvptx-surface.cu @@ -1,6 +1,6 @@ // RUN: %clang_cc1 -triple nvptx-unknown-unknown -fcuda-is-device -O3 -o - %s -emit-llvm | FileCheck %s // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device -O3 -o - %s -emit-llvm | FileCheck %s -#include "../Headers/Inputs/include/cuda.h" +#include "include/cuda.h" #include "__clang_cuda_texture_intrinsics.h" diff --git a/clang/test/Headers/Inputs/include/cuda.h b/clang/test/Headers/Inputs/include/cuda.h index 40a00b5af295a..4e8a2f46d27b3 100644 --- a/clang/test/Headers/Inputs/include/cuda.h +++ b/clang/test/Headers/Inputs/include/cuda.h @@ -1,4 +1,6 @@ -/* Minimal declarations for CUDA support. Testing purposes only. */ +/* Minimal declarations for CUDA support. Testing purposes only. + * This should stay in sync with clang/test/CodeGen/include/cuda.h + */ #pragma once // Make this file work with nvcc, for testing compatibility. _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
