forked from bartvdbraak/blender
Cycles: Remove fermi related defines from the code.
Did not touch Texture related defines, that comes next.
This commit is contained in:
parent
2eaf90b305
commit
e1ef902058
@ -18,12 +18,6 @@ CCL_NAMESPACE_BEGIN
|
||||
|
||||
#ifdef __HAIR__
|
||||
|
||||
#if defined(__KERNEL_CUDA__) && (__CUDA_ARCH__ < 300)
|
||||
# define ccl_device_curveintersect ccl_device
|
||||
#else
|
||||
# define ccl_device_curveintersect ccl_device_forceinline
|
||||
#endif
|
||||
|
||||
#ifdef __KERNEL_SSE2__
|
||||
ccl_device_inline ssef transform_point_T3(const ssef t[3], const ssef &a)
|
||||
{
|
||||
@ -32,7 +26,7 @@ ccl_device_inline ssef transform_point_T3(const ssef t[3], const ssef &a)
|
||||
#endif
|
||||
|
||||
/* On CPU pass P and dir by reference to aligned vector. */
|
||||
ccl_device_curveintersect bool cardinal_curve_intersect(
|
||||
ccl_device_forceinline bool cardinal_curve_intersect(
|
||||
KernelGlobals *kg,
|
||||
Intersection *isect,
|
||||
const float3 ccl_ref P,
|
||||
@ -505,7 +499,7 @@ ccl_device_curveintersect bool cardinal_curve_intersect(
|
||||
return hit;
|
||||
}
|
||||
|
||||
ccl_device_curveintersect bool curve_intersect(KernelGlobals *kg,
|
||||
ccl_device_forceinline bool curve_intersect(KernelGlobals *kg,
|
||||
Intersection *isect,
|
||||
float3 P,
|
||||
float3 direction,
|
||||
|
@ -50,10 +50,7 @@ __device__ half __float2half(const float f)
|
||||
/* Qualifier wrappers for different names on different devices */
|
||||
|
||||
#define ccl_device __device__ __inline__
|
||||
#if __CUDA_ARCH__ < 300
|
||||
# define ccl_device_inline __device__ __inline__
|
||||
# define ccl_device_forceinline __device__ __forceinline__
|
||||
#elif __CUDA_ARCH__ < 500
|
||||
#if __CUDA_ARCH__ < 500
|
||||
# define ccl_device_inline __device__ __forceinline__
|
||||
# define ccl_device_forceinline __device__ __forceinline__
|
||||
#else
|
||||
|
@ -16,20 +16,8 @@
|
||||
|
||||
/* device data taken from CUDA occupancy calculator */
|
||||
|
||||
/* 2.0 and 2.1 */
|
||||
#if __CUDA_ARCH__ == 200 || __CUDA_ARCH__ == 210
|
||||
# define CUDA_MULTIPRESSOR_MAX_REGISTERS 32768
|
||||
# define CUDA_MULTIPROCESSOR_MAX_BLOCKS 8
|
||||
# define CUDA_BLOCK_MAX_THREADS 1024
|
||||
# define CUDA_THREAD_MAX_REGISTERS 63
|
||||
|
||||
/* tunable parameters */
|
||||
# define CUDA_THREADS_BLOCK_WIDTH 16
|
||||
# define CUDA_KERNEL_MAX_REGISTERS 32
|
||||
# define CUDA_KERNEL_BRANCHED_MAX_REGISTERS 40
|
||||
|
||||
/* 3.0 and 3.5 */
|
||||
#elif __CUDA_ARCH__ == 300 || __CUDA_ARCH__ == 350
|
||||
#if __CUDA_ARCH__ == 300 || __CUDA_ARCH__ == 350
|
||||
# define CUDA_MULTIPRESSOR_MAX_REGISTERS 65536
|
||||
# define CUDA_MULTIPROCESSOR_MAX_BLOCKS 16
|
||||
# define CUDA_BLOCK_MAX_THREADS 1024
|
||||
|
@ -79,12 +79,7 @@ ccl_device bool ray_aligned_disk_intersect(
|
||||
return true;
|
||||
}
|
||||
|
||||
#if defined(__KERNEL_CUDA__) && __CUDA_ARCH__ < 300
|
||||
ccl_device_inline
|
||||
#else
|
||||
ccl_device_forceinline
|
||||
#endif
|
||||
bool ray_triangle_intersect(
|
||||
ccl_device_forceinline bool ray_triangle_intersect(
|
||||
float3 ray_P, float3 ray_dir, float ray_t,
|
||||
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
|
||||
const ssef *ssef_verts,
|
||||
|
Loading…
Reference in New Issue
Block a user