2 * Copyright 1993-2012 NVIDIA Corporation. All rights reserved.
6 * This source code and/or documentation ("Licensed Deliverables") are
7 * subject to NVIDIA intellectual property rights under U.S. and
8 * international Copyright laws.
10 * These Licensed Deliverables contained herein is PROPRIETARY and
11 * CONFIDENTIAL to NVIDIA and is being provided under the terms and
12 * conditions of a form of NVIDIA software license agreement by and
13 * between NVIDIA and Licensee ("License Agreement") or electronically
14 * accepted by Licensee. Notwithstanding any terms or conditions to
15 * the contrary in the License Agreement, reproduction or disclosure
16 * of the Licensed Deliverables to any third party without the express
17 * written consent of NVIDIA is prohibited.
19 * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
20 * LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE
21 * SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS
22 * PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND.
23 * NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED
24 * DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY,
25 * NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
26 * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
27 * LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY
28 * SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY
29 * DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,
30 * WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS
31 * ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE
32 * OF THESE LICENSED DELIVERABLES.
34 * U.S. Government End Users. These Licensed Deliverables are a
35 * "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT
36 * 1995), consisting of "commercial computer software" and "commercial
37 * computer software documentation" as such terms are used in 48
38 * C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government
39 * only as a commercial end item. Consistent with 48 C.F.R.12.212 and
40 * 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all
41 * U.S. Government End Users acquire the Licensed Deliverables with
42 * only those rights set forth herein.
44 * Any use of the Licensed Deliverables in individual and commercial
45 * software must include, in the user documentation and internal
46 * comments to the code, the above Disclaimer and U.S. Government End
50 #if !defined(__SURFACE_FUNCTIONS_H__)
51 #define __SURFACE_FUNCTIONS_H__
53 #if defined(__cplusplus) && defined(__CUDACC__)
55 /*******************************************************************************
59 *******************************************************************************/
61 #include "builtin_types.h"
62 #include "cuda_surface_types.h"
63 #include "host_defines.h"
64 #include "surface_types.h"
65 #include "vector_functions.h"
66 #include "vector_types.h"
68 /*******************************************************************************
72 *******************************************************************************/
73 extern __device__ __device_builtin__ uchar1 __surf1Dreadc1(surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
74 extern __device__ __device_builtin__ uchar2 __surf1Dreadc2(surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
75 extern __device__ __device_builtin__ uchar4 __surf1Dreadc4(surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
76 extern __device__ __device_builtin__ ushort1 __surf1Dreads1(surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
77 extern __device__ __device_builtin__ ushort2 __surf1Dreads2(surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
78 extern __device__ __device_builtin__ ushort4 __surf1Dreads4(surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
79 extern __device__ __device_builtin__ uint1 __surf1Dreadu1(surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
80 extern __device__ __device_builtin__ uint2 __surf1Dreadu2(surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
81 extern __device__ __device_builtin__ uint4 __surf1Dreadu4(surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
82 extern __device__ __device_builtin__ ulonglong1 __surf1Dreadl1(surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
83 extern __device__ __device_builtin__ ulonglong2 __surf1Dreadl2(surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
85 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
87 #define __surfModeSwitch(surf, x, mode, type) \
88 ((mode == cudaBoundaryModeZero) ? __surf1Dread##type(surf, x, cudaBoundaryModeZero ) : \
89 (mode == cudaBoundaryModeClamp) ? __surf1Dread##type(surf, x, cudaBoundaryModeClamp) : \
90 __surf1Dread##type(surf, x, cudaBoundaryModeTrap ))
92 #else /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
94 #define __surfModeSwitch(surf, x, mode, type) \
95 __surf1Dread##type(surf, x, cudaBoundaryModeTrap)
97 #endif /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
100 static __forceinline__ __device__ void surf1Dread(T *res, surface<void, cudaSurfaceType1D> surf, int x, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
102 (s == 1) ? (void)(*(uchar1 *)res = __surfModeSwitch(surf, x, mode, c1)) :
103 (s == 2) ? (void)(*(ushort1*)res = __surfModeSwitch(surf, x, mode, s1)) :
104 (s == 4) ? (void)(*(uint1 *)res = __surfModeSwitch(surf, x, mode, u1)) :
105 (s == 8) ? (void)(*(uint2 *)res = __surfModeSwitch(surf, x, mode, u2)) :
106 (s == 16) ? (void)(*(uint4 *)res = __surfModeSwitch(surf, x, mode, u4)) :
111 static __forceinline__ __device__ T surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
115 surf1Dread(&tmp, surf, x, (int)sizeof(T), mode);
121 static __forceinline__ __device__ void surf1Dread(T *res, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
123 *res = surf1Dread<T>(surf, x, mode);
127 __forceinline__ __device__ char surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
129 return (char)__surfModeSwitch(surf, x, mode, c1).x;
133 __forceinline__ __device__ signed char surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
135 return (signed char)__surfModeSwitch(surf, x, mode, c1).x;
139 __forceinline__ __device__ unsigned char surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
141 return __surfModeSwitch(surf, x, mode, c1).x;
145 __forceinline__ __device__ char1 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
147 return make_char1((signed char)__surfModeSwitch(surf, x, mode, c1).x);
151 __forceinline__ __device__ uchar1 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
153 return __surfModeSwitch(surf, x, mode, c1);
157 __forceinline__ __device__ char2 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
159 uchar2 tmp = __surfModeSwitch(surf, x, mode, c2);
161 return make_char2((signed char)tmp.x, (signed char)tmp.y);
165 __forceinline__ __device__ uchar2 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
167 return __surfModeSwitch(surf, x, mode, c2);
171 __forceinline__ __device__ char4 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
173 uchar4 tmp = __surfModeSwitch(surf, x, mode, c4);
175 return make_char4((signed char)tmp.x, (signed char)tmp.y, (signed char)tmp.z, (signed char)tmp.w);
179 __forceinline__ __device__ uchar4 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
181 return __surfModeSwitch(surf, x, mode, c4);
185 __forceinline__ __device__ short surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
187 return (short)__surfModeSwitch(surf, x, mode, s1).x;
191 __forceinline__ __device__ unsigned short surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
193 return __surfModeSwitch(surf, x, mode, s1).x;
197 __forceinline__ __device__ short1 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
199 return make_short1((signed short)__surfModeSwitch(surf, x, mode, s1).x);
203 __forceinline__ __device__ ushort1 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
205 return __surfModeSwitch(surf, x, mode, s1);
209 __forceinline__ __device__ short2 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
211 ushort2 tmp = __surfModeSwitch(surf, x, mode, s2);
213 return make_short2((signed short)tmp.x, (signed short)tmp.y);
217 __forceinline__ __device__ ushort2 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
219 return __surfModeSwitch(surf, x, mode, s2);
223 __forceinline__ __device__ short4 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
225 ushort4 tmp = __surfModeSwitch(surf, x, mode, s4);
227 return make_short4((signed short)tmp.x, (signed short)tmp.y, (signed short)tmp.z, (signed short)tmp.w);
231 __forceinline__ __device__ ushort4 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
233 return __surfModeSwitch(surf, x, mode, s4);
237 __forceinline__ __device__ int surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
239 return (int)__surfModeSwitch(surf, x, mode, u1).x;
243 __forceinline__ __device__ unsigned int surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
245 return __surfModeSwitch(surf, x, mode, u1).x;
249 __forceinline__ __device__ int1 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
251 return make_int1((signed int)__surfModeSwitch(surf, x, mode, u1).x);
255 __forceinline__ __device__ uint1 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
257 return __surfModeSwitch(surf, x, mode, u1);
261 __forceinline__ __device__ int2 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
263 uint2 tmp = __surfModeSwitch(surf, x, mode, u2);
265 return make_int2((int)tmp.x, (int)tmp.y);
269 __forceinline__ __device__ uint2 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
271 return __surfModeSwitch(surf, x, mode, u2);
275 __forceinline__ __device__ int4 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
277 uint4 tmp = __surfModeSwitch(surf, x, mode, u4);
279 return make_int4((int)tmp.x, (int)tmp.y, (int)tmp.z, (int)tmp.w);
283 __forceinline__ __device__ uint4 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
285 return __surfModeSwitch(surf, x, mode, u4);
289 __forceinline__ __device__ long long int surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
291 return (long long int)__surfModeSwitch(surf, x, mode, l1).x;
295 __forceinline__ __device__ unsigned long long int surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
297 return __surfModeSwitch(surf, x, mode, l1).x;
301 __forceinline__ __device__ longlong1 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
303 return make_longlong1((long long int)__surfModeSwitch(surf, x, mode, l1).x);
307 __forceinline__ __device__ ulonglong1 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
309 return __surfModeSwitch(surf, x, mode, l1);
313 __forceinline__ __device__ longlong2 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
315 ulonglong2 tmp = __surfModeSwitch(surf, x, mode, l2);
317 return make_longlong2((long long int)tmp.x, (long long int)tmp.y);
321 __forceinline__ __device__ ulonglong2 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
323 return __surfModeSwitch(surf, x, mode, l2);
326 #if !defined(__LP64__)
329 __forceinline__ __device__ long int surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
331 return (long int)__surfModeSwitch(surf, x, mode, u1).x;
335 __forceinline__ __device__ unsigned long int surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
337 return (unsigned long int)__surfModeSwitch(surf, x, mode, u1).x;
341 __forceinline__ __device__ long1 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
343 return make_long1((long int)__surfModeSwitch(surf, x, mode, u1).x);
347 __forceinline__ __device__ ulong1 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
349 return make_ulong1((unsigned long int)__surfModeSwitch(surf, x, mode, u1).x);
353 __forceinline__ __device__ long2 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
355 uint2 tmp = __surfModeSwitch(surf, x, mode, u2);
357 return make_long2((long int)tmp.x, (long int)tmp.y);
361 __forceinline__ __device__ ulong2 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
363 uint2 tmp = __surfModeSwitch(surf, x, mode, u2);
365 return make_ulong2((unsigned long int)tmp.x, (unsigned long int)tmp.y);
369 __forceinline__ __device__ long4 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
371 uint4 tmp = __surfModeSwitch(surf, x, mode, u4);
373 return make_long4((long int)tmp.x, (long int)tmp.y, (long int)tmp.z, (long int)tmp.w);
377 __forceinline__ __device__ ulong4 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
379 uint4 tmp = __surfModeSwitch(surf, x, mode, u4);
381 return make_ulong4((unsigned long int)tmp.x, (unsigned long int)tmp.y, (unsigned long int)tmp.z, (unsigned long int)tmp.w);
384 #endif /* !__LP64__ */
387 __forceinline__ __device__ float surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
389 return __int_as_float((int)__surfModeSwitch(surf, x, mode, u1).x);
393 __forceinline__ __device__ float1 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
395 return make_float1(__int_as_float((int)__surfModeSwitch(surf, x, mode, u1).x));
399 __forceinline__ __device__ float2 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
401 uint2 tmp = __surfModeSwitch(surf, x, mode, u2);
403 return make_float2(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y));
407 __forceinline__ __device__ float4 surf1Dread(surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode)
409 uint4 tmp = __surfModeSwitch(surf, x, mode, u4);
411 return make_float4(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y), __int_as_float((int)tmp.z), __int_as_float((int)tmp.w));
414 #undef __surfModeSwitch
416 /*******************************************************************************
420 *******************************************************************************/
421 extern __device__ __device_builtin__ uchar1 __surf2Dreadc1(surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
422 extern __device__ __device_builtin__ uchar2 __surf2Dreadc2(surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
423 extern __device__ __device_builtin__ uchar4 __surf2Dreadc4(surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
424 extern __device__ __device_builtin__ ushort1 __surf2Dreads1(surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
425 extern __device__ __device_builtin__ ushort2 __surf2Dreads2(surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
426 extern __device__ __device_builtin__ ushort4 __surf2Dreads4(surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
427 extern __device__ __device_builtin__ uint1 __surf2Dreadu1(surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
428 extern __device__ __device_builtin__ uint2 __surf2Dreadu2(surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
429 extern __device__ __device_builtin__ uint4 __surf2Dreadu4(surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
430 extern __device__ __device_builtin__ ulonglong1 __surf2Dreadl1(surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
431 extern __device__ __device_builtin__ ulonglong2 __surf2Dreadl2(surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
433 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
435 #define __surfModeSwitch(surf, x, y, mode, type) \
436 ((mode == cudaBoundaryModeZero) ? __surf2Dread##type(surf, x, y, cudaBoundaryModeZero ) : \
437 (mode == cudaBoundaryModeClamp) ? __surf2Dread##type(surf, x, y, cudaBoundaryModeClamp) : \
438 __surf2Dread##type(surf, x, y, cudaBoundaryModeTrap ))
440 #else /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
442 #define __surfModeSwitch(surf, x, y, mode, type) \
443 __surf2Dread##type(surf, x, y, cudaBoundaryModeTrap)
445 #endif /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
448 static __forceinline__ __device__ void surf2Dread(T *res, surface<void, cudaSurfaceType2D> surf, int x, int y, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
450 (s == 1) ? (void)(*(uchar1 *)res = __surfModeSwitch(surf, x, y, mode, c1)) :
451 (s == 2) ? (void)(*(ushort1*)res = __surfModeSwitch(surf, x, y, mode, s1)) :
452 (s == 4) ? (void)(*(uint1 *)res = __surfModeSwitch(surf, x, y, mode, u1)) :
453 (s == 8) ? (void)(*(uint2 *)res = __surfModeSwitch(surf, x, y, mode, u2)) :
454 (s == 16) ? (void)(*(uint4 *)res = __surfModeSwitch(surf, x, y, mode, u4)) :
459 static __forceinline__ __device__ T surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
463 surf2Dread(&tmp, surf, x, y, (int)sizeof(T), mode);
469 static __forceinline__ __device__ void surf2Dread(T *res, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
471 *res = surf2Dread<T>(surf, x, y, mode);
475 __forceinline__ __device__ char surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
477 return (char)__surfModeSwitch(surf, x, y, mode, c1).x;
481 __forceinline__ __device__ signed char surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
483 return (signed char)__surfModeSwitch(surf, x, y, mode, c1).x;
487 __forceinline__ __device__ unsigned char surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
489 return __surfModeSwitch(surf, x, y, mode, c1).x;
493 __forceinline__ __device__ char1 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
495 return make_char1((signed char)__surfModeSwitch(surf, x, y, mode, c1).x);
499 __forceinline__ __device__ uchar1 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
501 return __surfModeSwitch(surf, x, y, mode, c1);
505 __forceinline__ __device__ char2 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
507 uchar2 tmp = __surfModeSwitch(surf, x, y, mode, c2);
509 return make_char2((signed char)tmp.x, (signed char)tmp.y);
513 __forceinline__ __device__ uchar2 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
515 return __surfModeSwitch(surf, x, y, mode, c2);
519 __forceinline__ __device__ char4 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
521 uchar4 tmp = __surfModeSwitch(surf, x, y, mode, c4);
523 return make_char4((signed char)tmp.x, (signed char)tmp.y, (signed char)tmp.z, (signed char)tmp.w);
527 __forceinline__ __device__ uchar4 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
529 return __surfModeSwitch(surf, x, y, mode, c4);
533 __forceinline__ __device__ short surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
535 return (short)__surfModeSwitch(surf, x, y, mode, s1).x;
539 __forceinline__ __device__ unsigned short surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
541 return __surfModeSwitch(surf, x, y, mode, s1).x;
545 __forceinline__ __device__ short1 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
547 return make_short1((signed short)__surfModeSwitch(surf, x, y, mode, s1).x);
551 __forceinline__ __device__ ushort1 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
553 return __surfModeSwitch(surf, x, y, mode, s1);
557 __forceinline__ __device__ short2 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
559 ushort2 tmp = __surfModeSwitch(surf, x, y, mode, s2);
561 return make_short2((signed short)tmp.x, (signed short)tmp.y);
565 __forceinline__ __device__ ushort2 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
567 return __surfModeSwitch(surf, x, y, mode, s2);
571 __forceinline__ __device__ short4 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
573 ushort4 tmp = __surfModeSwitch(surf, x, y, mode, s4);
575 return make_short4((signed short)tmp.x, (signed short)tmp.y, (signed short)tmp.z, (signed short)tmp.w);
579 __forceinline__ __device__ ushort4 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
581 return __surfModeSwitch(surf, x, y, mode, s4);
585 __forceinline__ __device__ int surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
587 return (int)__surfModeSwitch(surf, x, y, mode, u1).x;
591 __forceinline__ __device__ unsigned int surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
593 return __surfModeSwitch(surf, x, y, mode, u1).x;
597 __forceinline__ __device__ int1 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
599 return make_int1((signed int)__surfModeSwitch(surf, x, y, mode, u1).x);
603 __forceinline__ __device__ uint1 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
605 return __surfModeSwitch(surf, x, y, mode, u1);
609 __forceinline__ __device__ int2 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
611 uint2 tmp = __surfModeSwitch(surf, x, y, mode, u2);
613 return make_int2((int)tmp.x, (int)tmp.y);
617 __forceinline__ __device__ uint2 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
619 return __surfModeSwitch(surf, x, y, mode, u2);
623 __forceinline__ __device__ int4 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
625 uint4 tmp = __surfModeSwitch(surf, x, y, mode, u4);
627 return make_int4((int)tmp.x, (int)tmp.y, (int)tmp.z, (int)tmp.w);
631 __forceinline__ __device__ uint4 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
633 return __surfModeSwitch(surf, x, y, mode, u4);
637 __forceinline__ __device__ long long int surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
639 return (long long int)__surfModeSwitch(surf, x, y, mode, l1).x;
643 __forceinline__ __device__ unsigned long long int surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
645 return __surfModeSwitch(surf, x, y, mode, l1).x;
649 __forceinline__ __device__ longlong1 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
651 return make_longlong1((long long int)__surfModeSwitch(surf, x, y, mode, l1).x);
655 __forceinline__ __device__ ulonglong1 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
657 return __surfModeSwitch(surf, x, y, mode, l1);
661 __forceinline__ __device__ longlong2 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
663 ulonglong2 tmp = __surfModeSwitch(surf, x, y, mode, l2);
665 return make_longlong2((long long int)tmp.x, (long long int)tmp.y);
669 __forceinline__ __device__ ulonglong2 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
671 return __surfModeSwitch(surf, x, y, mode, l2);
674 #if !defined(__LP64__)
677 __forceinline__ __device__ long int surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
679 return (long int)__surfModeSwitch(surf, x, y, mode, u1).x;
683 __forceinline__ __device__ unsigned long int surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
685 return (unsigned long int)__surfModeSwitch(surf, x, y, mode, u1).x;
689 __forceinline__ __device__ long1 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
691 return make_long1((long int)__surfModeSwitch(surf, x, y, mode, u1).x);
695 __forceinline__ __device__ ulong1 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
697 return make_ulong1((unsigned long int)__surfModeSwitch(surf, x, y, mode, u1).x);
701 __forceinline__ __device__ long2 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
703 uint2 tmp = __surfModeSwitch(surf, x, y, mode, u2);
705 return make_long2((long int)tmp.x, (long int)tmp.y);
709 __forceinline__ __device__ ulong2 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
711 uint2 tmp = __surfModeSwitch(surf, x, y, mode, u2);
713 return make_ulong2((unsigned long int)tmp.x, (unsigned long int)tmp.y);
717 __forceinline__ __device__ long4 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
719 uint4 tmp = __surfModeSwitch(surf, x, y, mode, u4);
721 return make_long4((long int)tmp.x, (long int)tmp.y, (long int)tmp.z, (long int)tmp.w);
725 __forceinline__ __device__ ulong4 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
727 uint4 tmp = __surfModeSwitch(surf, x, y, mode, u4);
729 return make_ulong4((unsigned long int)tmp.x, (unsigned long int)tmp.y, (unsigned long int)tmp.z, (unsigned long int)tmp.w);
732 #endif /* !__LP64__ */
735 __forceinline__ __device__ float surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
737 return __int_as_float((int)__surfModeSwitch(surf, x, y, mode, u1).x);
741 __forceinline__ __device__ float1 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
743 return make_float1(__int_as_float((int)__surfModeSwitch(surf, x, y, mode, u1).x));
747 __forceinline__ __device__ float2 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
749 uint2 tmp = __surfModeSwitch(surf, x, y, mode, u2);
751 return make_float2(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y));
755 __forceinline__ __device__ float4 surf2Dread(surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode)
757 uint4 tmp = __surfModeSwitch(surf, x, y, mode, u4);
759 return make_float4(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y), __int_as_float((int)tmp.z), __int_as_float((int)tmp.w));
762 #undef __surfModeSwitch
764 /*******************************************************************************
768 *******************************************************************************/
769 extern __device__ __device_builtin__ uchar1 __surf3Dreadc1(surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
770 extern __device__ __device_builtin__ uchar2 __surf3Dreadc2(surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
771 extern __device__ __device_builtin__ uchar4 __surf3Dreadc4(surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
772 extern __device__ __device_builtin__ ushort1 __surf3Dreads1(surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
773 extern __device__ __device_builtin__ ushort2 __surf3Dreads2(surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
774 extern __device__ __device_builtin__ ushort4 __surf3Dreads4(surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
775 extern __device__ __device_builtin__ uint1 __surf3Dreadu1(surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
776 extern __device__ __device_builtin__ uint2 __surf3Dreadu2(surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
777 extern __device__ __device_builtin__ uint4 __surf3Dreadu4(surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
778 extern __device__ __device_builtin__ ulonglong1 __surf3Dreadl1(surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
779 extern __device__ __device_builtin__ ulonglong2 __surf3Dreadl2(surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
781 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
783 #define __surfModeSwitch(surf, x, y, z, mode, type) \
784 ((mode == cudaBoundaryModeZero) ? __surf3Dread##type(surf, x, y, z, cudaBoundaryModeZero ) : \
785 (mode == cudaBoundaryModeClamp) ? __surf3Dread##type(surf, x, y, z, cudaBoundaryModeClamp) : \
786 __surf3Dread##type(surf, x, y, z, cudaBoundaryModeTrap ))
788 #else /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
790 #define __surfModeSwitch(surf, x, y, z, mode, type) \
791 __surf3Dread##type(surf, x, y, z, cudaBoundaryModeTrap)
793 #endif /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
796 static __forceinline__ __device__ void surf3Dread(T *res, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
798 (s == 1) ? (void)(*(uchar1 *)res = __surfModeSwitch(surf, x, y, z, mode, c1)) :
799 (s == 2) ? (void)(*(ushort1*)res = __surfModeSwitch(surf, x, y, z, mode, s1)) :
800 (s == 4) ? (void)(*(uint1 *)res = __surfModeSwitch(surf, x, y, z, mode, u1)) :
801 (s == 8) ? (void)(*(uint2 *)res = __surfModeSwitch(surf, x, y, z, mode, u2)) :
802 (s == 16) ? (void)(*(uint4 *)res = __surfModeSwitch(surf, x, y, z, mode, u4)) :
807 static __forceinline__ __device__ T surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
811 surf3Dread(&tmp, surf, x, y, z, (int)sizeof(T), mode);
817 static __forceinline__ __device__ void surf3Dread(T *res, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
819 *res = surf3Dread<T>(surf, x, y, z, mode);
823 __forceinline__ __device__ char surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
825 return (char)__surfModeSwitch(surf, x, y, z, mode, c1).x;
829 __forceinline__ __device__ signed char surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
831 return (signed char)__surfModeSwitch(surf, x, y, z, mode, c1).x;
835 __forceinline__ __device__ unsigned char surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
837 return __surfModeSwitch(surf, x, y, z, mode, c1).x;
841 __forceinline__ __device__ char1 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
843 return make_char1((signed char)__surfModeSwitch(surf, x, y, z, mode, c1).x);
847 __forceinline__ __device__ uchar1 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
849 return __surfModeSwitch(surf, x, y, z, mode, c1);
853 __forceinline__ __device__ char2 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
855 uchar2 tmp = __surfModeSwitch(surf, x, y, z, mode, c2);
857 return make_char2((signed char)tmp.x, (signed char)tmp.y);
861 __forceinline__ __device__ uchar2 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
863 return __surfModeSwitch(surf, x, y, z, mode, c2);
867 __forceinline__ __device__ char4 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
869 uchar4 tmp = __surfModeSwitch(surf, x, y, z, mode, c4);
871 return make_char4((signed char)tmp.x, (signed char)tmp.y, (signed char)tmp.z, (signed char)tmp.w);
875 __forceinline__ __device__ uchar4 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
877 return __surfModeSwitch(surf, x, y, z, mode, c4);
881 __forceinline__ __device__ short surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
883 return (short)__surfModeSwitch(surf, x, y, z, mode, s1).x;
887 __forceinline__ __device__ unsigned short surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
889 return __surfModeSwitch(surf, x, y, z, mode, s1).x;
893 __forceinline__ __device__ short1 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
895 return make_short1((signed short)__surfModeSwitch(surf, x, y, z, mode, s1).x);
899 __forceinline__ __device__ ushort1 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
901 return __surfModeSwitch(surf, x, y, z, mode, s1);
905 __forceinline__ __device__ short2 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
907 ushort2 tmp = __surfModeSwitch(surf, x, y, z, mode, s2);
909 return make_short2((signed short)tmp.x, (signed short)tmp.y);
913 __forceinline__ __device__ ushort2 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
915 return __surfModeSwitch(surf, x, y, z, mode, s2);
919 __forceinline__ __device__ short4 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
921 ushort4 tmp = __surfModeSwitch(surf, x, y, z, mode, s4);
923 return make_short4((signed short)tmp.x, (signed short)tmp.y, (signed short)tmp.z, (signed short)tmp.w);
927 __forceinline__ __device__ ushort4 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
929 return __surfModeSwitch(surf, x, y, z, mode, s4);
933 __forceinline__ __device__ int surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
935 return (int)__surfModeSwitch(surf, x, y, z, mode, u1).x;
939 __forceinline__ __device__ unsigned int surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
941 return __surfModeSwitch(surf, x, y, z, mode, u1).x;
945 __forceinline__ __device__ int1 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
947 return make_int1((signed int)__surfModeSwitch(surf, x, y, z, mode, u1).x);
951 __forceinline__ __device__ uint1 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
953 return __surfModeSwitch(surf, x, y, z, mode, u1);
957 __forceinline__ __device__ int2 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
959 uint2 tmp = __surfModeSwitch(surf, x, y, z, mode, u2);
961 return make_int2((int)tmp.x, (int)tmp.y);
965 __forceinline__ __device__ uint2 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
967 return __surfModeSwitch(surf, x, y, z, mode, u2);
971 __forceinline__ __device__ int4 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
973 uint4 tmp = __surfModeSwitch(surf, x, y, z, mode, u4);
975 return make_int4((int)tmp.x, (int)tmp.y, (int)tmp.z, (int)tmp.w);
979 __forceinline__ __device__ uint4 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
981 return __surfModeSwitch(surf, x, y, z, mode, u4);
985 __forceinline__ __device__ long long int surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
987 return (long long int)__surfModeSwitch(surf, x, y, z, mode, l1).x;
991 __forceinline__ __device__ unsigned long long int surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
993 return __surfModeSwitch(surf, x, y, z, mode, l1).x;
997 __forceinline__ __device__ longlong1 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
999 return make_longlong1((long long int)__surfModeSwitch(surf, x, y, z, mode, l1).x);
1003 __forceinline__ __device__ ulonglong1 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1005 return __surfModeSwitch(surf, x, y, z, mode, l1);
1009 __forceinline__ __device__ longlong2 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1011 ulonglong2 tmp = __surfModeSwitch(surf, x, y, z, mode, l2);
1013 return make_longlong2((long long int)tmp.x, (long long int)tmp.y);
1017 __forceinline__ __device__ ulonglong2 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1019 return __surfModeSwitch(surf, x, y, z, mode, l2);
1022 #if !defined(__LP64__)
1025 __forceinline__ __device__ long int surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1027 return (long int)__surfModeSwitch(surf, x, y, z, mode, u1).x;
1031 __forceinline__ __device__ unsigned long int surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1033 return (unsigned long int)__surfModeSwitch(surf, x, y, z, mode, u1).x;
1037 __forceinline__ __device__ long1 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1039 return make_long1((long int)__surfModeSwitch(surf, x, y, z, mode, u1).x);
1043 __forceinline__ __device__ ulong1 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1045 return make_ulong1((unsigned long int)__surfModeSwitch(surf, x, y, z, mode, u1).x);
1049 __forceinline__ __device__ long2 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1051 uint2 tmp = __surfModeSwitch(surf, x, y, z, mode, u2);
1053 return make_long2((long int)tmp.x, (long int)tmp.y);
1057 __forceinline__ __device__ ulong2 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1059 uint2 tmp = __surfModeSwitch(surf, x, y, z, mode, u2);
1061 return make_ulong2((unsigned long int)tmp.x, (unsigned long int)tmp.y);
1065 __forceinline__ __device__ long4 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1067 uint4 tmp = __surfModeSwitch(surf, x, y, z, mode, u4);
1069 return make_long4((long int)tmp.x, (long int)tmp.y, (long int)tmp.z, (long int)tmp.w);
1073 __forceinline__ __device__ ulong4 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1075 uint4 tmp = __surfModeSwitch(surf, x, y, z, mode, u4);
1077 return make_ulong4((unsigned long int)tmp.x, (unsigned long int)tmp.y, (unsigned long int)tmp.z, (unsigned long int)tmp.w);
1080 #endif /* !__LP64__ */
1083 __forceinline__ __device__ float surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1085 return __int_as_float((int)__surfModeSwitch(surf, x, y, z, mode, u1).x);
1089 __forceinline__ __device__ float1 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1091 return make_float1(__int_as_float((int)__surfModeSwitch(surf, x, y, z, mode, u1).x));
1095 __forceinline__ __device__ float2 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1097 uint2 tmp = __surfModeSwitch(surf, x, y, z, mode, u2);
1099 return make_float2(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y));
1103 __forceinline__ __device__ float4 surf3Dread(surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode)
1105 uint4 tmp = __surfModeSwitch(surf, x, y, z, mode, u4);
1107 return make_float4(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y), __int_as_float((int)tmp.z), __int_as_float((int)tmp.w));
1110 #undef __surfModeSwitch
1112 /*******************************************************************************
1116 *******************************************************************************/
1117 extern __device__ __device_builtin__ uchar1 __surf1DLayeredreadc1(surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
1118 extern __device__ __device_builtin__ uchar2 __surf1DLayeredreadc2(surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
1119 extern __device__ __device_builtin__ uchar4 __surf1DLayeredreadc4(surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
1120 extern __device__ __device_builtin__ ushort1 __surf1DLayeredreads1(surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
1121 extern __device__ __device_builtin__ ushort2 __surf1DLayeredreads2(surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
1122 extern __device__ __device_builtin__ ushort4 __surf1DLayeredreads4(surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
1123 extern __device__ __device_builtin__ uint1 __surf1DLayeredreadu1(surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
1124 extern __device__ __device_builtin__ uint2 __surf1DLayeredreadu2(surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
1125 extern __device__ __device_builtin__ uint4 __surf1DLayeredreadu4(surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
1126 extern __device__ __device_builtin__ ulonglong1 __surf1DLayeredreadl1(surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
1127 extern __device__ __device_builtin__ ulonglong2 __surf1DLayeredreadl2(surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
1129 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
1131 #define __surfModeSwitch(surf, x, layer, mode, type) \
1132 ((mode == cudaBoundaryModeZero) ? __surf1DLayeredread##type(surf, x, layer, cudaBoundaryModeZero ) : \
1133 (mode == cudaBoundaryModeClamp) ? __surf1DLayeredread##type(surf, x, layer, cudaBoundaryModeClamp) : \
1134 __surf1DLayeredread##type(surf, x, layer, cudaBoundaryModeTrap ))
1136 #else /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
1138 #define __surfModeSwitch(surf, x, layer, mode, type) \
1139 __surf1DLayeredread##type(surf, x, layer, cudaBoundaryModeTrap)
1141 #endif /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
1144 static __forceinline__ __device__ void surf1DLayeredread(T *res, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
1146 (s == 1) ? (void)(*(uchar1 *)res = __surfModeSwitch(surf, x, layer, mode, c1)) :
1147 (s == 2) ? (void)(*(ushort1*)res = __surfModeSwitch(surf, x, layer, mode, s1)) :
1148 (s == 4) ? (void)(*(uint1 *)res = __surfModeSwitch(surf, x, layer, mode, u1)) :
1149 (s == 8) ? (void)(*(uint2 *)res = __surfModeSwitch(surf, x, layer, mode, u2)) :
1150 (s == 16) ? (void)(*(uint4 *)res = __surfModeSwitch(surf, x, layer, mode, u4)) :
1155 static __forceinline__ __device__ T surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
1159 surf1DLayeredread(&tmp, surf, x, layer, (int)sizeof(T), mode);
1165 static __forceinline__ __device__ void surf1DLayeredread(T *res, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
1167 *res = surf1DLayeredread<T>(surf, x, layer, mode);
1171 __forceinline__ __device__ char surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1173 return (char)__surfModeSwitch(surf, x, layer, mode, c1).x;
1177 __forceinline__ __device__ signed char surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1179 return (signed char)__surfModeSwitch(surf, x, layer, mode, c1).x;
1183 __forceinline__ __device__ unsigned char surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1185 return __surfModeSwitch(surf, x, layer, mode, c1).x;
1189 __forceinline__ __device__ char1 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1191 return make_char1((signed char)__surfModeSwitch(surf, x, layer, mode, c1).x);
1195 __forceinline__ __device__ uchar1 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1197 return __surfModeSwitch(surf, x, layer, mode, c1);
1201 __forceinline__ __device__ char2 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1203 uchar2 tmp = __surfModeSwitch(surf, x, layer, mode, c2);
1205 return make_char2((signed char)tmp.x, (signed char)tmp.y);
1209 __forceinline__ __device__ uchar2 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1211 return __surfModeSwitch(surf, x, layer, mode, c2);
1215 __forceinline__ __device__ char4 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1217 uchar4 tmp = __surfModeSwitch(surf, x, layer, mode, c4);
1219 return make_char4((signed char)tmp.x, (signed char)tmp.y, (signed char)tmp.z, (signed char)tmp.w);
1223 __forceinline__ __device__ uchar4 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1225 return __surfModeSwitch(surf, x, layer, mode, c4);
1229 __forceinline__ __device__ short surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1231 return (short)__surfModeSwitch(surf, x, layer, mode, s1).x;
1235 __forceinline__ __device__ unsigned short surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1237 return __surfModeSwitch(surf, x, layer, mode, s1).x;
1241 __forceinline__ __device__ short1 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1243 return make_short1((signed short)__surfModeSwitch(surf, x, layer, mode, s1).x);
1247 __forceinline__ __device__ ushort1 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1249 return __surfModeSwitch(surf, x, layer, mode, s1);
1253 __forceinline__ __device__ short2 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1255 ushort2 tmp = __surfModeSwitch(surf, x, layer, mode, s2);
1257 return make_short2((signed short)tmp.x, (signed short)tmp.y);
1261 __forceinline__ __device__ ushort2 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1263 return __surfModeSwitch(surf, x, layer, mode, s2);
1267 __forceinline__ __device__ short4 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1269 ushort4 tmp = __surfModeSwitch(surf, x, layer, mode, s4);
1271 return make_short4((signed short)tmp.x, (signed short)tmp.y, (signed short)tmp.z, (signed short)tmp.w);
1275 __forceinline__ __device__ ushort4 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1277 return __surfModeSwitch(surf, x, layer, mode, s4);
1281 __forceinline__ __device__ int surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1283 return (int)__surfModeSwitch(surf, x, layer, mode, u1).x;
1287 __forceinline__ __device__ unsigned int surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1289 return __surfModeSwitch(surf, x, layer, mode, u1).x;
1293 __forceinline__ __device__ int1 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1295 return make_int1((signed int)__surfModeSwitch(surf, x, layer, mode, u1).x);
1299 __forceinline__ __device__ uint1 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1301 return __surfModeSwitch(surf, x, layer, mode, u1);
1305 __forceinline__ __device__ int2 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1307 uint2 tmp = __surfModeSwitch(surf, x, layer, mode, u2);
1309 return make_int2((int)tmp.x, (int)tmp.y);
1313 __forceinline__ __device__ uint2 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1315 return __surfModeSwitch(surf, x, layer, mode, u2);
1319 __forceinline__ __device__ int4 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1321 uint4 tmp = __surfModeSwitch(surf, x, layer, mode, u4);
1323 return make_int4((int)tmp.x, (int)tmp.y, (int)tmp.z, (int)tmp.w);
1327 __forceinline__ __device__ uint4 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1329 return __surfModeSwitch(surf, x, layer, mode, u4);
1333 __forceinline__ __device__ long long int surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1335 return (long long int)__surfModeSwitch(surf, x, layer, mode, l1).x;
1339 __forceinline__ __device__ unsigned long long int surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1341 return __surfModeSwitch(surf, x, layer, mode, l1).x;
1345 __forceinline__ __device__ longlong1 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1347 return make_longlong1((long long int)__surfModeSwitch(surf, x, layer, mode, l1).x);
1351 __forceinline__ __device__ ulonglong1 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1353 return __surfModeSwitch(surf, x, layer, mode, l1);
1357 __forceinline__ __device__ longlong2 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1359 ulonglong2 tmp = __surfModeSwitch(surf, x, layer, mode, l2);
1361 return make_longlong2((long long int)tmp.x, (long long int)tmp.y);
1365 __forceinline__ __device__ ulonglong2 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1367 return __surfModeSwitch(surf, x, layer, mode, l2);
1370 #if !defined(__LP64__)
1373 __forceinline__ __device__ long int surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1375 return (long int)__surfModeSwitch(surf, x, layer, mode, u1).x;
1379 __forceinline__ __device__ unsigned long int surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1381 return (unsigned long int)__surfModeSwitch(surf, x, layer, mode, u1).x;
1385 __forceinline__ __device__ long1 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1387 return make_long1((long int)__surfModeSwitch(surf, x, layer, mode, u1).x);
1391 __forceinline__ __device__ ulong1 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1393 return make_ulong1((unsigned long int)__surfModeSwitch(surf, x, layer, mode, u1).x);
1397 __forceinline__ __device__ long2 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1399 uint2 tmp = __surfModeSwitch(surf, x, layer, mode, u2);
1401 return make_long2((long int)tmp.x, (long int)tmp.y);
1405 __forceinline__ __device__ ulong2 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1407 uint2 tmp = __surfModeSwitch(surf, x, layer, mode, u2);
1409 return make_ulong2((unsigned long int)tmp.x, (unsigned long int)tmp.y);
1413 __forceinline__ __device__ long4 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1415 uint4 tmp = __surfModeSwitch(surf, x, layer, mode, u4);
1417 return make_long4((long int)tmp.x, (long int)tmp.y, (long int)tmp.z, (long int)tmp.w);
1421 __forceinline__ __device__ ulong4 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1423 uint4 tmp = __surfModeSwitch(surf, x, layer, mode, u4);
1425 return make_ulong4((unsigned long int)tmp.x, (unsigned long int)tmp.y, (unsigned long int)tmp.z, (unsigned long int)tmp.w);
1428 #endif /* !__LP64__ */
1431 __forceinline__ __device__ float surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1433 return __int_as_float((int)__surfModeSwitch(surf, x, layer, mode, u1).x);
1437 __forceinline__ __device__ float1 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1439 return make_float1(__int_as_float((int)__surfModeSwitch(surf, x, layer, mode, u1).x));
1443 __forceinline__ __device__ float2 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1445 uint2 tmp = __surfModeSwitch(surf, x, layer, mode, u2);
1447 return make_float2(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y));
1451 __forceinline__ __device__ float4 surf1DLayeredread(surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode)
1453 uint4 tmp = __surfModeSwitch(surf, x, layer, mode, u4);
1455 return make_float4(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y), __int_as_float((int)tmp.z), __int_as_float((int)tmp.w));
1458 #undef __surfModeSwitch
1460 /*******************************************************************************
1464 *******************************************************************************/
1465 extern __device__ __device_builtin__ uchar1 __surf2DLayeredreadc1(surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
1466 extern __device__ __device_builtin__ uchar2 __surf2DLayeredreadc2(surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
1467 extern __device__ __device_builtin__ uchar4 __surf2DLayeredreadc4(surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
1468 extern __device__ __device_builtin__ ushort1 __surf2DLayeredreads1(surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
1469 extern __device__ __device_builtin__ ushort2 __surf2DLayeredreads2(surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
1470 extern __device__ __device_builtin__ ushort4 __surf2DLayeredreads4(surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
1471 extern __device__ __device_builtin__ uint1 __surf2DLayeredreadu1(surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
1472 extern __device__ __device_builtin__ uint2 __surf2DLayeredreadu2(surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
1473 extern __device__ __device_builtin__ uint4 __surf2DLayeredreadu4(surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
1474 extern __device__ __device_builtin__ ulonglong1 __surf2DLayeredreadl1(surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
1475 extern __device__ __device_builtin__ ulonglong2 __surf2DLayeredreadl2(surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
1477 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
1479 #define __surfModeSwitch(surf, x, y, layer, mode, type) \
1480 ((mode == cudaBoundaryModeZero) ? __surf2DLayeredread##type(surf, x, y, layer, cudaBoundaryModeZero ) : \
1481 (mode == cudaBoundaryModeClamp) ? __surf2DLayeredread##type(surf, x, y, layer, cudaBoundaryModeClamp) : \
1482 __surf2DLayeredread##type(surf, x, y, layer, cudaBoundaryModeTrap ))
1484 #else /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
1486 #define __surfModeSwitch(surf, x, y, layer, mode, type) \
1487 __surf2DLayeredread##type(surf, x, y, layer, cudaBoundaryModeTrap)
1489 #endif /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
1492 static __forceinline__ __device__ void surf2DLayeredread(T *res, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
1494 (s == 1) ? (void)(*(uchar1 *)res = __surfModeSwitch(surf, x, y, layer, mode, c1)) :
1495 (s == 2) ? (void)(*(ushort1*)res = __surfModeSwitch(surf, x, y, layer, mode, s1)) :
1496 (s == 4) ? (void)(*(uint1 *)res = __surfModeSwitch(surf, x, y, layer, mode, u1)) :
1497 (s == 8) ? (void)(*(uint2 *)res = __surfModeSwitch(surf, x, y, layer, mode, u2)) :
1498 (s == 16) ? (void)(*(uint4 *)res = __surfModeSwitch(surf, x, y, layer, mode, u4)) :
1503 static __forceinline__ __device__ T surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
1507 surf2DLayeredread(&tmp, surf, x, y, layer, (int)sizeof(T), mode);
1513 static __forceinline__ __device__ void surf2DLayeredread(T *res, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
1515 *res = surf2DLayeredread<T>(surf, x, y, layer, mode);
1519 __forceinline__ __device__ char surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1521 return (char)__surfModeSwitch(surf, x, y, layer, mode, c1).x;
1525 __forceinline__ __device__ signed char surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1527 return (signed char)__surfModeSwitch(surf, x, y, layer, mode, c1).x;
1531 __forceinline__ __device__ unsigned char surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1533 return __surfModeSwitch(surf, x, y, layer, mode, c1).x;
1537 __forceinline__ __device__ char1 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1539 return make_char1((signed char)__surfModeSwitch(surf, x, y, layer, mode, c1).x);
1543 __forceinline__ __device__ uchar1 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1545 return __surfModeSwitch(surf, x, y, layer, mode, c1);
1549 __forceinline__ __device__ char2 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1551 uchar2 tmp = __surfModeSwitch(surf, x, y, layer, mode, c2);
1553 return make_char2((signed char)tmp.x, (signed char)tmp.y);
1557 __forceinline__ __device__ uchar2 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1559 return __surfModeSwitch(surf, x, y, layer, mode, c2);
1563 __forceinline__ __device__ char4 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1565 uchar4 tmp = __surfModeSwitch(surf, x, y, layer, mode, c4);
1567 return make_char4((signed char)tmp.x, (signed char)tmp.y, (signed char)tmp.z, (signed char)tmp.w);
1571 __forceinline__ __device__ uchar4 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1573 return __surfModeSwitch(surf, x, y, layer, mode, c4);
1577 __forceinline__ __device__ short surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1579 return (short)__surfModeSwitch(surf, x, y, layer, mode, s1).x;
1583 __forceinline__ __device__ unsigned short surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1585 return __surfModeSwitch(surf, x, y, layer, mode, s1).x;
1589 __forceinline__ __device__ short1 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1591 return make_short1((signed short)__surfModeSwitch(surf, x, y, layer, mode, s1).x);
1595 __forceinline__ __device__ ushort1 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1597 return __surfModeSwitch(surf, x, y, layer, mode, s1);
1601 __forceinline__ __device__ short2 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1603 ushort2 tmp = __surfModeSwitch(surf, x, y, layer, mode, s2);
1605 return make_short2((signed short)tmp.x, (signed short)tmp.y);
1609 __forceinline__ __device__ ushort2 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1611 return __surfModeSwitch(surf, x, y, layer, mode, s2);
1615 __forceinline__ __device__ short4 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1617 ushort4 tmp = __surfModeSwitch(surf, x, y, layer, mode, s4);
1619 return make_short4((signed short)tmp.x, (signed short)tmp.y, (signed short)tmp.z, (signed short)tmp.w);
1623 __forceinline__ __device__ ushort4 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1625 return __surfModeSwitch(surf, x, y, layer, mode, s4);
1629 __forceinline__ __device__ int surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1631 return (int)__surfModeSwitch(surf, x, y, layer, mode, u1).x;
1635 __forceinline__ __device__ unsigned int surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1637 return __surfModeSwitch(surf, x, y, layer, mode, u1).x;
1641 __forceinline__ __device__ int1 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1643 return make_int1((signed int)__surfModeSwitch(surf, x, y, layer, mode, u1).x);
1647 __forceinline__ __device__ uint1 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1649 return __surfModeSwitch(surf, x, y, layer, mode, u1);
1653 __forceinline__ __device__ int2 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1655 uint2 tmp = __surfModeSwitch(surf, x, y, layer, mode, u2);
1657 return make_int2((int)tmp.x, (int)tmp.y);
1661 __forceinline__ __device__ uint2 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1663 return __surfModeSwitch(surf, x, y, layer, mode, u2);
1667 __forceinline__ __device__ int4 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1669 uint4 tmp = __surfModeSwitch(surf, x, y, layer, mode, u4);
1671 return make_int4((int)tmp.x, (int)tmp.y, (int)tmp.z, (int)tmp.w);
1675 __forceinline__ __device__ uint4 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1677 return __surfModeSwitch(surf, x, y, layer, mode, u4);
1681 __forceinline__ __device__ long long int surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1683 return (long long int)__surfModeSwitch(surf, x, y, layer, mode, l1).x;
1687 __forceinline__ __device__ unsigned long long int surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1689 return __surfModeSwitch(surf, x, y, layer, mode, l1).x;
1693 __forceinline__ __device__ longlong1 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1695 return make_longlong1((long long int)__surfModeSwitch(surf, x, y, layer, mode, l1).x);
1699 __forceinline__ __device__ ulonglong1 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1701 return __surfModeSwitch(surf, x, y, layer, mode, l1);
1705 __forceinline__ __device__ longlong2 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1707 ulonglong2 tmp = __surfModeSwitch(surf, x, y, layer, mode, l2);
1709 return make_longlong2((long long int)tmp.x, (long long int)tmp.y);
1713 __forceinline__ __device__ ulonglong2 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1715 return __surfModeSwitch(surf, x, y, layer, mode, l2);
1718 #if !defined(__LP64__)
1721 __forceinline__ __device__ long int surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1723 return (long int)__surfModeSwitch(surf, x, y, layer, mode, u1).x;
1727 __forceinline__ __device__ unsigned long int surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1729 return (unsigned long int)__surfModeSwitch(surf, x, y, layer, mode, u1).x;
1733 __forceinline__ __device__ long1 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1735 return make_long1((long int)__surfModeSwitch(surf, x, y, layer, mode, u1).x);
1739 __forceinline__ __device__ ulong1 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1741 return make_ulong1((unsigned long int)__surfModeSwitch(surf, x, y, layer, mode, u1).x);
1745 __forceinline__ __device__ long2 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1747 uint2 tmp = __surfModeSwitch(surf, x, y, layer, mode, u2);
1749 return make_long2((long int)tmp.x, (long int)tmp.y);
1753 __forceinline__ __device__ ulong2 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1755 uint2 tmp = __surfModeSwitch(surf, x, y, layer, mode, u2);
1757 return make_ulong2((unsigned long int)tmp.x, (unsigned long int)tmp.y);
1761 __forceinline__ __device__ long4 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1763 uint4 tmp = __surfModeSwitch(surf, x, y, layer, mode, u4);
1765 return make_long4((long int)tmp.x, (long int)tmp.y, (long int)tmp.z, (long int)tmp.w);
1769 __forceinline__ __device__ ulong4 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1771 uint4 tmp = __surfModeSwitch(surf, x, y, layer, mode, u4);
1773 return make_ulong4((unsigned long int)tmp.x, (unsigned long int)tmp.y, (unsigned long int)tmp.z, (unsigned long int)tmp.w);
1776 #endif /* !__LP64__ */
1779 __forceinline__ __device__ float surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1781 return __int_as_float((int)__surfModeSwitch(surf, x, y, layer, mode, u1).x);
1785 __forceinline__ __device__ float1 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1787 return make_float1(__int_as_float((int)__surfModeSwitch(surf, x, y, layer, mode, u1).x));
1791 __forceinline__ __device__ float2 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1793 uint2 tmp = __surfModeSwitch(surf, x, y, layer, mode, u2);
1795 return make_float2(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y));
1799 __forceinline__ __device__ float4 surf2DLayeredread(surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode)
1801 uint4 tmp = __surfModeSwitch(surf, x, y, layer, mode, u4);
1803 return make_float4(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y), __int_as_float((int)tmp.z), __int_as_float((int)tmp.w));
1806 #undef __surfModeSwitch
1808 /*******************************************************************************
1812 *******************************************************************************/
1813 extern __device__ __device_builtin__ uchar1 __surfCubemapreadc1(surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
1814 extern __device__ __device_builtin__ uchar2 __surfCubemapreadc2(surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
1815 extern __device__ __device_builtin__ uchar4 __surfCubemapreadc4(surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
1816 extern __device__ __device_builtin__ ushort1 __surfCubemapreads1(surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
1817 extern __device__ __device_builtin__ ushort2 __surfCubemapreads2(surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
1818 extern __device__ __device_builtin__ ushort4 __surfCubemapreads4(surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
1819 extern __device__ __device_builtin__ uint1 __surfCubemapreadu1(surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
1820 extern __device__ __device_builtin__ uint2 __surfCubemapreadu2(surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
1821 extern __device__ __device_builtin__ uint4 __surfCubemapreadu4(surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
1822 extern __device__ __device_builtin__ ulonglong1 __surfCubemapreadl1(surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
1823 extern __device__ __device_builtin__ ulonglong2 __surfCubemapreadl2(surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
1825 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
1827 #define __surfModeSwitch(surf, x, y, face, mode, type) \
1828 ((mode == cudaBoundaryModeZero) ? __surfCubemapread##type(surf, x, y, face, cudaBoundaryModeZero ) : \
1829 (mode == cudaBoundaryModeClamp) ? __surfCubemapread##type(surf, x, y, face, cudaBoundaryModeClamp) : \
1830 __surfCubemapread##type(surf, x, y, face, cudaBoundaryModeTrap ))
1832 #else /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
1834 #define __surfModeSwitch(surf, x, y, face, mode, type) \
1835 __surfCubemapread##type(surf, x, y, face, cudaBoundaryModeTrap)
1837 #endif /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
1840 static __forceinline__ __device__ void surfCubemapread(T *res, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
1842 (s == 1) ? (void)(*(uchar1 *)res = __surfModeSwitch(surf, x, y, face, mode, c1)) :
1843 (s == 2) ? (void)(*(ushort1*)res = __surfModeSwitch(surf, x, y, face, mode, s1)) :
1844 (s == 4) ? (void)(*(uint1 *)res = __surfModeSwitch(surf, x, y, face, mode, u1)) :
1845 (s == 8) ? (void)(*(uint2 *)res = __surfModeSwitch(surf, x, y, face, mode, u2)) :
1846 (s == 16) ? (void)(*(uint4 *)res = __surfModeSwitch(surf, x, y, face, mode, u4)) :
1851 static __forceinline__ __device__ T surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
1855 surfCubemapread(&tmp, surf, x, y, face, (int)sizeof(T), mode);
1861 static __forceinline__ __device__ void surfCubemapread(T *res, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
1863 *res = surfCubemapread<T>(surf, x, y, face, mode);
1867 __forceinline__ __device__ char surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1869 return (char)__surfModeSwitch(surf, x, y, face, mode, c1).x;
1873 __forceinline__ __device__ signed char surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1875 return (signed char)__surfModeSwitch(surf, x, y, face, mode, c1).x;
1879 __forceinline__ __device__ unsigned char surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1881 return __surfModeSwitch(surf, x, y, face, mode, c1).x;
1885 __forceinline__ __device__ char1 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1887 return make_char1((signed char)__surfModeSwitch(surf, x, y, face, mode, c1).x);
1891 __forceinline__ __device__ uchar1 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1893 return __surfModeSwitch(surf, x, y, face, mode, c1);
1897 __forceinline__ __device__ char2 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1899 uchar2 tmp = __surfModeSwitch(surf, x, y, face, mode, c2);
1901 return make_char2((signed char)tmp.x, (signed char)tmp.y);
1905 __forceinline__ __device__ uchar2 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1907 return __surfModeSwitch(surf, x, y, face, mode, c2);
1911 __forceinline__ __device__ char4 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1913 uchar4 tmp = __surfModeSwitch(surf, x, y, face, mode, c4);
1915 return make_char4((signed char)tmp.x, (signed char)tmp.y, (signed char)tmp.z, (signed char)tmp.w);
1919 __forceinline__ __device__ uchar4 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1921 return __surfModeSwitch(surf, x, y, face, mode, c4);
1925 __forceinline__ __device__ short surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1927 return (short)__surfModeSwitch(surf, x, y, face, mode, s1).x;
1931 __forceinline__ __device__ unsigned short surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1933 return __surfModeSwitch(surf, x, y, face, mode, s1).x;
1937 __forceinline__ __device__ short1 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1939 return make_short1((signed short)__surfModeSwitch(surf, x, y, face, mode, s1).x);
1943 __forceinline__ __device__ ushort1 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1945 return __surfModeSwitch(surf, x, y, face, mode, s1);
1949 __forceinline__ __device__ short2 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1951 ushort2 tmp = __surfModeSwitch(surf, x, y, face, mode, s2);
1953 return make_short2((signed short)tmp.x, (signed short)tmp.y);
1957 __forceinline__ __device__ ushort2 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1959 return __surfModeSwitch(surf, x, y, face, mode, s2);
1963 __forceinline__ __device__ short4 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1965 ushort4 tmp = __surfModeSwitch(surf, x, y, face, mode, s4);
1967 return make_short4((signed short)tmp.x, (signed short)tmp.y, (signed short)tmp.z, (signed short)tmp.w);
1971 __forceinline__ __device__ ushort4 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1973 return __surfModeSwitch(surf, x, y, face, mode, s4);
1977 __forceinline__ __device__ int surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1979 return (int)__surfModeSwitch(surf, x, y, face, mode, u1).x;
1983 __forceinline__ __device__ unsigned int surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1985 return __surfModeSwitch(surf, x, y, face, mode, u1).x;
1989 __forceinline__ __device__ int1 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1991 return make_int1((signed int)__surfModeSwitch(surf, x, y, face, mode, u1).x);
1995 __forceinline__ __device__ uint1 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
1997 return __surfModeSwitch(surf, x, y, face, mode, u1);
2001 __forceinline__ __device__ int2 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2003 uint2 tmp = __surfModeSwitch(surf, x, y, face, mode, u2);
2005 return make_int2((int)tmp.x, (int)tmp.y);
2009 __forceinline__ __device__ uint2 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2011 return __surfModeSwitch(surf, x, y, face, mode, u2);
2015 __forceinline__ __device__ int4 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2017 uint4 tmp = __surfModeSwitch(surf, x, y, face, mode, u4);
2019 return make_int4((int)tmp.x, (int)tmp.y, (int)tmp.z, (int)tmp.w);
2023 __forceinline__ __device__ uint4 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2025 return __surfModeSwitch(surf, x, y, face, mode, u4);
2029 __forceinline__ __device__ long long int surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2031 return (long long int)__surfModeSwitch(surf, x, y, face, mode, l1).x;
2035 __forceinline__ __device__ unsigned long long int surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2037 return __surfModeSwitch(surf, x, y, face, mode, l1).x;
2041 __forceinline__ __device__ longlong1 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2043 return make_longlong1((long long int)__surfModeSwitch(surf, x, y, face, mode, l1).x);
2047 __forceinline__ __device__ ulonglong1 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2049 return __surfModeSwitch(surf, x, y, face, mode, l1);
2053 __forceinline__ __device__ longlong2 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2055 ulonglong2 tmp = __surfModeSwitch(surf, x, y, face, mode, l2);
2057 return make_longlong2((long long int)tmp.x, (long long int)tmp.y);
2061 __forceinline__ __device__ ulonglong2 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2063 return __surfModeSwitch(surf, x, y, face, mode, l2);
2066 #if !defined(__LP64__)
2069 __forceinline__ __device__ long int surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2071 return (long int)__surfModeSwitch(surf, x, y, face, mode, u1).x;
2075 __forceinline__ __device__ unsigned long int surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2077 return (unsigned long int)__surfModeSwitch(surf, x, y, face, mode, u1).x;
2081 __forceinline__ __device__ long1 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2083 return make_long1((long int)__surfModeSwitch(surf, x, y, face, mode, u1).x);
2087 __forceinline__ __device__ ulong1 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2089 return make_ulong1((unsigned long int)__surfModeSwitch(surf, x, y, face, mode, u1).x);
2093 __forceinline__ __device__ long2 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2095 uint2 tmp = __surfModeSwitch(surf, x, y, face, mode, u2);
2097 return make_long2((long int)tmp.x, (long int)tmp.y);
2101 __forceinline__ __device__ ulong2 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2103 uint2 tmp = __surfModeSwitch(surf, x, y, face, mode, u2);
2105 return make_ulong2((unsigned long int)tmp.x, (unsigned long int)tmp.y);
2109 __forceinline__ __device__ long4 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2111 uint4 tmp = __surfModeSwitch(surf, x, y, face, mode, u4);
2113 return make_long4((long int)tmp.x, (long int)tmp.y, (long int)tmp.z, (long int)tmp.w);
2117 __forceinline__ __device__ ulong4 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2119 uint4 tmp = __surfModeSwitch(surf, x, y, face, mode, u4);
2121 return make_ulong4((unsigned long int)tmp.x, (unsigned long int)tmp.y, (unsigned long int)tmp.z, (unsigned long int)tmp.w);
2124 #endif /* !__LP64__ */
2127 __forceinline__ __device__ float surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2129 return __int_as_float((int)__surfModeSwitch(surf, x, y, face, mode, u1).x);
2133 __forceinline__ __device__ float1 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2135 return make_float1(__int_as_float((int)__surfModeSwitch(surf, x, y, face, mode, u1).x));
2139 __forceinline__ __device__ float2 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2141 uint2 tmp = __surfModeSwitch(surf, x, y, face, mode, u2);
2143 return make_float2(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y));
2147 __forceinline__ __device__ float4 surfCubemapread(surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode)
2149 uint4 tmp = __surfModeSwitch(surf, x, y, face, mode, u4);
2151 return make_float4(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y), __int_as_float((int)tmp.z), __int_as_float((int)tmp.w));
2154 #undef __surfModeSwitch
2156 /*******************************************************************************
2160 *******************************************************************************/
2161 extern __device__ __device_builtin__ uchar1 __surfCubemapLayeredreadc1(surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
2162 extern __device__ __device_builtin__ uchar2 __surfCubemapLayeredreadc2(surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
2163 extern __device__ __device_builtin__ uchar4 __surfCubemapLayeredreadc4(surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
2164 extern __device__ __device_builtin__ ushort1 __surfCubemapLayeredreads1(surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
2165 extern __device__ __device_builtin__ ushort2 __surfCubemapLayeredreads2(surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
2166 extern __device__ __device_builtin__ ushort4 __surfCubemapLayeredreads4(surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
2167 extern __device__ __device_builtin__ uint1 __surfCubemapLayeredreadu1(surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
2168 extern __device__ __device_builtin__ uint2 __surfCubemapLayeredreadu2(surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
2169 extern __device__ __device_builtin__ uint4 __surfCubemapLayeredreadu4(surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
2170 extern __device__ __device_builtin__ ulonglong1 __surfCubemapLayeredreadl1(surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
2171 extern __device__ __device_builtin__ ulonglong2 __surfCubemapLayeredreadl2(surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
2173 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
2175 #define __surfModeSwitch(surf, x, y, layerFace, mode, type) \
2176 ((mode == cudaBoundaryModeZero) ? __surfCubemapLayeredread##type(surf, x, y, layerFace, cudaBoundaryModeZero ) : \
2177 (mode == cudaBoundaryModeClamp) ? __surfCubemapLayeredread##type(surf, x, y, layerFace, cudaBoundaryModeClamp) : \
2178 __surfCubemapLayeredread##type(surf, x, y, layerFace, cudaBoundaryModeTrap ))
2180 #else /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
2182 #define __surfModeSwitch(surf, x, y, layerFace, mode, type) \
2183 __surfCubemapLayeredread##type(surf, x, y, layerFace, cudaBoundaryModeTrap)
2186 #endif /* CUDA_ARCH && __CUDA_ARCH__ >= 200 */
2189 static __forceinline__ __device__ void surfCubemapLayeredread(T *res, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2191 (s == 1) ? (void)(*(uchar1 *)res = __surfModeSwitch(surf, x, y, layerFace, mode, c1)) :
2192 (s == 2) ? (void)(*(ushort1*)res = __surfModeSwitch(surf, x, y, layerFace, mode, s1)) :
2193 (s == 4) ? (void)(*(uint1 *)res = __surfModeSwitch(surf, x, y, layerFace, mode, u1)) :
2194 (s == 8) ? (void)(*(uint2 *)res = __surfModeSwitch(surf, x, y, layerFace, mode, u2)) :
2195 (s == 16) ? (void)(*(uint4 *)res = __surfModeSwitch(surf, x, y, layerFace, mode, u4)) :
2200 static __forceinline__ __device__ T surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2204 surfCubemapLayeredread(&tmp, surf, x, y, layerFace, (int)sizeof(T), mode);
2210 static __forceinline__ __device__ void surfCubemapLayeredread(T *res, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2212 *res = surfCubemapLayeredread<T>(surf, x, y, layerFace, mode);
2216 __forceinline__ __device__ char surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2218 return (char)__surfModeSwitch(surf, x, y, layerFace, mode, c1).x;
2222 __forceinline__ __device__ signed char surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2224 return (signed char)__surfModeSwitch(surf, x, y, layerFace, mode, c1).x;
2228 __forceinline__ __device__ unsigned char surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2230 return __surfModeSwitch(surf, x, y, layerFace, mode, c1).x;
2234 __forceinline__ __device__ char1 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2236 return make_char1((signed char)__surfModeSwitch(surf, x, y, layerFace, mode, c1).x);
2240 __forceinline__ __device__ uchar1 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2242 return __surfModeSwitch(surf, x, y, layerFace, mode, c1);
2246 __forceinline__ __device__ char2 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2248 uchar2 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, c2);
2250 return make_char2((signed char)tmp.x, (signed char)tmp.y);
2254 __forceinline__ __device__ uchar2 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2256 return __surfModeSwitch(surf, x, y, layerFace, mode, c2);
2260 __forceinline__ __device__ char4 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2262 uchar4 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, c4);
2264 return make_char4((signed char)tmp.x, (signed char)tmp.y, (signed char)tmp.z, (signed char)tmp.w);
2268 __forceinline__ __device__ uchar4 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2270 return __surfModeSwitch(surf, x, y, layerFace, mode, c4);
2274 __forceinline__ __device__ short surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2276 return (short)__surfModeSwitch(surf, x, y, layerFace, mode, s1).x;
2280 __forceinline__ __device__ unsigned short surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2282 return __surfModeSwitch(surf, x, y, layerFace, mode, s1).x;
2286 __forceinline__ __device__ short1 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2288 return make_short1((signed short)__surfModeSwitch(surf, x, y, layerFace, mode, s1).x);
2292 __forceinline__ __device__ ushort1 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2294 return __surfModeSwitch(surf, x, y, layerFace, mode, s1);
2298 __forceinline__ __device__ short2 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2300 ushort2 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, s2);
2302 return make_short2((signed short)tmp.x, (signed short)tmp.y);
2306 __forceinline__ __device__ ushort2 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2308 return __surfModeSwitch(surf, x, y, layerFace, mode, s2);
2312 __forceinline__ __device__ short4 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2314 ushort4 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, s4);
2316 return make_short4((signed short)tmp.x, (signed short)tmp.y, (signed short)tmp.z, (signed short)tmp.w);
2320 __forceinline__ __device__ ushort4 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2322 return __surfModeSwitch(surf, x, y, layerFace, mode, s4);
2326 __forceinline__ __device__ int surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2328 return (int)__surfModeSwitch(surf, x, y, layerFace, mode, u1).x;
2332 __forceinline__ __device__ unsigned int surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2334 return __surfModeSwitch(surf, x, y, layerFace, mode, u1).x;
2338 __forceinline__ __device__ int1 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2340 return make_int1((signed int)__surfModeSwitch(surf, x, y, layerFace, mode, u1).x);
2344 __forceinline__ __device__ uint1 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2346 return __surfModeSwitch(surf, x, y, layerFace, mode, u1);
2350 __forceinline__ __device__ int2 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2352 uint2 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, u2);
2354 return make_int2((int)tmp.x, (int)tmp.y);
2358 __forceinline__ __device__ uint2 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2360 return __surfModeSwitch(surf, x, y, layerFace, mode, u2);
2364 __forceinline__ __device__ int4 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2366 uint4 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, u4);
2368 return make_int4((int)tmp.x, (int)tmp.y, (int)tmp.z, (int)tmp.w);
2372 __forceinline__ __device__ uint4 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2374 return __surfModeSwitch(surf, x, y, layerFace, mode, u4);
2378 __forceinline__ __device__ long long int surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2380 return (long long int)__surfModeSwitch(surf, x, y, layerFace, mode, l1).x;
2384 __forceinline__ __device__ unsigned long long int surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2386 return __surfModeSwitch(surf, x, y, layerFace, mode, l1).x;
2390 __forceinline__ __device__ longlong1 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2392 return make_longlong1((long long int)__surfModeSwitch(surf, x, y, layerFace, mode, l1).x);
2396 __forceinline__ __device__ ulonglong1 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2398 return __surfModeSwitch(surf, x, y, layerFace, mode, l1);
2402 __forceinline__ __device__ longlong2 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2404 ulonglong2 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, l2);
2406 return make_longlong2((long long int)tmp.x, (long long int)tmp.y);
2410 __forceinline__ __device__ ulonglong2 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2412 return __surfModeSwitch(surf, x, y, layerFace, mode, l2);
2415 #if !defined(__LP64__)
2418 __forceinline__ __device__ long int surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2420 return (long int)__surfModeSwitch(surf, x, y, layerFace, mode, u1).x;
2424 __forceinline__ __device__ unsigned long int surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2426 return (unsigned long int)__surfModeSwitch(surf, x, y, layerFace, mode, u1).x;
2430 __forceinline__ __device__ long1 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2432 return make_long1((long int)__surfModeSwitch(surf, x, y, layerFace, mode, u1).x);
2436 __forceinline__ __device__ ulong1 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2438 return make_ulong1((unsigned long int)__surfModeSwitch(surf, x, y, layerFace, mode, u1).x);
2442 __forceinline__ __device__ long2 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2444 uint2 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, u2);
2446 return make_long2((long int)tmp.x, (long int)tmp.y);
2450 __forceinline__ __device__ ulong2 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2452 uint2 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, u2);
2454 return make_ulong2((unsigned long int)tmp.x, (unsigned long int)tmp.y);
2458 __forceinline__ __device__ long4 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2460 uint4 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, u4);
2462 return make_long4((long int)tmp.x, (long int)tmp.y, (long int)tmp.z, (long int)tmp.w);
2466 __forceinline__ __device__ ulong4 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2468 uint4 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, u4);
2470 return make_ulong4((unsigned long int)tmp.x, (unsigned long int)tmp.y, (unsigned long int)tmp.z, (unsigned long int)tmp.w);
2473 #endif /* !__LP64__ */
2476 __forceinline__ __device__ float surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2478 return __int_as_float((int)__surfModeSwitch(surf, x, y, layerFace, mode, u1).x);
2482 __forceinline__ __device__ float1 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2484 return make_float1(__int_as_float((int)__surfModeSwitch(surf, x, y, layerFace, mode, u1).x));
2488 __forceinline__ __device__ float2 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2490 uint2 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, u2);
2492 return make_float2(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y));
2496 __forceinline__ __device__ float4 surfCubemapLayeredread(surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode)
2498 uint4 tmp = __surfModeSwitch(surf, x, y, layerFace, mode, u4);
2500 return make_float4(__int_as_float((int)tmp.x), __int_as_float((int)tmp.y), __int_as_float((int)tmp.z), __int_as_float((int)tmp.w));
2503 #undef __surfModeSwitch
2505 /*******************************************************************************
2509 *******************************************************************************/
2511 extern __device__ __device_builtin__ void __surf1Dwritec1( uchar1 val, surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
2512 extern __device__ __device_builtin__ void __surf1Dwritec2( uchar2 val, surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
2513 extern __device__ __device_builtin__ void __surf1Dwritec4( uchar4 val, surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
2514 extern __device__ __device_builtin__ void __surf1Dwrites1( ushort1 val, surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
2515 extern __device__ __device_builtin__ void __surf1Dwrites2( ushort2 val, surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
2516 extern __device__ __device_builtin__ void __surf1Dwrites4( ushort4 val, surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
2517 extern __device__ __device_builtin__ void __surf1Dwriteu1( uint1 val, surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
2518 extern __device__ __device_builtin__ void __surf1Dwriteu2( uint2 val, surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
2519 extern __device__ __device_builtin__ void __surf1Dwriteu4( uint4 val, surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
2520 extern __device__ __device_builtin__ void __surf1Dwritel1(ulonglong1 val, surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
2521 extern __device__ __device_builtin__ void __surf1Dwritel2(ulonglong2 val, surface<void, cudaSurfaceType1D> t, int x, enum cudaSurfaceBoundaryMode mode);
2523 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
2525 #define __surfModeSwitch(val, surf, x, mode, type) \
2526 ((mode == cudaBoundaryModeZero) ? __surf1Dwrite##type(val, surf, x, cudaBoundaryModeZero ) : \
2527 (mode == cudaBoundaryModeClamp) ? __surf1Dwrite##type(val, surf, x, cudaBoundaryModeClamp) : \
2528 __surf1Dwrite##type(val, surf, x, cudaBoundaryModeTrap ))
2530 #else /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
2532 #define __surfModeSwitch(val, surf, x, mode, type) \
2533 __surf1Dwrite##type(val, surf, x, cudaBoundaryModeTrap)
2535 #endif /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
2538 static __forceinline__ __device__ void surf1Dwrite(T val, surface<void, cudaSurfaceType1D> surf, int x, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2551 (s == 1) ? (void)(__surfModeSwitch(tmp.c1, surf, x, mode, c1)) :
2552 (s == 2) ? (void)(__surfModeSwitch(tmp.s1, surf, x, mode, s1)) :
2553 (s == 4) ? (void)(__surfModeSwitch(tmp.u1, surf, x, mode, u1)) :
2554 (s == 8) ? (void)(__surfModeSwitch(tmp.u2, surf, x, mode, u2)) :
2555 (s == 16) ? (void)(__surfModeSwitch(tmp.u4, surf, x, mode, u4)) :
2560 static __forceinline__ __device__ void surf1Dwrite(T val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2562 surf1Dwrite(val, surf, x, (int)sizeof(T), mode);
2566 static __forceinline__ __device__ void surf1Dwrite(char val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2568 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, mode, c1);
2571 static __forceinline__ __device__ void surf1Dwrite(signed char val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2573 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, mode, c1);
2576 static __forceinline__ __device__ void surf1Dwrite(unsigned char val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2578 __surfModeSwitch(make_uchar1(val), surf, x, mode, c1);
2581 static __forceinline__ __device__ void surf1Dwrite(char1 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2583 __surfModeSwitch(make_uchar1((unsigned char)val.x), surf, x, mode, c1);
2586 static __forceinline__ __device__ void surf1Dwrite(uchar1 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2588 __surfModeSwitch(val, surf, x, mode, c1);
2591 static __forceinline__ __device__ void surf1Dwrite(char2 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2593 __surfModeSwitch(make_uchar2((unsigned char)val.x, (unsigned char)val.y), surf, x, mode, c2);
2596 static __forceinline__ __device__ void surf1Dwrite(uchar2 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2598 __surfModeSwitch(val, surf, x, mode, c2);
2601 static __forceinline__ __device__ void surf1Dwrite(char4 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2603 __surfModeSwitch(make_uchar4((unsigned char)val.x, (unsigned char)val.y, (unsigned char)val.z, (unsigned char)val.w), surf, x, mode, c4);
2606 static __forceinline__ __device__ void surf1Dwrite(uchar4 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2608 __surfModeSwitch(val, surf, x, mode, c4);
2611 static __forceinline__ __device__ void surf1Dwrite(short val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2613 __surfModeSwitch(make_ushort1((unsigned short)val), surf, x, mode, s1);
2616 static __forceinline__ __device__ void surf1Dwrite(unsigned short val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2618 __surfModeSwitch(make_ushort1(val), surf, x, mode, s1);
2621 static __forceinline__ __device__ void surf1Dwrite(short1 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2623 __surfModeSwitch(make_ushort1((unsigned short)val.x), surf, x, mode, s1);
2626 static __forceinline__ __device__ void surf1Dwrite(ushort1 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2628 __surfModeSwitch(val, surf, x, mode, s1);
2631 static __forceinline__ __device__ void surf1Dwrite(short2 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2633 __surfModeSwitch(make_ushort2((unsigned short)val.x, (unsigned short)val.y), surf, x, mode, s2);
2636 static __forceinline__ __device__ void surf1Dwrite(ushort2 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2638 __surfModeSwitch(val, surf, x, mode, s2);
2641 static __forceinline__ __device__ void surf1Dwrite(short4 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2643 __surfModeSwitch(make_ushort4((unsigned short)val.x, (unsigned short)val.y, (unsigned short)val.z, (unsigned short)val.w), surf, x, mode, s4);
2646 static __forceinline__ __device__ void surf1Dwrite(ushort4 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2648 __surfModeSwitch(val, surf, x, mode, s4);
2651 static __forceinline__ __device__ void surf1Dwrite(int val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2653 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, mode, u1);
2656 static __forceinline__ __device__ void surf1Dwrite(unsigned int val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2658 __surfModeSwitch(make_uint1(val), surf, x, mode, u1);
2661 static __forceinline__ __device__ void surf1Dwrite(int1 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2663 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, mode, u1);
2666 static __forceinline__ __device__ void surf1Dwrite(uint1 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2668 __surfModeSwitch(val, surf, x, mode, u1);
2671 static __forceinline__ __device__ void surf1Dwrite(int2 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2673 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, mode, u2);
2676 static __forceinline__ __device__ void surf1Dwrite(uint2 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2678 __surfModeSwitch(val, surf, x, mode, u2);
2681 static __forceinline__ __device__ void surf1Dwrite(int4 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2683 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, mode, u4);
2686 static __forceinline__ __device__ void surf1Dwrite(uint4 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2688 __surfModeSwitch(val, surf, x, mode, u4);
2691 static __forceinline__ __device__ void surf1Dwrite(long long int val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2693 __surfModeSwitch(make_ulonglong1((unsigned long long int)val), surf, x, mode, l1);
2696 static __forceinline__ __device__ void surf1Dwrite(unsigned long long int val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2698 __surfModeSwitch(make_ulonglong1(val), surf, x, mode, l1);
2701 static __forceinline__ __device__ void surf1Dwrite(longlong1 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2703 __surfModeSwitch(make_ulonglong1((unsigned long long int)val.x), surf, x, mode, l1);
2706 static __forceinline__ __device__ void surf1Dwrite(ulonglong1 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2708 __surfModeSwitch(val, surf, x, mode, l1);
2711 static __forceinline__ __device__ void surf1Dwrite(longlong2 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2713 __surfModeSwitch(make_ulonglong2((unsigned long long int)val.x, (unsigned long long int)val.y), surf, x, mode, l2);
2716 static __forceinline__ __device__ void surf1Dwrite(ulonglong2 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2718 __surfModeSwitch(val, surf, x, mode, l2);
2721 #if !defined(__LP64__)
2723 static __forceinline__ __device__ void surf1Dwrite(long int val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2725 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, mode, u1);
2728 static __forceinline__ __device__ void surf1Dwrite(unsigned long int val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2730 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, mode, u1);
2733 static __forceinline__ __device__ void surf1Dwrite(long1 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2735 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, mode, u1);
2738 static __forceinline__ __device__ void surf1Dwrite(ulong1 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2740 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, mode, u1);
2743 static __forceinline__ __device__ void surf1Dwrite(long2 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2745 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, mode, u2);
2748 static __forceinline__ __device__ void surf1Dwrite(ulong2 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2750 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, mode, u2);
2753 static __forceinline__ __device__ void surf1Dwrite(long4 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2755 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, mode, u4);
2758 static __forceinline__ __device__ void surf1Dwrite(ulong4 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2760 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, mode, u4);
2763 #endif /* !__LP64__ */
2765 static __forceinline__ __device__ void surf1Dwrite(float val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2767 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val)), surf, x, mode, u1);
2770 static __forceinline__ __device__ void surf1Dwrite(float1 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2772 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val.x)), surf, x, mode, u1);
2775 static __forceinline__ __device__ void surf1Dwrite(float2 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2777 __surfModeSwitch(make_uint2((unsigned int)__float_as_int(val.x), __float_as_int((unsigned int)val.y)), surf, x, mode, u2);
2780 static __forceinline__ __device__ void surf1Dwrite(float4 val, surface<void, cudaSurfaceType1D> surf, int x, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2782 __surfModeSwitch(make_uint4((unsigned int)__float_as_int(val.x), (unsigned int)__float_as_int(val.y), (unsigned int)__float_as_int(val.z), (unsigned int)__float_as_int(val.w)), surf, x, mode, u4);
2785 #undef __surfModeSwitch
2787 /*******************************************************************************
2791 *******************************************************************************/
2793 extern __device__ __device_builtin__ void __surf2Dwritec1( uchar1 val, surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
2794 extern __device__ __device_builtin__ void __surf2Dwritec2( uchar2 val, surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
2795 extern __device__ __device_builtin__ void __surf2Dwritec4( uchar4 val, surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
2796 extern __device__ __device_builtin__ void __surf2Dwrites1( ushort1 val, surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
2797 extern __device__ __device_builtin__ void __surf2Dwrites2( ushort2 val, surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
2798 extern __device__ __device_builtin__ void __surf2Dwrites4( ushort4 val, surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
2799 extern __device__ __device_builtin__ void __surf2Dwriteu1( uint1 val, surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
2800 extern __device__ __device_builtin__ void __surf2Dwriteu2( uint2 val, surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
2801 extern __device__ __device_builtin__ void __surf2Dwriteu4( uint4 val, surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
2802 extern __device__ __device_builtin__ void __surf2Dwritel1(ulonglong1 val, surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
2803 extern __device__ __device_builtin__ void __surf2Dwritel2(ulonglong2 val, surface<void, cudaSurfaceType2D> t, int x, int y, enum cudaSurfaceBoundaryMode mode);
2805 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
2807 #define __surfModeSwitch(val, surf, x, y, mode, type) \
2808 ((mode == cudaBoundaryModeZero) ? __surf2Dwrite##type(val, surf, x, y, cudaBoundaryModeZero ) : \
2809 (mode == cudaBoundaryModeClamp) ? __surf2Dwrite##type(val, surf, x, y, cudaBoundaryModeClamp) : \
2810 __surf2Dwrite##type(val, surf, x, y, cudaBoundaryModeTrap ))
2812 #else /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
2814 #define __surfModeSwitch(val, surf, x, y, mode, type) \
2815 __surf2Dwrite##type(val, surf, x, y, cudaBoundaryModeTrap)
2817 #endif /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
2820 static __forceinline__ __device__ void surf2Dwrite(T val, surface<void, cudaSurfaceType2D> surf, int x, int y, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2833 (s == 1) ? (void)(__surfModeSwitch(tmp.c1, surf, x, y, mode, c1)) :
2834 (s == 2) ? (void)(__surfModeSwitch(tmp.s1, surf, x, y, mode, s1)) :
2835 (s == 4) ? (void)(__surfModeSwitch(tmp.u1, surf, x, y, mode, u1)) :
2836 (s == 8) ? (void)(__surfModeSwitch(tmp.u2, surf, x, y, mode, u2)) :
2837 (s == 16) ? (void)(__surfModeSwitch(tmp.u4, surf, x, y, mode, u4)) :
2842 static __forceinline__ __device__ void surf2Dwrite(T val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2844 surf2Dwrite(val, surf, x, y, (int)sizeof(T), mode);
2848 static __forceinline__ __device__ void surf2Dwrite(char val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2850 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, y, mode, c1);
2853 static __forceinline__ __device__ void surf2Dwrite(signed char val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2855 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, y, mode, c1);
2858 static __forceinline__ __device__ void surf2Dwrite(unsigned char val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2860 __surfModeSwitch(make_uchar1(val), surf, x, y, mode, c1);
2863 static __forceinline__ __device__ void surf2Dwrite(char1 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2865 __surfModeSwitch(make_uchar1((unsigned char)val.x), surf, x, y, mode, c1);
2868 static __forceinline__ __device__ void surf2Dwrite(uchar1 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2870 __surfModeSwitch(val, surf, x, y, mode, c1);
2873 static __forceinline__ __device__ void surf2Dwrite(char2 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2875 __surfModeSwitch(make_uchar2((unsigned char)val.x, (unsigned char)val.y), surf, x, y, mode, c2);
2878 static __forceinline__ __device__ void surf2Dwrite(uchar2 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2880 __surfModeSwitch(val, surf, x, y, mode, c2);
2883 static __forceinline__ __device__ void surf2Dwrite(char4 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2885 __surfModeSwitch(make_uchar4((unsigned char)val.x, (unsigned char)val.y, (unsigned char)val.z, (unsigned char)val.w), surf, x, y, mode, c4);
2888 static __forceinline__ __device__ void surf2Dwrite(uchar4 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2890 __surfModeSwitch(val, surf, x, y, mode, c4);
2893 static __forceinline__ __device__ void surf2Dwrite(short val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2895 __surfModeSwitch(make_ushort1((unsigned short)val), surf, x, y, mode, s1);
2898 static __forceinline__ __device__ void surf2Dwrite(unsigned short val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2900 __surfModeSwitch(make_ushort1(val), surf, x, y, mode, s1);
2903 static __forceinline__ __device__ void surf2Dwrite(short1 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2905 __surfModeSwitch(make_ushort1((unsigned short)val.x), surf, x, y, mode, s1);
2908 static __forceinline__ __device__ void surf2Dwrite(ushort1 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2910 __surfModeSwitch(val, surf, x, y, mode, s1);
2913 static __forceinline__ __device__ void surf2Dwrite(short2 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2915 __surfModeSwitch(make_ushort2((unsigned short)val.x, (unsigned short)val.y), surf, x, y, mode, s2);
2918 static __forceinline__ __device__ void surf2Dwrite(ushort2 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2920 __surfModeSwitch(val, surf, x, y, mode, s2);
2923 static __forceinline__ __device__ void surf2Dwrite(short4 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2925 __surfModeSwitch(make_ushort4((unsigned short)val.x, (unsigned short)val.y, (unsigned short)val.z, (unsigned short)val.w), surf, x, y, mode, s4);
2928 static __forceinline__ __device__ void surf2Dwrite(ushort4 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2930 __surfModeSwitch(val, surf, x, y, mode, s4);
2933 static __forceinline__ __device__ void surf2Dwrite(int val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2935 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, mode, u1);
2938 static __forceinline__ __device__ void surf2Dwrite(unsigned int val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2940 __surfModeSwitch(make_uint1(val), surf, x, y, mode, u1);
2943 static __forceinline__ __device__ void surf2Dwrite(int1 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2945 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, mode, u1);
2948 static __forceinline__ __device__ void surf2Dwrite(uint1 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2950 __surfModeSwitch(val, surf, x, y, mode, u1);
2953 static __forceinline__ __device__ void surf2Dwrite(int2 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2955 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, mode, u2);
2958 static __forceinline__ __device__ void surf2Dwrite(uint2 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2960 __surfModeSwitch(val, surf, x, y, mode, u2);
2963 static __forceinline__ __device__ void surf2Dwrite(int4 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2965 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, mode, u4);
2968 static __forceinline__ __device__ void surf2Dwrite(uint4 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2970 __surfModeSwitch(val, surf, x, y, mode, u4);
2973 static __forceinline__ __device__ void surf2Dwrite(long long int val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2975 __surfModeSwitch(make_ulonglong1((unsigned long long int)val), surf, x, y, mode, l1);
2978 static __forceinline__ __device__ void surf2Dwrite(unsigned long long int val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2980 __surfModeSwitch(make_ulonglong1(val), surf, x, y, mode, l1);
2983 static __forceinline__ __device__ void surf2Dwrite(longlong1 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2985 __surfModeSwitch(make_ulonglong1((unsigned long long int)val.x), surf, x, y, mode, l1);
2988 static __forceinline__ __device__ void surf2Dwrite(ulonglong1 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2990 __surfModeSwitch(val, surf, x, y, mode, l1);
2993 static __forceinline__ __device__ void surf2Dwrite(longlong2 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
2995 __surfModeSwitch(make_ulonglong2((unsigned long long int)val.x, (unsigned long long int)val.y), surf, x, y, mode, l2);
2998 static __forceinline__ __device__ void surf2Dwrite(ulonglong2 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3000 __surfModeSwitch(val, surf, x, y, mode, l2);
3003 #if !defined(__LP64__)
3005 static __forceinline__ __device__ void surf2Dwrite(long int val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3007 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, mode, u1);
3010 static __forceinline__ __device__ void surf2Dwrite(unsigned long int val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3012 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, mode, u1);
3015 static __forceinline__ __device__ void surf2Dwrite(long1 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3017 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, mode, u1);
3020 static __forceinline__ __device__ void surf2Dwrite(ulong1 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3022 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, mode, u1);
3025 static __forceinline__ __device__ void surf2Dwrite(long2 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3027 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, mode, u2);
3030 static __forceinline__ __device__ void surf2Dwrite(ulong2 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3032 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, mode, u2);
3035 static __forceinline__ __device__ void surf2Dwrite(long4 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3037 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, mode, u4);
3040 static __forceinline__ __device__ void surf2Dwrite(ulong4 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3042 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, mode, u4);
3045 #endif /* !__LP64__ */
3047 static __forceinline__ __device__ void surf2Dwrite(float val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3049 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val)), surf, x, y, mode, u1);
3052 static __forceinline__ __device__ void surf2Dwrite(float1 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3054 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val.x)), surf, x, y, mode, u1);
3057 static __forceinline__ __device__ void surf2Dwrite(float2 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3059 __surfModeSwitch(make_uint2((unsigned int)__float_as_int(val.x), __float_as_int((unsigned int)val.y)), surf, x, y, mode, u2);
3062 static __forceinline__ __device__ void surf2Dwrite(float4 val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3064 __surfModeSwitch(make_uint4((unsigned int)__float_as_int(val.x), (unsigned int)__float_as_int(val.y), (unsigned int)__float_as_int(val.z), (unsigned int)__float_as_int(val.w)), surf, x, y, mode, u4);
3067 #undef __surfModeSwitch
3069 /*******************************************************************************
3073 *******************************************************************************/
3075 extern __device__ __device_builtin__ void __surf3Dwritec1( uchar1 val, surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
3076 extern __device__ __device_builtin__ void __surf3Dwritec2( uchar2 val, surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
3077 extern __device__ __device_builtin__ void __surf3Dwritec4( uchar4 val, surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
3078 extern __device__ __device_builtin__ void __surf3Dwrites1( ushort1 val, surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
3079 extern __device__ __device_builtin__ void __surf3Dwrites2( ushort2 val, surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
3080 extern __device__ __device_builtin__ void __surf3Dwrites4( ushort4 val, surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
3081 extern __device__ __device_builtin__ void __surf3Dwriteu1( uint1 val, surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
3082 extern __device__ __device_builtin__ void __surf3Dwriteu2( uint2 val, surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
3083 extern __device__ __device_builtin__ void __surf3Dwriteu4( uint4 val, surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
3084 extern __device__ __device_builtin__ void __surf3Dwritel1(ulonglong1 val, surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
3085 extern __device__ __device_builtin__ void __surf3Dwritel2(ulonglong2 val, surface<void, cudaSurfaceType3D> t, int x, int y, int z, enum cudaSurfaceBoundaryMode mode);
3087 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
3089 #define __surfModeSwitch(val, surf, x, y, z, mode, type) \
3090 ((mode == cudaBoundaryModeZero) ? __surf3Dwrite##type(val, surf, x, y, z, cudaBoundaryModeZero ) : \
3091 (mode == cudaBoundaryModeClamp) ? __surf3Dwrite##type(val, surf, x, y, z, cudaBoundaryModeClamp) : \
3092 __surf3Dwrite##type(val, surf, x, y, z, cudaBoundaryModeTrap ))
3094 #else /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
3096 #define __surfModeSwitch(val, surf, x, y, z, mode, type) \
3097 __surf3Dwrite##type(val, surf, x, y, z, cudaBoundaryModeTrap)
3099 #endif /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
3102 static __forceinline__ __device__ void surf3Dwrite(T val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3115 (s == 1) ? (void)(__surfModeSwitch(tmp.c1, surf, x, y, z, mode, c1)) :
3116 (s == 2) ? (void)(__surfModeSwitch(tmp.s1, surf, x, y, z, mode, s1)) :
3117 (s == 4) ? (void)(__surfModeSwitch(tmp.u1, surf, x, y, z, mode, u1)) :
3118 (s == 8) ? (void)(__surfModeSwitch(tmp.u2, surf, x, y, z, mode, u2)) :
3119 (s == 16) ? (void)(__surfModeSwitch(tmp.u4, surf, x, y, z, mode, u4)) :
3124 static __forceinline__ __device__ void surf3Dwrite(T val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3126 surf3Dwrite(val, surf, x, y, z, (int)sizeof(T), mode);
3130 static __forceinline__ __device__ void surf3Dwrite(char val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3132 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, y, z, mode, c1);
3135 static __forceinline__ __device__ void surf3Dwrite(signed char val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3137 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, y, z, mode, c1);
3140 static __forceinline__ __device__ void surf3Dwrite(unsigned char val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3142 __surfModeSwitch(make_uchar1(val), surf, x, y, z, mode, c1);
3145 static __forceinline__ __device__ void surf3Dwrite(char1 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3147 __surfModeSwitch(make_uchar1((unsigned char)val.x), surf, x, y, z, mode, c1);
3150 static __forceinline__ __device__ void surf3Dwrite(uchar1 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3152 __surfModeSwitch(val, surf, x, y, z, mode, c1);
3155 static __forceinline__ __device__ void surf3Dwrite(char2 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3157 __surfModeSwitch(make_uchar2((unsigned char)val.x, (unsigned char)val.y), surf, x, y, z, mode, c2);
3160 static __forceinline__ __device__ void surf3Dwrite(uchar2 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3162 __surfModeSwitch(val, surf, x, y, z, mode, c2);
3165 static __forceinline__ __device__ void surf3Dwrite(char4 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3167 __surfModeSwitch(make_uchar4((unsigned char)val.x, (unsigned char)val.y, (unsigned char)val.z, (unsigned char)val.w), surf, x, y, z, mode, c4);
3170 static __forceinline__ __device__ void surf3Dwrite(uchar4 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3172 __surfModeSwitch(val, surf, x, y, z, mode, c4);
3175 static __forceinline__ __device__ void surf3Dwrite(short val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3177 __surfModeSwitch(make_ushort1((unsigned short)val), surf, x, y, z, mode, s1);
3180 static __forceinline__ __device__ void surf3Dwrite(unsigned short val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3182 __surfModeSwitch(make_ushort1(val), surf, x, y, z, mode, s1);
3185 static __forceinline__ __device__ void surf3Dwrite(short1 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3187 __surfModeSwitch(make_ushort1((unsigned short)val.x), surf, x, y, z, mode, s1);
3190 static __forceinline__ __device__ void surf3Dwrite(ushort1 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3192 __surfModeSwitch(val, surf, x, y, z, mode, s1);
3195 static __forceinline__ __device__ void surf3Dwrite(short2 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3197 __surfModeSwitch(make_ushort2((unsigned short)val.x, (unsigned short)val.y), surf, x, y, z, mode, s2);
3200 static __forceinline__ __device__ void surf3Dwrite(ushort2 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3202 __surfModeSwitch(val, surf, x, y, z, mode, s2);
3205 static __forceinline__ __device__ void surf3Dwrite(short4 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3207 __surfModeSwitch(make_ushort4((unsigned short)val.x, (unsigned short)val.y, (unsigned short)val.z, (unsigned short)val.w), surf, x, y, z, mode, s4);
3210 static __forceinline__ __device__ void surf3Dwrite(ushort4 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3212 __surfModeSwitch(val, surf, x, y, z, mode, s4);
3215 static __forceinline__ __device__ void surf3Dwrite(int val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3217 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, z, mode, u1);
3220 static __forceinline__ __device__ void surf3Dwrite(unsigned int val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3222 __surfModeSwitch(make_uint1(val), surf, x, y, z, mode, u1);
3225 static __forceinline__ __device__ void surf3Dwrite(int1 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3227 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, z, mode, u1);
3230 static __forceinline__ __device__ void surf3Dwrite(uint1 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3232 __surfModeSwitch(val, surf, x, y, z, mode, u1);
3235 static __forceinline__ __device__ void surf3Dwrite(int2 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3237 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, z, mode, u2);
3240 static __forceinline__ __device__ void surf3Dwrite(uint2 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3242 __surfModeSwitch(val, surf, x, y, z, mode, u2);
3245 static __forceinline__ __device__ void surf3Dwrite(int4 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3247 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, z, mode, u4);
3250 static __forceinline__ __device__ void surf3Dwrite(uint4 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3252 __surfModeSwitch(val, surf, x, y, z, mode, u4);
3255 static __forceinline__ __device__ void surf3Dwrite(long long int val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3257 __surfModeSwitch(make_ulonglong1((unsigned long long int)val), surf, x, y, z, mode, l1);
3260 static __forceinline__ __device__ void surf3Dwrite(unsigned long long int val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3262 __surfModeSwitch(make_ulonglong1(val), surf, x, y, z, mode, l1);
3265 static __forceinline__ __device__ void surf3Dwrite(longlong1 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3267 __surfModeSwitch(make_ulonglong1((unsigned long long int)val.x), surf, x, y, z, mode, l1);
3270 static __forceinline__ __device__ void surf3Dwrite(ulonglong1 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3272 __surfModeSwitch(val, surf, x, y, z, mode, l1);
3275 static __forceinline__ __device__ void surf3Dwrite(longlong2 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3277 __surfModeSwitch(make_ulonglong2((unsigned long long int)val.x, (unsigned long long int)val.y), surf, x, y, z, mode, l2);
3280 static __forceinline__ __device__ void surf3Dwrite(ulonglong2 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3282 __surfModeSwitch(val, surf, x, y, z, mode, l2);
3285 #if !defined(__LP64__)
3287 static __forceinline__ __device__ void surf3Dwrite(long int val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3289 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, z, mode, u1);
3292 static __forceinline__ __device__ void surf3Dwrite(unsigned long int val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3294 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, z, mode, u1);
3297 static __forceinline__ __device__ void surf3Dwrite(long1 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3299 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, z, mode, u1);
3302 static __forceinline__ __device__ void surf3Dwrite(ulong1 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3304 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, z, mode, u1);
3307 static __forceinline__ __device__ void surf3Dwrite(long2 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3309 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, z, mode, u2);
3312 static __forceinline__ __device__ void surf3Dwrite(ulong2 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3314 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, z, mode, u2);
3317 static __forceinline__ __device__ void surf3Dwrite(long4 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3319 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, z, mode, u4);
3322 static __forceinline__ __device__ void surf3Dwrite(ulong4 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3324 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, z, mode, u4);
3327 #endif /* !__LP64__ */
3329 static __forceinline__ __device__ void surf3Dwrite(float val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3331 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val)), surf, x, y, z, mode, u1);
3334 static __forceinline__ __device__ void surf3Dwrite(float1 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3336 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val.x)), surf, x, y, z, mode, u1);
3339 static __forceinline__ __device__ void surf3Dwrite(float2 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3341 __surfModeSwitch(make_uint2((unsigned int)__float_as_int(val.x), __float_as_int((unsigned int)val.y)), surf, x, y, z, mode, u2);
3344 static __forceinline__ __device__ void surf3Dwrite(float4 val, surface<void, cudaSurfaceType3D> surf, int x, int y, int z, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3346 __surfModeSwitch(make_uint4((unsigned int)__float_as_int(val.x), (unsigned int)__float_as_int(val.y), (unsigned int)__float_as_int(val.z), (unsigned int)__float_as_int(val.w)), surf, x, y, z, mode, u4);
3349 #undef __surfModeSwitch
3351 /*******************************************************************************
3355 *******************************************************************************/
3357 extern __device__ __device_builtin__ void __surf1DLayeredwritec1( uchar1 val, surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
3358 extern __device__ __device_builtin__ void __surf1DLayeredwritec2( uchar2 val, surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
3359 extern __device__ __device_builtin__ void __surf1DLayeredwritec4( uchar4 val, surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
3360 extern __device__ __device_builtin__ void __surf1DLayeredwrites1( ushort1 val, surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
3361 extern __device__ __device_builtin__ void __surf1DLayeredwrites2( ushort2 val, surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
3362 extern __device__ __device_builtin__ void __surf1DLayeredwrites4( ushort4 val, surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
3363 extern __device__ __device_builtin__ void __surf1DLayeredwriteu1( uint1 val, surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
3364 extern __device__ __device_builtin__ void __surf1DLayeredwriteu2( uint2 val, surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
3365 extern __device__ __device_builtin__ void __surf1DLayeredwriteu4( uint4 val, surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
3366 extern __device__ __device_builtin__ void __surf1DLayeredwritel1(ulonglong1 val, surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
3367 extern __device__ __device_builtin__ void __surf1DLayeredwritel2(ulonglong2 val, surface<void, cudaSurfaceType1DLayered> t, int x, int layer, enum cudaSurfaceBoundaryMode mode);
3369 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
3371 #define __surfModeSwitch(val, surf, x, layer, mode, type) \
3372 ((mode == cudaBoundaryModeZero) ? __surf1DLayeredwrite##type(val, surf, x, layer, cudaBoundaryModeZero ) : \
3373 (mode == cudaBoundaryModeClamp) ? __surf1DLayeredwrite##type(val, surf, x, layer, cudaBoundaryModeClamp) : \
3374 __surf1DLayeredwrite##type(val, surf, x, layer, cudaBoundaryModeTrap ))
3376 #else /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
3378 #define __surfModeSwitch(val, surf, x, layer, mode, type) \
3379 __surf1DLayeredwrite##type(val, surf, x, layer, cudaBoundaryModeTrap)
3381 #endif /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
3384 static __forceinline__ __device__ void surf1DLayeredwrite(T val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3397 (s == 1) ? (void)(__surfModeSwitch(tmp.c1, surf, x, layer, mode, c1)) :
3398 (s == 2) ? (void)(__surfModeSwitch(tmp.s1, surf, x, layer, mode, s1)) :
3399 (s == 4) ? (void)(__surfModeSwitch(tmp.u1, surf, x, layer, mode, u1)) :
3400 (s == 8) ? (void)(__surfModeSwitch(tmp.u2, surf, x, layer, mode, u2)) :
3401 (s == 16) ? (void)(__surfModeSwitch(tmp.u4, surf, x, layer, mode, u4)) :
3406 static __forceinline__ __device__ void surf1DLayeredwrite(T val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3408 surf1DLayeredwrite(val, surf, x, layer, (int)sizeof(T), mode);
3412 static __forceinline__ __device__ void surf1DLayeredwrite(char val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3414 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, layer, mode, c1);
3417 static __forceinline__ __device__ void surf1DLayeredwrite(signed char val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3419 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, layer, mode, c1);
3422 static __forceinline__ __device__ void surf1DLayeredwrite(unsigned char val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3424 __surfModeSwitch(make_uchar1(val), surf, x, layer, mode, c1);
3427 static __forceinline__ __device__ void surf1DLayeredwrite(char1 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3429 __surfModeSwitch(make_uchar1((unsigned char)val.x), surf, x, layer, mode, c1);
3432 static __forceinline__ __device__ void surf1DLayeredwrite(uchar1 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3434 __surfModeSwitch(val, surf, x, layer, mode, c1);
3437 static __forceinline__ __device__ void surf1DLayeredwrite(char2 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3439 __surfModeSwitch(make_uchar2((unsigned char)val.x, (unsigned char)val.y), surf, x, layer, mode, c2);
3442 static __forceinline__ __device__ void surf1DLayeredwrite(uchar2 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3444 __surfModeSwitch(val, surf, x, layer, mode, c2);
3447 static __forceinline__ __device__ void surf1DLayeredwrite(char4 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3449 __surfModeSwitch(make_uchar4((unsigned char)val.x, (unsigned char)val.y, (unsigned char)val.z, (unsigned char)val.w), surf, x, layer, mode, c4);
3452 static __forceinline__ __device__ void surf1DLayeredwrite(uchar4 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3454 __surfModeSwitch(val, surf, x, layer, mode, c4);
3457 static __forceinline__ __device__ void surf1DLayeredwrite(short val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3459 __surfModeSwitch(make_ushort1((unsigned short)val), surf, x, layer, mode, s1);
3462 static __forceinline__ __device__ void surf1DLayeredwrite(unsigned short val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3464 __surfModeSwitch(make_ushort1(val), surf, x, layer, mode, s1);
3467 static __forceinline__ __device__ void surf1DLayeredwrite(short1 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3469 __surfModeSwitch(make_ushort1((unsigned short)val.x), surf, x, layer, mode, s1);
3472 static __forceinline__ __device__ void surf1DLayeredwrite(ushort1 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3474 __surfModeSwitch(val, surf, x, layer, mode, s1);
3477 static __forceinline__ __device__ void surf1DLayeredwrite(short2 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3479 __surfModeSwitch(make_ushort2((unsigned short)val.x, (unsigned short)val.y), surf, x, layer, mode, s2);
3482 static __forceinline__ __device__ void surf1DLayeredwrite(ushort2 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3484 __surfModeSwitch(val, surf, x, layer, mode, s2);
3487 static __forceinline__ __device__ void surf1DLayeredwrite(short4 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3489 __surfModeSwitch(make_ushort4((unsigned short)val.x, (unsigned short)val.y, (unsigned short)val.z, (unsigned short)val.w), surf, x, layer, mode, s4);
3492 static __forceinline__ __device__ void surf1DLayeredwrite(ushort4 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3494 __surfModeSwitch(val, surf, x, layer, mode, s4);
3497 static __forceinline__ __device__ void surf1DLayeredwrite(int val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3499 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, layer, mode, u1);
3502 static __forceinline__ __device__ void surf1DLayeredwrite(unsigned int val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3504 __surfModeSwitch(make_uint1(val), surf, x, layer, mode, u1);
3507 static __forceinline__ __device__ void surf1DLayeredwrite(int1 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3509 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, layer, mode, u1);
3512 static __forceinline__ __device__ void surf1DLayeredwrite(uint1 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3514 __surfModeSwitch(val, surf, x, layer, mode, u1);
3517 static __forceinline__ __device__ void surf1DLayeredwrite(int2 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3519 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, layer, mode, u2);
3522 static __forceinline__ __device__ void surf1DLayeredwrite(uint2 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3524 __surfModeSwitch(val, surf, x, layer, mode, u2);
3527 static __forceinline__ __device__ void surf1DLayeredwrite(int4 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3529 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, layer, mode, u4);
3532 static __forceinline__ __device__ void surf1DLayeredwrite(uint4 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3534 __surfModeSwitch(val, surf, x, layer, mode, u4);
3537 static __forceinline__ __device__ void surf1DLayeredwrite(long long int val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3539 __surfModeSwitch(make_ulonglong1((unsigned long long int)val), surf, x, layer, mode, l1);
3542 static __forceinline__ __device__ void surf1DLayeredwrite(unsigned long long int val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3544 __surfModeSwitch(make_ulonglong1(val), surf, x, layer, mode, l1);
3547 static __forceinline__ __device__ void surf1DLayeredwrite(longlong1 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3549 __surfModeSwitch(make_ulonglong1((unsigned long long int)val.x), surf, x, layer, mode, l1);
3552 static __forceinline__ __device__ void surf1DLayeredwrite(ulonglong1 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3554 __surfModeSwitch(val, surf, x, layer, mode, l1);
3557 static __forceinline__ __device__ void surf1DLayeredwrite(longlong2 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3559 __surfModeSwitch(make_ulonglong2((unsigned long long int)val.x, (unsigned long long int)val.y), surf, x, layer, mode, l2);
3562 static __forceinline__ __device__ void surf1DLayeredwrite(ulonglong2 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3564 __surfModeSwitch(val, surf, x, layer, mode, l2);
3567 #if !defined(__LP64__)
3569 static __forceinline__ __device__ void surf1DLayeredwrite(long int val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3571 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, layer, mode, u1);
3574 static __forceinline__ __device__ void surf1DLayeredwrite(unsigned long int val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3576 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, layer, mode, u1);
3579 static __forceinline__ __device__ void surf1DLayeredwrite(long1 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3581 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, layer, mode, u1);
3584 static __forceinline__ __device__ void surf1DLayeredwrite(ulong1 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3586 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, layer, mode, u1);
3589 static __forceinline__ __device__ void surf1DLayeredwrite(long2 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3591 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, layer, mode, u2);
3594 static __forceinline__ __device__ void surf1DLayeredwrite(ulong2 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3596 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, layer, mode, u2);
3599 static __forceinline__ __device__ void surf1DLayeredwrite(long4 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3601 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, layer, mode, u4);
3604 static __forceinline__ __device__ void surf1DLayeredwrite(ulong4 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3606 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, layer, mode, u4);
3609 #endif /* !__LP64__ */
3611 static __forceinline__ __device__ void surf1DLayeredwrite(float val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3613 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val)), surf, x, layer, mode, u1);
3616 static __forceinline__ __device__ void surf1DLayeredwrite(float1 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3618 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val.x)), surf, x, layer, mode, u1);
3621 static __forceinline__ __device__ void surf1DLayeredwrite(float2 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3623 __surfModeSwitch(make_uint2((unsigned int)__float_as_int(val.x), __float_as_int((unsigned int)val.y)), surf, x, layer, mode, u2);
3626 static __forceinline__ __device__ void surf1DLayeredwrite(float4 val, surface<void, cudaSurfaceType1DLayered> surf, int x, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3628 __surfModeSwitch(make_uint4((unsigned int)__float_as_int(val.x), (unsigned int)__float_as_int(val.y), (unsigned int)__float_as_int(val.z), (unsigned int)__float_as_int(val.w)), surf, x, layer, mode, u4);
3631 #undef __surfModeSwitch
3633 /*******************************************************************************
3637 *******************************************************************************/
3639 extern __device__ __device_builtin__ void __surf2DLayeredwritec1( uchar1 val, surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
3640 extern __device__ __device_builtin__ void __surf2DLayeredwritec2( uchar2 val, surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
3641 extern __device__ __device_builtin__ void __surf2DLayeredwritec4( uchar4 val, surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
3642 extern __device__ __device_builtin__ void __surf2DLayeredwrites1( ushort1 val, surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
3643 extern __device__ __device_builtin__ void __surf2DLayeredwrites2( ushort2 val, surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
3644 extern __device__ __device_builtin__ void __surf2DLayeredwrites4( ushort4 val, surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
3645 extern __device__ __device_builtin__ void __surf2DLayeredwriteu1( uint1 val, surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
3646 extern __device__ __device_builtin__ void __surf2DLayeredwriteu2( uint2 val, surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
3647 extern __device__ __device_builtin__ void __surf2DLayeredwriteu4( uint4 val, surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
3648 extern __device__ __device_builtin__ void __surf2DLayeredwritel1(ulonglong1 val, surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
3649 extern __device__ __device_builtin__ void __surf2DLayeredwritel2(ulonglong2 val, surface<void, cudaSurfaceType2DLayered> t, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode);
3651 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
3653 #define __surfModeSwitch(val, surf, x, y, layer, mode, type) \
3654 ((mode == cudaBoundaryModeZero) ? __surf2DLayeredwrite##type(val, surf, x, y, layer, cudaBoundaryModeZero ) : \
3655 (mode == cudaBoundaryModeClamp) ? __surf2DLayeredwrite##type(val, surf, x, y, layer, cudaBoundaryModeClamp) : \
3656 __surf2DLayeredwrite##type(val, surf, x, y, layer, cudaBoundaryModeTrap ))
3658 #else /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
3660 #define __surfModeSwitch(val, surf, x, y, layer, mode, type) \
3661 __surf2DLayeredwrite##type(val, surf, x, y, layer, cudaBoundaryModeTrap)
3663 #endif /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
3666 static __forceinline__ __device__ void surf2DLayeredwrite(T val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3679 (s == 1) ? (void)(__surfModeSwitch(tmp.c1, surf, x, y, layer, mode, c1)) :
3680 (s == 2) ? (void)(__surfModeSwitch(tmp.s1, surf, x, y, layer, mode, s1)) :
3681 (s == 4) ? (void)(__surfModeSwitch(tmp.u1, surf, x, y, layer, mode, u1)) :
3682 (s == 8) ? (void)(__surfModeSwitch(tmp.u2, surf, x, y, layer, mode, u2)) :
3683 (s == 16) ? (void)(__surfModeSwitch(tmp.u4, surf, x, y, layer, mode, u4)) :
3688 static __forceinline__ __device__ void surf2DLayeredwrite(T val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3690 surf2DLayeredwrite(val, surf, x, y, layer, (int)sizeof(T), mode);
3694 static __forceinline__ __device__ void surf2DLayeredwrite(char val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3696 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, y, layer, mode, c1);
3699 static __forceinline__ __device__ void surf2DLayeredwrite(signed char val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3701 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, y, layer, mode, c1);
3704 static __forceinline__ __device__ void surf2DLayeredwrite(unsigned char val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3706 __surfModeSwitch(make_uchar1(val), surf, x, y, layer, mode, c1);
3709 static __forceinline__ __device__ void surf2DLayeredwrite(char1 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3711 __surfModeSwitch(make_uchar1((unsigned char)val.x), surf, x, y, layer, mode, c1);
3714 static __forceinline__ __device__ void surf2DLayeredwrite(uchar1 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3716 __surfModeSwitch(val, surf, x, y, layer, mode, c1);
3719 static __forceinline__ __device__ void surf2DLayeredwrite(char2 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3721 __surfModeSwitch(make_uchar2((unsigned char)val.x, (unsigned char)val.y), surf, x, y, layer, mode, c2);
3724 static __forceinline__ __device__ void surf2DLayeredwrite(uchar2 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3726 __surfModeSwitch(val, surf, x, y, layer, mode, c2);
3729 static __forceinline__ __device__ void surf2DLayeredwrite(char4 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3731 __surfModeSwitch(make_uchar4((unsigned char)val.x, (unsigned char)val.y, (unsigned char)val.z, (unsigned char)val.w), surf, x, y, layer, mode, c4);
3734 static __forceinline__ __device__ void surf2DLayeredwrite(uchar4 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3736 __surfModeSwitch(val, surf, x, y, layer, mode, c4);
3739 static __forceinline__ __device__ void surf2DLayeredwrite(short val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3741 __surfModeSwitch(make_ushort1((unsigned short)val), surf, x, y, layer, mode, s1);
3744 static __forceinline__ __device__ void surf2DLayeredwrite(unsigned short val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3746 __surfModeSwitch(make_ushort1(val), surf, x, y, layer, mode, s1);
3749 static __forceinline__ __device__ void surf2DLayeredwrite(short1 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3751 __surfModeSwitch(make_ushort1((unsigned short)val.x), surf, x, y, layer, mode, s1);
3754 static __forceinline__ __device__ void surf2DLayeredwrite(ushort1 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3756 __surfModeSwitch(val, surf, x, y, layer, mode, s1);
3759 static __forceinline__ __device__ void surf2DLayeredwrite(short2 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3761 __surfModeSwitch(make_ushort2((unsigned short)val.x, (unsigned short)val.y), surf, x, y, layer, mode, s2);
3764 static __forceinline__ __device__ void surf2DLayeredwrite(ushort2 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3766 __surfModeSwitch(val, surf, x, y, layer, mode, s2);
3769 static __forceinline__ __device__ void surf2DLayeredwrite(short4 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3771 __surfModeSwitch(make_ushort4((unsigned short)val.x, (unsigned short)val.y, (unsigned short)val.z, (unsigned short)val.w), surf, x, y, layer, mode, s4);
3774 static __forceinline__ __device__ void surf2DLayeredwrite(ushort4 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3776 __surfModeSwitch(val, surf, x, y, layer, mode, s4);
3779 static __forceinline__ __device__ void surf2DLayeredwrite(int val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3781 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, layer, mode, u1);
3784 static __forceinline__ __device__ void surf2DLayeredwrite(unsigned int val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3786 __surfModeSwitch(make_uint1(val), surf, x, y, layer, mode, u1);
3789 static __forceinline__ __device__ void surf2DLayeredwrite(int1 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3791 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, layer, mode, u1);
3794 static __forceinline__ __device__ void surf2DLayeredwrite(uint1 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3796 __surfModeSwitch(val, surf, x, y, layer, mode, u1);
3799 static __forceinline__ __device__ void surf2DLayeredwrite(int2 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3801 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, layer, mode, u2);
3804 static __forceinline__ __device__ void surf2DLayeredwrite(uint2 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3806 __surfModeSwitch(val, surf, x, y, layer, mode, u2);
3809 static __forceinline__ __device__ void surf2DLayeredwrite(int4 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3811 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, layer, mode, u4);
3814 static __forceinline__ __device__ void surf2DLayeredwrite(uint4 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3816 __surfModeSwitch(val, surf, x, y, layer, mode, u4);
3819 static __forceinline__ __device__ void surf2DLayeredwrite(long long int val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3821 __surfModeSwitch(make_ulonglong1((unsigned long long int)val), surf, x, y, layer, mode, l1);
3824 static __forceinline__ __device__ void surf2DLayeredwrite(unsigned long long int val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3826 __surfModeSwitch(make_ulonglong1(val), surf, x, y, layer, mode, l1);
3829 static __forceinline__ __device__ void surf2DLayeredwrite(longlong1 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3831 __surfModeSwitch(make_ulonglong1((unsigned long long int)val.x), surf, x, y, layer, mode, l1);
3834 static __forceinline__ __device__ void surf2DLayeredwrite(ulonglong1 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3836 __surfModeSwitch(val, surf, x, y, layer, mode, l1);
3839 static __forceinline__ __device__ void surf2DLayeredwrite(longlong2 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3841 __surfModeSwitch(make_ulonglong2((unsigned long long int)val.x, (unsigned long long int)val.y), surf, x, y, layer, mode, l2);
3844 static __forceinline__ __device__ void surf2DLayeredwrite(ulonglong2 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3846 __surfModeSwitch(val, surf, x, y, layer, mode, l2);
3849 #if !defined(__LP64__)
3851 static __forceinline__ __device__ void surf2DLayeredwrite(long int val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3853 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, layer, mode, u1);
3856 static __forceinline__ __device__ void surf2DLayeredwrite(unsigned long int val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3858 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, layer, mode, u1);
3861 static __forceinline__ __device__ void surf2DLayeredwrite(long1 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3863 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, layer, mode, u1);
3866 static __forceinline__ __device__ void surf2DLayeredwrite(ulong1 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3868 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, layer, mode, u1);
3871 static __forceinline__ __device__ void surf2DLayeredwrite(long2 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3873 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, layer, mode, u2);
3876 static __forceinline__ __device__ void surf2DLayeredwrite(ulong2 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3878 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, layer, mode, u2);
3881 static __forceinline__ __device__ void surf2DLayeredwrite(long4 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3883 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, layer, mode, u4);
3886 static __forceinline__ __device__ void surf2DLayeredwrite(ulong4 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3888 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, layer, mode, u4);
3891 #endif /* !__LP64__ */
3893 static __forceinline__ __device__ void surf2DLayeredwrite(float val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3895 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val)), surf, x, y, layer, mode, u1);
3898 static __forceinline__ __device__ void surf2DLayeredwrite(float1 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3900 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val.x)), surf, x, y, layer, mode, u1);
3903 static __forceinline__ __device__ void surf2DLayeredwrite(float2 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3905 __surfModeSwitch(make_uint2((unsigned int)__float_as_int(val.x), __float_as_int((unsigned int)val.y)), surf, x, y, layer, mode, u2);
3908 static __forceinline__ __device__ void surf2DLayeredwrite(float4 val, surface<void, cudaSurfaceType2DLayered> surf, int x, int y, int layer, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3910 __surfModeSwitch(make_uint4((unsigned int)__float_as_int(val.x), (unsigned int)__float_as_int(val.y), (unsigned int)__float_as_int(val.z), (unsigned int)__float_as_int(val.w)), surf, x, y, layer, mode, u4);
3913 #undef __surfModeSwitch
3915 /*******************************************************************************
3919 *******************************************************************************/
3920 extern __device__ __device_builtin__ void __surfCubemapwritec1( uchar1 val, surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
3921 extern __device__ __device_builtin__ void __surfCubemapwritec2( uchar2 val, surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
3922 extern __device__ __device_builtin__ void __surfCubemapwritec4( uchar4 val, surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
3923 extern __device__ __device_builtin__ void __surfCubemapwrites1( ushort1 val, surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
3924 extern __device__ __device_builtin__ void __surfCubemapwrites2( ushort2 val, surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
3925 extern __device__ __device_builtin__ void __surfCubemapwrites4( ushort4 val, surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
3926 extern __device__ __device_builtin__ void __surfCubemapwriteu1( uint1 val, surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
3927 extern __device__ __device_builtin__ void __surfCubemapwriteu2( uint2 val, surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
3928 extern __device__ __device_builtin__ void __surfCubemapwriteu4( uint4 val, surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
3929 extern __device__ __device_builtin__ void __surfCubemapwritel1(ulonglong1 val, surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
3930 extern __device__ __device_builtin__ void __surfCubemapwritel2(ulonglong2 val, surface<void, cudaSurfaceTypeCubemap> t, int x, int y, int face, enum cudaSurfaceBoundaryMode mode);
3932 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
3934 #define __surfModeSwitch(val, surf, x, y, face, mode, type) \
3935 ((mode == cudaBoundaryModeZero) ? __surfCubemapwrite##type(val, surf, x, y, face, cudaBoundaryModeZero ) : \
3936 (mode == cudaBoundaryModeClamp) ? __surfCubemapwrite##type(val, surf, x, y, face, cudaBoundaryModeClamp) : \
3937 __surfCubemapwrite##type(val, surf, x, y, face, cudaBoundaryModeTrap ))
3939 #else /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
3941 #define __surfModeSwitch(val, surf, x, y, face, mode, type) \
3942 __surfCubemapwrite##type(val, surf, x, y, face, cudaBoundaryModeTrap)
3945 #endif /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
3948 static __forceinline__ __device__ void surfCubemapwrite(T val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3961 (s == 1) ? (void)(__surfModeSwitch(tmp.c1, surf, x, y, face, mode, c1)) :
3962 (s == 2) ? (void)(__surfModeSwitch(tmp.s1, surf, x, y, face, mode, s1)) :
3963 (s == 4) ? (void)(__surfModeSwitch(tmp.u1, surf, x, y, face, mode, u1)) :
3964 (s == 8) ? (void)(__surfModeSwitch(tmp.u2, surf, x, y, face, mode, u2)) :
3965 (s == 16) ? (void)(__surfModeSwitch(tmp.u4, surf, x, y, face, mode, u4)) :
3970 static __forceinline__ __device__ void surfCubemapwrite(T val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3972 surfCubemapwrite(val, surf, x, y, face, (int)sizeof(T), mode);
3976 static __forceinline__ __device__ void surfCubemapwrite(char val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3978 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, y, face, mode, c1);
3981 static __forceinline__ __device__ void surfCubemapwrite(signed char val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3983 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, y, face, mode, c1);
3986 static __forceinline__ __device__ void surfCubemapwrite(unsigned char val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3988 __surfModeSwitch(make_uchar1(val), surf, x, y, face, mode, c1);
3991 static __forceinline__ __device__ void surfCubemapwrite(char1 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3993 __surfModeSwitch(make_uchar1((unsigned char)val.x), surf, x, y, face, mode, c1);
3996 static __forceinline__ __device__ void surfCubemapwrite(uchar1 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
3998 __surfModeSwitch(val, surf, x, y, face, mode, c1);
4001 static __forceinline__ __device__ void surfCubemapwrite(char2 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4003 __surfModeSwitch(make_uchar2((unsigned char)val.x, (unsigned char)val.y), surf, x, y, face, mode, c2);
4006 static __forceinline__ __device__ void surfCubemapwrite(uchar2 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4008 __surfModeSwitch(val, surf, x, y, face, mode, c2);
4011 static __forceinline__ __device__ void surfCubemapwrite(char4 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4013 __surfModeSwitch(make_uchar4((unsigned char)val.x, (unsigned char)val.y, (unsigned char)val.z, (unsigned char)val.w), surf, x, y, face, mode, c4);
4016 static __forceinline__ __device__ void surfCubemapwrite(uchar4 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4018 __surfModeSwitch(val, surf, x, y, face, mode, c4);
4021 static __forceinline__ __device__ void surfCubemapwrite(short val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4023 __surfModeSwitch(make_ushort1((unsigned short)val), surf, x, y, face, mode, s1);
4026 static __forceinline__ __device__ void surfCubemapwrite(unsigned short val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4028 __surfModeSwitch(make_ushort1(val), surf, x, y, face, mode, s1);
4031 static __forceinline__ __device__ void surfCubemapwrite(short1 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4033 __surfModeSwitch(make_ushort1((unsigned short)val.x), surf, x, y, face, mode, s1);
4036 static __forceinline__ __device__ void surfCubemapwrite(ushort1 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4038 __surfModeSwitch(val, surf, x, y, face, mode, s1);
4041 static __forceinline__ __device__ void surfCubemapwrite(short2 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4043 __surfModeSwitch(make_ushort2((unsigned short)val.x, (unsigned short)val.y), surf, x, y, face, mode, s2);
4046 static __forceinline__ __device__ void surfCubemapwrite(ushort2 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4048 __surfModeSwitch(val, surf, x, y, face, mode, s2);
4051 static __forceinline__ __device__ void surfCubemapwrite(short4 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4053 __surfModeSwitch(make_ushort4((unsigned short)val.x, (unsigned short)val.y, (unsigned short)val.z, (unsigned short)val.w), surf, x, y, face, mode, s4);
4056 static __forceinline__ __device__ void surfCubemapwrite(ushort4 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4058 __surfModeSwitch(val, surf, x, y, face, mode, s4);
4061 static __forceinline__ __device__ void surfCubemapwrite(int val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4063 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, face, mode, u1);
4066 static __forceinline__ __device__ void surfCubemapwrite(unsigned int val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4068 __surfModeSwitch(make_uint1(val), surf, x, y, face, mode, u1);
4071 static __forceinline__ __device__ void surfCubemapwrite(int1 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4073 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, face, mode, u1);
4076 static __forceinline__ __device__ void surfCubemapwrite(uint1 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4078 __surfModeSwitch(val, surf, x, y, face, mode, u1);
4081 static __forceinline__ __device__ void surfCubemapwrite(int2 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4083 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, face, mode, u2);
4086 static __forceinline__ __device__ void surfCubemapwrite(uint2 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4088 __surfModeSwitch(val, surf, x, y, face, mode, u2);
4091 static __forceinline__ __device__ void surfCubemapwrite(int4 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4093 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, face, mode, u4);
4096 static __forceinline__ __device__ void surfCubemapwrite(uint4 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4098 __surfModeSwitch(val, surf, x, y, face, mode, u4);
4101 static __forceinline__ __device__ void surfCubemapwrite(long long int val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4103 __surfModeSwitch(make_ulonglong1((unsigned long long int)val), surf, x, y, face, mode, l1);
4106 static __forceinline__ __device__ void surfCubemapwrite(unsigned long long int val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4108 __surfModeSwitch(make_ulonglong1(val), surf, x, y, face, mode, l1);
4111 static __forceinline__ __device__ void surfCubemapwrite(longlong1 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4113 __surfModeSwitch(make_ulonglong1((unsigned long long int)val.x), surf, x, y, face, mode, l1);
4116 static __forceinline__ __device__ void surfCubemapwrite(ulonglong1 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4118 __surfModeSwitch(val, surf, x, y, face, mode, l1);
4121 static __forceinline__ __device__ void surfCubemapwrite(longlong2 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4123 __surfModeSwitch(make_ulonglong2((unsigned long long int)val.x, (unsigned long long int)val.y), surf, x, y, face, mode, l2);
4126 static __forceinline__ __device__ void surfCubemapwrite(ulonglong2 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4128 __surfModeSwitch(val, surf, x, y, face, mode, l2);
4131 #if !defined(__LP64__)
4133 static __forceinline__ __device__ void surfCubemapwrite(long int val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4135 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, face, mode, u1);
4138 static __forceinline__ __device__ void surfCubemapwrite(unsigned long int val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4140 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, face, mode, u1);
4143 static __forceinline__ __device__ void surfCubemapwrite(long1 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4145 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, face, mode, u1);
4148 static __forceinline__ __device__ void surfCubemapwrite(ulong1 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4150 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, face, mode, u1);
4153 static __forceinline__ __device__ void surfCubemapwrite(long2 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4155 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, face, mode, u2);
4158 static __forceinline__ __device__ void surfCubemapwrite(ulong2 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4160 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, face, mode, u2);
4163 static __forceinline__ __device__ void surfCubemapwrite(long4 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4165 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, face, mode, u4);
4168 static __forceinline__ __device__ void surfCubemapwrite(ulong4 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4170 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, face, mode, u4);
4173 #endif /* !__LP64__ */
4175 static __forceinline__ __device__ void surfCubemapwrite(float val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4177 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val)), surf, x, y, face, mode, u1);
4180 static __forceinline__ __device__ void surfCubemapwrite(float1 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4182 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val.x)), surf, x, y, face, mode, u1);
4185 static __forceinline__ __device__ void surfCubemapwrite(float2 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4187 __surfModeSwitch(make_uint2((unsigned int)__float_as_int(val.x), __float_as_int((unsigned int)val.y)), surf, x, y, face, mode, u2);
4190 static __forceinline__ __device__ void surfCubemapwrite(float4 val, surface<void, cudaSurfaceTypeCubemap> surf, int x, int y, int face, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4192 __surfModeSwitch(make_uint4((unsigned int)__float_as_int(val.x), (unsigned int)__float_as_int(val.y), (unsigned int)__float_as_int(val.z), (unsigned int)__float_as_int(val.w)), surf, x, y, face, mode, u4);
4195 #undef __surfModeSwitch
4197 /*******************************************************************************
4201 *******************************************************************************/
4202 extern __device__ __device_builtin__ void __surfCubemapLayeredwritec1( uchar1 val, surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
4203 extern __device__ __device_builtin__ void __surfCubemapLayeredwritec2( uchar2 val, surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
4204 extern __device__ __device_builtin__ void __surfCubemapLayeredwritec4( uchar4 val, surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
4205 extern __device__ __device_builtin__ void __surfCubemapLayeredwrites1( ushort1 val, surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
4206 extern __device__ __device_builtin__ void __surfCubemapLayeredwrites2( ushort2 val, surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
4207 extern __device__ __device_builtin__ void __surfCubemapLayeredwrites4( ushort4 val, surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
4208 extern __device__ __device_builtin__ void __surfCubemapLayeredwriteu1( uint1 val, surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
4209 extern __device__ __device_builtin__ void __surfCubemapLayeredwriteu2( uint2 val, surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
4210 extern __device__ __device_builtin__ void __surfCubemapLayeredwriteu4( uint4 val, surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
4211 extern __device__ __device_builtin__ void __surfCubemapLayeredwritel1(ulonglong1 val, surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
4212 extern __device__ __device_builtin__ void __surfCubemapLayeredwritel2(ulonglong2 val, surface<void, cudaSurfaceTypeCubemapLayered> t, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode);
4214 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 200
4216 #define __surfModeSwitch(val, surf, x, y, layerFace, mode, type) \
4217 ((mode == cudaBoundaryModeZero) ? __surfCubemapLayeredwrite##type(val, surf, x, y, layerFace, cudaBoundaryModeZero ) : \
4218 (mode == cudaBoundaryModeClamp) ? __surfCubemapLayeredwrite##type(val, surf, x, y, layerFace, cudaBoundaryModeClamp) : \
4219 __surfCubemapLayeredwrite##type(val, surf, x, y, layerFace, cudaBoundaryModeTrap ))
4221 #else /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
4223 #define __surfModeSwitch(val, surf, x, y, layerFace, mode, type) \
4224 __surfCubemapLayeredwrite##type(val, surf, x, y, layerFace, cudaBoundaryModeTrap)
4227 #endif /* __CUDA_ARCH__ && __CUDA_ARCH__ >= 200 */
4230 static __forceinline__ __device__ void surfCubemapLayeredwrite(T val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, int s, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4243 (s == 1) ? (void)(__surfModeSwitch(tmp.c1, surf, x, y, layerFace, mode, c1)) :
4244 (s == 2) ? (void)(__surfModeSwitch(tmp.s1, surf, x, y, layerFace, mode, s1)) :
4245 (s == 4) ? (void)(__surfModeSwitch(tmp.u1, surf, x, y, layerFace, mode, u1)) :
4246 (s == 8) ? (void)(__surfModeSwitch(tmp.u2, surf, x, y, layerFace, mode, u2)) :
4247 (s == 16) ? (void)(__surfModeSwitch(tmp.u4, surf, x, y, layerFace, mode, u4)) :
4252 static __forceinline__ __device__ void surfCubemapLayeredwrite(T val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4254 surfCubemapLayeredwrite(val, surf, x, y, layerFace, (int)sizeof(T), mode);
4258 static __forceinline__ __device__ void surfCubemapLayeredwrite(char val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4260 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, y, layerFace, mode, c1);
4263 static __forceinline__ __device__ void surfCubemapLayeredwrite(signed char val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4265 __surfModeSwitch(make_uchar1((unsigned char)val), surf, x, y, layerFace, mode, c1);
4268 static __forceinline__ __device__ void surfCubemapLayeredwrite(unsigned char val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4270 __surfModeSwitch(make_uchar1(val), surf, x, y, layerFace, mode, c1);
4273 static __forceinline__ __device__ void surfCubemapLayeredwrite(char1 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4275 __surfModeSwitch(make_uchar1((unsigned char)val.x), surf, x, y, layerFace, mode, c1);
4278 static __forceinline__ __device__ void surfCubemapLayeredwrite(uchar1 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4280 __surfModeSwitch(val, surf, x, y, layerFace, mode, c1);
4283 static __forceinline__ __device__ void surfCubemapLayeredwrite(char2 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4285 __surfModeSwitch(make_uchar2((unsigned char)val.x, (unsigned char)val.y), surf, x, y, layerFace, mode, c2);
4288 static __forceinline__ __device__ void surfCubemapLayeredwrite(uchar2 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4290 __surfModeSwitch(val, surf, x, y, layerFace, mode, c2);
4293 static __forceinline__ __device__ void surfCubemapLayeredwrite(char4 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4295 __surfModeSwitch(make_uchar4((unsigned char)val.x, (unsigned char)val.y, (unsigned char)val.z, (unsigned char)val.w), surf, x, y, layerFace, mode, c4);
4298 static __forceinline__ __device__ void surfCubemapLayeredwrite(uchar4 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4300 __surfModeSwitch(val, surf, x, y, layerFace, mode, c4);
4303 static __forceinline__ __device__ void surfCubemapLayeredwrite(short val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4305 __surfModeSwitch(make_ushort1((unsigned short)val), surf, x, y, layerFace, mode, s1);
4308 static __forceinline__ __device__ void surfCubemapLayeredwrite(unsigned short val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4310 __surfModeSwitch(make_ushort1(val), surf, x, y, layerFace, mode, s1);
4313 static __forceinline__ __device__ void surfCubemapLayeredwrite(short1 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4315 __surfModeSwitch(make_ushort1((unsigned short)val.x), surf, x, y, layerFace, mode, s1);
4318 static __forceinline__ __device__ void surfCubemapLayeredwrite(ushort1 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4320 __surfModeSwitch(val, surf, x, y, layerFace, mode, s1);
4323 static __forceinline__ __device__ void surfCubemapLayeredwrite(short2 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4325 __surfModeSwitch(make_ushort2((unsigned short)val.x, (unsigned short)val.y), surf, x, y, layerFace, mode, s2);
4328 static __forceinline__ __device__ void surfCubemapLayeredwrite(ushort2 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4330 __surfModeSwitch(val, surf, x, y, layerFace, mode, s2);
4333 static __forceinline__ __device__ void surfCubemapLayeredwrite(short4 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4335 __surfModeSwitch(make_ushort4((unsigned short)val.x, (unsigned short)val.y, (unsigned short)val.z, (unsigned short)val.w), surf, x, y, layerFace, mode, s4);
4338 static __forceinline__ __device__ void surfCubemapLayeredwrite(ushort4 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4340 __surfModeSwitch(val, surf, x, y, layerFace, mode, s4);
4343 static __forceinline__ __device__ void surfCubemapLayeredwrite(int val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4345 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, layerFace, mode, u1);
4348 static __forceinline__ __device__ void surfCubemapLayeredwrite(unsigned int val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4350 __surfModeSwitch(make_uint1(val), surf, x, y, layerFace, mode, u1);
4353 static __forceinline__ __device__ void surfCubemapLayeredwrite(int1 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4355 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, layerFace, mode, u1);
4358 static __forceinline__ __device__ void surfCubemapLayeredwrite(uint1 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4360 __surfModeSwitch(val, surf, x, y, layerFace, mode, u1);
4363 static __forceinline__ __device__ void surfCubemapLayeredwrite(int2 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4365 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, layerFace, mode, u2);
4368 static __forceinline__ __device__ void surfCubemapLayeredwrite(uint2 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4370 __surfModeSwitch(val, surf, x, y, layerFace, mode, u2);
4373 static __forceinline__ __device__ void surfCubemapLayeredwrite(int4 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4375 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, layerFace, mode, u4);
4378 static __forceinline__ __device__ void surfCubemapLayeredwrite(uint4 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4380 __surfModeSwitch(val, surf, x, y, layerFace, mode, u4);
4383 static __forceinline__ __device__ void surfCubemapLayeredwrite(long long int val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4385 __surfModeSwitch(make_ulonglong1((unsigned long long int)val), surf, x, y, layerFace, mode, l1);
4388 static __forceinline__ __device__ void surfCubemapLayeredwrite(unsigned long long int val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4390 __surfModeSwitch(make_ulonglong1(val), surf, x, y, layerFace, mode, l1);
4393 static __forceinline__ __device__ void surfCubemapLayeredwrite(longlong1 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4395 __surfModeSwitch(make_ulonglong1((unsigned long long int)val.x), surf, x, y, layerFace, mode, l1);
4398 static __forceinline__ __device__ void surfCubemapLayeredwrite(ulonglong1 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4400 __surfModeSwitch(val, surf, x, y, layerFace, mode, l1);
4403 static __forceinline__ __device__ void surfCubemapLayeredwrite(longlong2 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4405 __surfModeSwitch(make_ulonglong2((unsigned long long int)val.x, (unsigned long long int)val.y), surf, x, y, layerFace, mode, l2);
4408 static __forceinline__ __device__ void surfCubemapLayeredwrite(ulonglong2 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4410 __surfModeSwitch(val, surf, x, y, layerFace, mode, l2);
4413 #if !defined(__LP64__)
4415 static __forceinline__ __device__ void surfCubemapLayeredwrite(long int val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4417 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, layerFace, mode, u1);
4420 static __forceinline__ __device__ void surfCubemapLayeredwrite(unsigned long int val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4422 __surfModeSwitch(make_uint1((unsigned int)val), surf, x, y, layerFace, mode, u1);
4425 static __forceinline__ __device__ void surfCubemapLayeredwrite(long1 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4427 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, layerFace, mode, u1);
4430 static __forceinline__ __device__ void surfCubemapLayeredwrite(ulong1 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4432 __surfModeSwitch(make_uint1((unsigned int)val.x), surf, x, y, layerFace, mode, u1);
4435 static __forceinline__ __device__ void surfCubemapLayeredwrite(long2 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4437 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, layerFace, mode, u2);
4440 static __forceinline__ __device__ void surfCubemapLayeredwrite(ulong2 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4442 __surfModeSwitch(make_uint2((unsigned int)val.x, (unsigned int)val.y), surf, x, y, layerFace, mode, u2);
4445 static __forceinline__ __device__ void surfCubemapLayeredwrite(long4 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4447 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, layerFace, mode, u4);
4450 static __forceinline__ __device__ void surfCubemapLayeredwrite(ulong4 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4452 __surfModeSwitch(make_uint4((unsigned int)val.x, (unsigned int)val.y, (unsigned int)val.z, (unsigned int)val.w), surf, x, y, layerFace, mode, u4);
4455 #endif /* !__LP64__ */
4457 static __forceinline__ __device__ void surfCubemapLayeredwrite(float val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4459 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val)), surf, x, y, layerFace, mode, u1);
4462 static __forceinline__ __device__ void surfCubemapLayeredwrite(float1 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4464 __surfModeSwitch(make_uint1((unsigned int)__float_as_int(val.x)), surf, x, y, layerFace, mode, u1);
4467 static __forceinline__ __device__ void surfCubemapLayeredwrite(float2 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4469 __surfModeSwitch(make_uint2((unsigned int)__float_as_int(val.x), __float_as_int((unsigned int)val.y)), surf, x, y, layerFace, mode, u2);
4472 static __forceinline__ __device__ void surfCubemapLayeredwrite(float4 val, surface<void, cudaSurfaceTypeCubemapLayered> surf, int x, int y, int layerFace, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
4474 __surfModeSwitch(make_uint4((unsigned int)__float_as_int(val.x), (unsigned int)__float_as_int(val.y), (unsigned int)__float_as_int(val.z), (unsigned int)__float_as_int(val.w)), surf, x, y, layerFace, mode, u4);
4477 #undef __surfModeSwitch
4479 /*******************************************************************************
4483 *******************************************************************************/
4485 #elif defined(__CUDABE__)
4487 #if defined(__CUDANVVM__)
4488 extern uchar1 __surf1Dreadc1(unsigned long long, int, enum cudaSurfaceBoundaryMode);
4489 extern uchar2 __surf1Dreadc2(unsigned long long, int, enum cudaSurfaceBoundaryMode);
4490 extern uchar4 __surf1Dreadc4(unsigned long long, int, enum cudaSurfaceBoundaryMode);
4491 extern ushort1 __surf1Dreads1(unsigned long long, int, enum cudaSurfaceBoundaryMode);
4492 extern ushort2 __surf1Dreads2(unsigned long long, int, enum cudaSurfaceBoundaryMode);
4493 extern ushort4 __surf1Dreads4(unsigned long long, int, enum cudaSurfaceBoundaryMode);
4494 extern uint1 __surf1Dreadu1(unsigned long long, int, enum cudaSurfaceBoundaryMode);
4495 extern uint2 __surf1Dreadu2(unsigned long long, int, enum cudaSurfaceBoundaryMode);
4496 extern uint4 __surf1Dreadu4(unsigned long long, int, enum cudaSurfaceBoundaryMode);
4497 extern ulonglong1 __surf1Dreadl1(unsigned long long, int, enum cudaSurfaceBoundaryMode);
4498 extern ulonglong2 __surf1Dreadl2(unsigned long long, int, enum cudaSurfaceBoundaryMode);
4499 extern uchar1 __surf2Dreadc1(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4500 extern uchar2 __surf2Dreadc2(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4501 extern uchar4 __surf2Dreadc4(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4502 extern ushort1 __surf2Dreads1(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4503 extern ushort2 __surf2Dreads2(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4504 extern ushort4 __surf2Dreads4(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4505 extern uint1 __surf2Dreadu1(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4506 extern uint2 __surf2Dreadu2(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4507 extern uint4 __surf2Dreadu4(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4508 extern ulonglong1 __surf2Dreadl1(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4509 extern ulonglong2 __surf2Dreadl2(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4510 extern uchar1 __surf3Dreadc1(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4511 extern uchar2 __surf3Dreadc2(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4512 extern uchar4 __surf3Dreadc4(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4513 extern ushort1 __surf3Dreads1(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4514 extern ushort2 __surf3Dreads2(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4515 extern ushort4 __surf3Dreads4(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4516 extern uint1 __surf3Dreadu1(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4517 extern uint2 __surf3Dreadu2(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4518 extern uint4 __surf3Dreadu4(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4519 extern ulonglong1 __surf3Dreadl1(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4520 extern ulonglong2 __surf3Dreadl2(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4521 extern uchar1 __surf1DLayeredreadc1(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4522 extern uchar2 __surf1DLayeredreadc2(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4523 extern uchar4 __surf1DLayeredreadc4(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4524 extern ushort1 __surf1DLayeredreads1(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4525 extern ushort2 __surf1DLayeredreads2(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4526 extern ushort4 __surf1DLayeredreads4(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4527 extern uint1 __surf1DLayeredreadu1(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4528 extern uint2 __surf1DLayeredreadu2(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4529 extern uint4 __surf1DLayeredreadu4(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4530 extern ulonglong1 __surf1DLayeredreadl1(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4531 extern ulonglong2 __surf1DLayeredreadl2(unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4532 extern uchar1 __surf2DLayeredreadc1(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4533 extern uchar2 __surf2DLayeredreadc2(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4534 extern uchar4 __surf2DLayeredreadc4(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4535 extern ushort1 __surf2DLayeredreads1(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4536 extern ushort2 __surf2DLayeredreads2(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4537 extern ushort4 __surf2DLayeredreads4(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4538 extern uint1 __surf2DLayeredreadu1(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4539 extern uint2 __surf2DLayeredreadu2(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4540 extern uint4 __surf2DLayeredreadu4(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4541 extern ulonglong1 __surf2DLayeredreadl1(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4542 extern ulonglong2 __surf2DLayeredreadl2(unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4543 extern void __surf1Dwritec1( uchar1, unsigned long long, int, enum cudaSurfaceBoundaryMode);
4544 extern void __surf1Dwritec2( uchar2, unsigned long long, int, enum cudaSurfaceBoundaryMode);
4545 extern void __surf1Dwritec4( uchar4, unsigned long long, int, enum cudaSurfaceBoundaryMode);
4546 extern void __surf1Dwrites1( ushort1, unsigned long long, int, enum cudaSurfaceBoundaryMode);
4547 extern void __surf1Dwrites2( ushort2, unsigned long long, int, enum cudaSurfaceBoundaryMode);
4548 extern void __surf1Dwrites4( ushort4, unsigned long long, int, enum cudaSurfaceBoundaryMode);
4549 extern void __surf1Dwriteu1( uint1, unsigned long long, int, enum cudaSurfaceBoundaryMode);
4550 extern void __surf1Dwriteu2( uint2, unsigned long long, int, enum cudaSurfaceBoundaryMode);
4551 extern void __surf1Dwriteu4( uint4, unsigned long long, int, enum cudaSurfaceBoundaryMode);
4552 extern void __surf1Dwritel1(ulonglong1, unsigned long long, int, enum cudaSurfaceBoundaryMode);
4553 extern void __surf1Dwritel2(ulonglong2, unsigned long long, int, enum cudaSurfaceBoundaryMode);
4554 extern void __surf2Dwritec1( uchar1, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4555 extern void __surf2Dwritec2( uchar2, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4556 extern void __surf2Dwritec4( uchar4, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4557 extern void __surf2Dwrites1( ushort1, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4558 extern void __surf2Dwrites2( ushort2, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4559 extern void __surf2Dwrites4( ushort4, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4560 extern void __surf2Dwriteu1( uint1, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4561 extern void __surf2Dwriteu2( uint2, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4562 extern void __surf2Dwriteu4( uint4, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4563 extern void __surf2Dwritel1(ulonglong1, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4564 extern void __surf2Dwritel2(ulonglong2, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4565 extern void __surf3Dwritec1( uchar1 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4566 extern void __surf3Dwritec2( uchar2 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4567 extern void __surf3Dwritec4( uchar4 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4568 extern void __surf3Dwrites1( ushort1 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4569 extern void __surf3Dwrites2( ushort2 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4570 extern void __surf3Dwrites4( ushort4 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4571 extern void __surf3Dwriteu1( uint1 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4572 extern void __surf3Dwriteu2( uint2 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4573 extern void __surf3Dwriteu4( uint4 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4574 extern void __surf3Dwritel1(ulonglong1 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4575 extern void __surf3Dwritel2(ulonglong2 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4576 extern void __surf1DLayeredwritec1( uchar1 val, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4577 extern void __surf1DLayeredwritec2( uchar2 val, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4578 extern void __surf1DLayeredwritec4( uchar4 val, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4579 extern void __surf1DLayeredwrites1( ushort1 val, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4580 extern void __surf1DLayeredwrites2( ushort2 val, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4581 extern void __surf1DLayeredwrites4( ushort4 val, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4582 extern void __surf1DLayeredwriteu1( uint1 val, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4583 extern void __surf1DLayeredwriteu2( uint2 val, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4584 extern void __surf1DLayeredwriteu4( uint4 val, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4585 extern void __surf1DLayeredwritel1(ulonglong1 val, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4586 extern void __surf1DLayeredwritel2(ulonglong2 val, unsigned long long, int, int, enum cudaSurfaceBoundaryMode);
4587 extern void __surf2DLayeredwritec1( uchar1 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4588 extern void __surf2DLayeredwritec2( uchar2 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4589 extern void __surf2DLayeredwritec4( uchar4 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4590 extern void __surf2DLayeredwrites1( ushort1 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4591 extern void __surf2DLayeredwrites2( ushort2 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4592 extern void __surf2DLayeredwrites4( ushort4 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4593 extern void __surf2DLayeredwriteu1( uint1 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4594 extern void __surf2DLayeredwriteu2( uint2 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4595 extern void __surf2DLayeredwriteu4( uint4 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4596 extern void __surf2DLayeredwritel1(ulonglong1 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4597 extern void __surf2DLayeredwritel2(ulonglong2 val, unsigned long long, int, int, int, enum cudaSurfaceBoundaryMode);
4598 #else /* __CUDANVVM__ */
4599 extern uchar1 __surf1Dreadc1(const void*, int, enum cudaSurfaceBoundaryMode);
4600 extern uchar2 __surf1Dreadc2(const void*, int, enum cudaSurfaceBoundaryMode);
4601 extern uchar4 __surf1Dreadc4(const void*, int, enum cudaSurfaceBoundaryMode);
4602 extern ushort1 __surf1Dreads1(const void*, int, enum cudaSurfaceBoundaryMode);
4603 extern ushort2 __surf1Dreads2(const void*, int, enum cudaSurfaceBoundaryMode);
4604 extern ushort4 __surf1Dreads4(const void*, int, enum cudaSurfaceBoundaryMode);
4605 extern uint1 __surf1Dreadu1(const void*, int, enum cudaSurfaceBoundaryMode);
4606 extern uint2 __surf1Dreadu2(const void*, int, enum cudaSurfaceBoundaryMode);
4607 extern uint4 __surf1Dreadu4(const void*, int, enum cudaSurfaceBoundaryMode);
4608 extern ulonglong1 __surf1Dreadl1(const void*, int, enum cudaSurfaceBoundaryMode);
4609 extern ulonglong2 __surf1Dreadl2(const void*, int, enum cudaSurfaceBoundaryMode);
4610 extern uchar1 __surf2Dreadc1(const void*, int, int, enum cudaSurfaceBoundaryMode);
4611 extern uchar2 __surf2Dreadc2(const void*, int, int, enum cudaSurfaceBoundaryMode);
4612 extern uchar4 __surf2Dreadc4(const void*, int, int, enum cudaSurfaceBoundaryMode);
4613 extern ushort1 __surf2Dreads1(const void*, int, int, enum cudaSurfaceBoundaryMode);
4614 extern ushort2 __surf2Dreads2(const void*, int, int, enum cudaSurfaceBoundaryMode);
4615 extern ushort4 __surf2Dreads4(const void*, int, int, enum cudaSurfaceBoundaryMode);
4616 extern uint1 __surf2Dreadu1(const void*, int, int, enum cudaSurfaceBoundaryMode);
4617 extern uint2 __surf2Dreadu2(const void*, int, int, enum cudaSurfaceBoundaryMode);
4618 extern uint4 __surf2Dreadu4(const void*, int, int, enum cudaSurfaceBoundaryMode);
4619 extern ulonglong1 __surf2Dreadl1(const void*, int, int, enum cudaSurfaceBoundaryMode);
4620 extern ulonglong2 __surf2Dreadl2(const void*, int, int, enum cudaSurfaceBoundaryMode);
4621 extern uchar1 __surf3Dreadc1(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4622 extern uchar2 __surf3Dreadc2(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4623 extern uchar4 __surf3Dreadc4(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4624 extern ushort1 __surf3Dreads1(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4625 extern ushort2 __surf3Dreads2(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4626 extern ushort4 __surf3Dreads4(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4627 extern uint1 __surf3Dreadu1(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4628 extern uint2 __surf3Dreadu2(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4629 extern uint4 __surf3Dreadu4(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4630 extern ulonglong1 __surf3Dreadl1(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4631 extern ulonglong2 __surf3Dreadl2(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4632 extern uchar1 __surf1DLayeredreadc1(const void*, int, int, enum cudaSurfaceBoundaryMode);
4633 extern uchar2 __surf1DLayeredreadc2(const void*, int, int, enum cudaSurfaceBoundaryMode);
4634 extern uchar4 __surf1DLayeredreadc4(const void*, int, int, enum cudaSurfaceBoundaryMode);
4635 extern ushort1 __surf1DLayeredreads1(const void*, int, int, enum cudaSurfaceBoundaryMode);
4636 extern ushort2 __surf1DLayeredreads2(const void*, int, int, enum cudaSurfaceBoundaryMode);
4637 extern ushort4 __surf1DLayeredreads4(const void*, int, int, enum cudaSurfaceBoundaryMode);
4638 extern uint1 __surf1DLayeredreadu1(const void*, int, int, enum cudaSurfaceBoundaryMode);
4639 extern uint2 __surf1DLayeredreadu2(const void*, int, int, enum cudaSurfaceBoundaryMode);
4640 extern uint4 __surf1DLayeredreadu4(const void*, int, int, enum cudaSurfaceBoundaryMode);
4641 extern ulonglong1 __surf1DLayeredreadl1(const void*, int, int, enum cudaSurfaceBoundaryMode);
4642 extern ulonglong2 __surf1DLayeredreadl2(const void*, int, int, enum cudaSurfaceBoundaryMode);
4643 extern uchar1 __surf2DLayeredreadc1(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4644 extern uchar2 __surf2DLayeredreadc2(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4645 extern uchar4 __surf2DLayeredreadc4(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4646 extern ushort1 __surf2DLayeredreads1(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4647 extern ushort2 __surf2DLayeredreads2(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4648 extern ushort4 __surf2DLayeredreads4(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4649 extern uint1 __surf2DLayeredreadu1(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4650 extern uint2 __surf2DLayeredreadu2(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4651 extern uint4 __surf2DLayeredreadu4(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4652 extern ulonglong1 __surf2DLayeredreadl1(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4653 extern ulonglong2 __surf2DLayeredreadl2(const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4654 extern void __surf1Dwritec1( uchar1, const void*, int, enum cudaSurfaceBoundaryMode);
4655 extern void __surf1Dwritec2( uchar2, const void*, int, enum cudaSurfaceBoundaryMode);
4656 extern void __surf1Dwritec4( uchar4, const void*, int, enum cudaSurfaceBoundaryMode);
4657 extern void __surf1Dwrites1( ushort1, const void*, int, enum cudaSurfaceBoundaryMode);
4658 extern void __surf1Dwrites2( ushort2, const void*, int, enum cudaSurfaceBoundaryMode);
4659 extern void __surf1Dwrites4( ushort4, const void*, int, enum cudaSurfaceBoundaryMode);
4660 extern void __surf1Dwriteu1( uint1, const void*, int, enum cudaSurfaceBoundaryMode);
4661 extern void __surf1Dwriteu2( uint2, const void*, int, enum cudaSurfaceBoundaryMode);
4662 extern void __surf1Dwriteu4( uint4, const void*, int, enum cudaSurfaceBoundaryMode);
4663 extern void __surf1Dwritel1(ulonglong1, const void*, int, enum cudaSurfaceBoundaryMode);
4664 extern void __surf1Dwritel2(ulonglong2, const void*, int, enum cudaSurfaceBoundaryMode);
4665 extern void __surf2Dwritec1( uchar1, const void*, int, int, enum cudaSurfaceBoundaryMode);
4666 extern void __surf2Dwritec2( uchar2, const void*, int, int, enum cudaSurfaceBoundaryMode);
4667 extern void __surf2Dwritec4( uchar4, const void*, int, int, enum cudaSurfaceBoundaryMode);
4668 extern void __surf2Dwrites1( ushort1, const void*, int, int, enum cudaSurfaceBoundaryMode);
4669 extern void __surf2Dwrites2( ushort2, const void*, int, int, enum cudaSurfaceBoundaryMode);
4670 extern void __surf2Dwrites4( ushort4, const void*, int, int, enum cudaSurfaceBoundaryMode);
4671 extern void __surf2Dwriteu1( uint1, const void*, int, int, enum cudaSurfaceBoundaryMode);
4672 extern void __surf2Dwriteu2( uint2, const void*, int, int, enum cudaSurfaceBoundaryMode);
4673 extern void __surf2Dwriteu4( uint4, const void*, int, int, enum cudaSurfaceBoundaryMode);
4674 extern void __surf2Dwritel1(ulonglong1, const void*, int, int, enum cudaSurfaceBoundaryMode);
4675 extern void __surf2Dwritel2(ulonglong2, const void*, int, int, enum cudaSurfaceBoundaryMode);
4676 extern void __surf3Dwritec1( uchar1 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4677 extern void __surf3Dwritec2( uchar2 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4678 extern void __surf3Dwritec4( uchar4 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4679 extern void __surf3Dwrites1( ushort1 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4680 extern void __surf3Dwrites2( ushort2 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4681 extern void __surf3Dwrites4( ushort4 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4682 extern void __surf3Dwriteu1( uint1 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4683 extern void __surf3Dwriteu2( uint2 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4684 extern void __surf3Dwriteu4( uint4 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4685 extern void __surf3Dwritel1(ulonglong1 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4686 extern void __surf3Dwritel2(ulonglong2 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4687 extern void __surf1DLayeredwritec1( uchar1 val, const void*, int, int, enum cudaSurfaceBoundaryMode);
4688 extern void __surf1DLayeredwritec2( uchar2 val, const void*, int, int, enum cudaSurfaceBoundaryMode);
4689 extern void __surf1DLayeredwritec4( uchar4 val, const void*, int, int, enum cudaSurfaceBoundaryMode);
4690 extern void __surf1DLayeredwrites1( ushort1 val, const void*, int, int, enum cudaSurfaceBoundaryMode);
4691 extern void __surf1DLayeredwrites2( ushort2 val, const void*, int, int, enum cudaSurfaceBoundaryMode);
4692 extern void __surf1DLayeredwrites4( ushort4 val, const void*, int, int, enum cudaSurfaceBoundaryMode);
4693 extern void __surf1DLayeredwriteu1( uint1 val, const void*, int, int, enum cudaSurfaceBoundaryMode);
4694 extern void __surf1DLayeredwriteu2( uint2 val, const void*, int, int, enum cudaSurfaceBoundaryMode);
4695 extern void __surf1DLayeredwriteu4( uint4 val, const void*, int, int, enum cudaSurfaceBoundaryMode);
4696 extern void __surf1DLayeredwritel1(ulonglong1 val, const void*, int, int, enum cudaSurfaceBoundaryMode);
4697 extern void __surf1DLayeredwritel2(ulonglong2 val, const void*, int, int, enum cudaSurfaceBoundaryMode);
4698 extern void __surf2DLayeredwritec1( uchar1 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4699 extern void __surf2DLayeredwritec2( uchar2 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4700 extern void __surf2DLayeredwritec4( uchar4 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4701 extern void __surf2DLayeredwrites1( ushort1 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4702 extern void __surf2DLayeredwrites2( ushort2 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4703 extern void __surf2DLayeredwrites4( ushort4 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4704 extern void __surf2DLayeredwriteu1( uint1 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4705 extern void __surf2DLayeredwriteu2( uint2 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4706 extern void __surf2DLayeredwriteu4( uint4 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4707 extern void __surf2DLayeredwritel1(ulonglong1 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4708 extern void __surf2DLayeredwritel2(ulonglong2 val, const void*, int, int, int, enum cudaSurfaceBoundaryMode);
4709 #endif /* __CUDANVVM__ */
4711 // Cubemap and cubemap layered surfaces use 2D Layered instrinsics
4712 #define __surfCubemapreadc1 __surf2DLayeredreadc1
4713 #define __surfCubemapreadc2 __surf2DLayeredreadc2
4714 #define __surfCubemapreadc4 __surf2DLayeredreadc4
4715 #define __surfCubemapreads1 __surf2DLayeredreads1
4716 #define __surfCubemapreads2 __surf2DLayeredreads2
4717 #define __surfCubemapreads4 __surf2DLayeredreads4
4718 #define __surfCubemapreadu1 __surf2DLayeredreadu1
4719 #define __surfCubemapreadu2 __surf2DLayeredreadu2
4720 #define __surfCubemapreadu4 __surf2DLayeredreadu4
4721 #define __surfCubemapreadl1 __surf2DLayeredreadl1
4722 #define __surfCubemapreadl2 __surf2DLayeredreadl2
4723 #define __surfCubemapLayeredreadc1 __surf2DLayeredreadc1
4724 #define __surfCubemapLayeredreadc2 __surf2DLayeredreadc2
4725 #define __surfCubemapLayeredreadc4 __surf2DLayeredreadc4
4726 #define __surfCubemapLayeredreads1 __surf2DLayeredreads1
4727 #define __surfCubemapLayeredreads2 __surf2DLayeredreads2
4728 #define __surfCubemapLayeredreads4 __surf2DLayeredreads4
4729 #define __surfCubemapLayeredreadu1 __surf2DLayeredreadu1
4730 #define __surfCubemapLayeredreadu2 __surf2DLayeredreadu2
4731 #define __surfCubemapLayeredreadu4 __surf2DLayeredreadu4
4732 #define __surfCubemapLayeredreadl1 __surf2DLayeredreadl1
4733 #define __surfCubemapLayeredreadl2 __surf2DLayeredreadl2
4735 #define __surfCubemapwritec1 __surf2DLayeredwritec1
4736 #define __surfCubemapwritec2 __surf2DLayeredwritec2
4737 #define __surfCubemapwritec4 __surf2DLayeredwritec4
4738 #define __surfCubemapwrites1 __surf2DLayeredwrites1
4739 #define __surfCubemapwrites2 __surf2DLayeredwrites2
4740 #define __surfCubemapwrites4 __surf2DLayeredwrites4
4741 #define __surfCubemapwriteu1 __surf2DLayeredwriteu1
4742 #define __surfCubemapwriteu2 __surf2DLayeredwriteu2
4743 #define __surfCubemapwriteu4 __surf2DLayeredwriteu4
4744 #define __surfCubemapwritel1 __surf2DLayeredwritel1
4745 #define __surfCubemapwritel2 __surf2DLayeredwritel2
4746 #define __surfCubemapLayeredwritec1 __surf2DLayeredwritec1
4747 #define __surfCubemapLayeredwritec2 __surf2DLayeredwritec2
4748 #define __surfCubemapLayeredwritec4 __surf2DLayeredwritec4
4749 #define __surfCubemapLayeredwrites1 __surf2DLayeredwrites1
4750 #define __surfCubemapLayeredwrites2 __surf2DLayeredwrites2
4751 #define __surfCubemapLayeredwrites4 __surf2DLayeredwrites4
4752 #define __surfCubemapLayeredwriteu1 __surf2DLayeredwriteu1
4753 #define __surfCubemapLayeredwriteu2 __surf2DLayeredwriteu2
4754 #define __surfCubemapLayeredwriteu4 __surf2DLayeredwriteu4
4755 #define __surfCubemapLayeredwritel1 __surf2DLayeredwritel1
4756 #define __surfCubemapLayeredwritel2 __surf2DLayeredwritel2
4758 #endif /* __cplusplus && __CUDACC__ */
4760 #endif /* !__SURFACE_FUNCTIONS_H__ */