254 lines · c
1/* Minimal declarations for CUDA support. Testing purposes only. */2 3#include <stddef.h>4 5#if __HIP__ || __CUDA__6#define __constant__ __attribute__((constant))7#define __device__ __attribute__((device))8#define __global__ __attribute__((global))9#define __host__ __attribute__((host))10#define __shared__ __attribute__((shared))11#if __HIP__12#define __managed__ __attribute__((managed))13#endif14#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))15#define __grid_constant__ __attribute__((grid_constant))16#define __cluster_dims__(...) __attribute__((cluster_dims(__VA_ARGS__)))17#define __no_cluster__ __attribute__((no_cluster))18#else19#define __constant__20#define __device__21#define __global__22#define __host__23#define __shared__24#define __managed__25#define __launch_bounds__(...)26#define __grid_constant__27#define __cluster_dims__(...)28#define __no_cluster__29#endif30 31struct dim3 {32 unsigned x, y, z;33 __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}34};35 36#if __HIP__ || HIP_PLATFORM37typedef struct hipStream *hipStream_t;38typedef enum hipError {} hipError_t;39int hipConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,40 hipStream_t stream = 0);41extern "C" hipError_t __hipPushCallConfiguration(dim3 gridSize, dim3 blockSize,42 size_t sharedSize = 0,43 hipStream_t stream = 0);44#ifndef __HIP_API_PER_THREAD_DEFAULT_STREAM__45extern "C" hipError_t hipLaunchKernel(const void *func, dim3 gridDim,46 dim3 blockDim, void **args,47 size_t sharedMem,48 hipStream_t stream);49#else50extern "C" hipError_t hipLaunchKernel_spt(const void *func, dim3 gridDim,51 dim3 blockDim, void **args,52 size_t sharedMem,53 hipStream_t stream);54#endif // __HIP_API_PER_THREAD_DEFAULT_STREAM__55#elif __OFFLOAD_VIA_LLVM__56extern "C" unsigned __llvmPushCallConfiguration(dim3 gridDim, dim3 blockDim,57 size_t sharedMem = 0, void *stream = 0);58extern "C" unsigned llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,59 void **args, size_t sharedMem = 0, void *stream = 0);60#else61typedef struct cudaStream *cudaStream_t;62typedef enum cudaError {} cudaError_t;63extern "C" int cudaConfigureCall(dim3 gridSize, dim3 blockSize,64 size_t sharedSize = 0,65 cudaStream_t stream = 0);66extern "C" int __cudaPushCallConfiguration(dim3 gridSize, dim3 blockSize,67 size_t sharedSize = 0,68 cudaStream_t stream = 0);69extern "C" cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim,70 dim3 blockDim, void **args,71 size_t sharedMem, cudaStream_t stream);72extern "C" cudaError_t cudaLaunchKernel_ptsz(const void *func, dim3 gridDim,73 dim3 blockDim, void **args,74 size_t sharedMem, cudaStream_t stream);75extern "C" __device__ cudaError_t cudaLaunchDevice(void *func,76 void *parameterBuffer,77 dim3 gridDim, dim3 blockDim,78 unsigned int sharedMem,79 cudaStream_t stream);80extern "C" __device__ void *cudaGetParameterBuffer(size_t alignment,81 size_t size);82#endif83 84extern "C" __device__ int printf(const char*, ...);85 86struct char1 {87 char x;88 __host__ __device__ char1(char x = 0) : x(x) {}89};90struct char2 {91 char x, y;92 __host__ __device__ char2(char x = 0, char y = 0) : x(x), y(y) {}93};94struct char4 {95 char x, y, z, w;96 __host__ __device__ char4(char x = 0, char y = 0, char z = 0, char w = 0) : x(x), y(y), z(z), w(w) {}97};98 99struct uchar1 {100 unsigned char x;101 __host__ __device__ uchar1(unsigned char x = 0) : x(x) {}102};103struct uchar2 {104 unsigned char x, y;105 __host__ __device__ uchar2(unsigned char x = 0, unsigned char y = 0) : x(x), y(y) {}106};107struct uchar4 {108 unsigned char x, y, z, w;109 __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) {}110};111 112struct short1 {113 short x;114 __host__ __device__ short1(short x = 0) : x(x) {}115};116struct short2 {117 short x, y;118 __host__ __device__ short2(short x = 0, short y = 0) : x(x), y(y) {}119};120struct short4 {121 short x, y, z, w;122 __host__ __device__ short4(short x = 0, short y = 0, short z = 0, short w = 0) : x(x), y(y), z(z), w(w) {}123};124 125struct ushort1 {126 unsigned short x;127 __host__ __device__ ushort1(unsigned short x = 0) : x(x) {}128};129struct ushort2 {130 unsigned short x, y;131 __host__ __device__ ushort2(unsigned short x = 0, unsigned short y = 0) : x(x), y(y) {}132};133struct ushort4 {134 unsigned short x, y, z, w;135 __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) {}136};137 138struct int1 {139 int x;140 __host__ __device__ int1(int x = 0) : x(x) {}141};142struct int2 {143 int x, y;144 __host__ __device__ int2(int x = 0, int y = 0) : x(x), y(y) {}145};146struct int4 {147 int x, y, z, w;148 __host__ __device__ int4(int x = 0, int y = 0, int z = 0, int w = 0) : x(x), y(y), z(z), w(w) {}149};150 151struct uint1 {152 unsigned x;153 __host__ __device__ uint1(unsigned x = 0) : x(x) {}154};155struct uint2 {156 unsigned x, y;157 __host__ __device__ uint2(unsigned x = 0, unsigned y = 0) : x(x), y(y) {}158};159struct uint3 {160 unsigned x, y, z;161 __host__ __device__ uint3(unsigned x = 0, unsigned y = 0, unsigned z = 0) : x(x), y(y), z(z) {}162};163struct uint4 {164 unsigned x, y, z, w;165 __host__ __device__ uint4(unsigned x = 0, unsigned y = 0, unsigned z = 0, unsigned w = 0) : x(x), y(y), z(z), w(w) {}166};167 168struct longlong1 {169 long long x;170 __host__ __device__ longlong1(long long x = 0) : x(x) {}171};172struct longlong2 {173 long long x, y;174 __host__ __device__ longlong2(long long x = 0, long long y = 0) : x(x), y(y) {}175};176struct longlong4 {177 long long x, y, z, w;178 __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) {}179};180 181struct ulonglong1 {182 unsigned long long x;183 __host__ __device__ ulonglong1(unsigned long long x = 0) : x(x) {}184};185struct ulonglong2 {186 unsigned long long x, y;187 __host__ __device__ ulonglong2(unsigned long long x = 0, unsigned long long y = 0) : x(x), y(y) {}188};189struct ulonglong4 {190 unsigned long long x, y, z, w;191 __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) {}192};193 194struct float1 {195 float x;196 __host__ __device__ float1(float x = 0) : x(x) {}197};198struct float2 {199 float x, y;200 __host__ __device__ float2(float x = 0, float y = 0) : x(x), y(y) {}201};202struct float4 {203 float x, y, z, w;204 __host__ __device__ float4(float x = 0, float y = 0, float z = 0, float w = 0) : x(x), y(y), z(z), w(w) {}205};206 207struct double1 {208 double x;209 __host__ __device__ double1(double x = 0) : x(x) {}210};211struct double2 {212 double x, y;213 __host__ __device__ double2(double x = 0, double y = 0) : x(x), y(y) {}214};215struct double4 {216 double x, y, z, w;217 __host__ __device__ double4(double x = 0, double y = 0, double z = 0, double w = 0) : x(x), y(y), z(z), w(w) {}218};219 220typedef unsigned long long cudaTextureObject_t;221typedef unsigned long long cudaSurfaceObject_t;222 223enum cudaTextureReadMode {224 cudaReadModeNormalizedFloat,225 cudaReadModeElementType226};227 228enum cudaSurfaceBoundaryMode {229 cudaBoundaryModeZero,230 cudaBoundaryModeClamp,231 cudaBoundaryModeTrap232};233 234enum {235 cudaTextureType1D,236 cudaTextureType2D,237 cudaTextureType3D,238 cudaTextureTypeCubemap,239 cudaTextureType1DLayered,240 cudaTextureType2DLayered,241 cudaTextureTypeCubemapLayered242};243 244struct textureReference { };245template <class T, int texType = cudaTextureType1D,246 enum cudaTextureReadMode mode = cudaReadModeElementType>247struct __attribute__((device_builtin_texture_type)) texture248 : public textureReference {};249 250struct surfaceReference { int desc; };251 252template <typename T, int dim = 1>253struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {};254