Changeset View
Changeset View
Standalone View
Standalone View
intern/cycles/kernel/kernel_globals.h
| Show First 20 Lines • Show All 70 Lines • ▼ Show 20 Lines | |||||
| /* For CUDA, constant memory textures must be globals, so we can't put them | /* For CUDA, constant memory textures must be globals, so we can't put them | ||||
| * into a struct. As a result we don't actually use this struct and use actual | * into a struct. As a result we don't actually use this struct and use actual | ||||
| * globals and simply pass along a NULL pointer everywhere, which we hope gets | * globals and simply pass along a NULL pointer everywhere, which we hope gets | ||||
| * optimized out. */ | * optimized out. */ | ||||
| #ifdef __KERNEL_CUDA__ | #ifdef __KERNEL_CUDA__ | ||||
| __constant__ KernelData __data; | __constant__ KernelData __data; | ||||
| typedef struct KernelGlobals {} KernelGlobals; | typedef struct KernelGlobals { | ||||
| /* NOTE: Keep the size in sync with SHADOW_STACK_MAX_HITS. */ | |||||
| Intersection hits_stack[64]; | |||||
| } KernelGlobals; | |||||
| # ifdef __KERNEL_CUDA_TEX_STORAGE__ | # ifdef __KERNEL_CUDA_TEX_STORAGE__ | ||||
| # define KERNEL_TEX(type, ttype, name) ttype name; | # define KERNEL_TEX(type, ttype, name) ttype name; | ||||
| # else | # else | ||||
| # define KERNEL_TEX(type, ttype, name) const __constant__ __device__ type *name; | # define KERNEL_TEX(type, ttype, name) const __constant__ __device__ type *name; | ||||
| # endif | # endif | ||||
| # define KERNEL_IMAGE_TEX(type, ttype, name) ttype name; | # define KERNEL_IMAGE_TEX(type, ttype, name) ttype name; | ||||
| # include "kernel_textures.h" | # include "kernel_textures.h" | ||||
| ▲ Show 20 Lines • Show All 58 Lines • Show Last 20 Lines | |||||