[clang] cuda clang: Clean up test dependency for CUDA surfaces (PR #134459)
Austin Schuh via cfe-commits
cfe-commits at lists.llvm.org
Fri Apr 4 15:26:01 PDT 2025
https://github.com/AustinSchuh created https://github.com/llvm/llvm-project/pull/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.
>From 52d4f06540761507c4d81cead1278bff783fc54d Mon Sep 17 00:00:00 2001
From: Austin Schuh <austin.linux at gmail.com>
Date: Fri, 4 Apr 2025 15:23:15 -0700
Subject: [PATCH] cuda clang: Clean up test dependency for CUDA surfaces
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.linux at gmail.com>
---
clang/test/CodeGen/include/cuda.h | 194 +++++++++++++++++++++++
clang/test/CodeGen/nvptx-surface.cu | 2 +-
clang/test/Headers/Inputs/include/cuda.h | 4 +-
3 files changed, 198 insertions(+), 2 deletions(-)
create mode 100644 clang/test/CodeGen/include/cuda.h
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.
More information about the cfe-commits
mailing list