Welcome to mirror list, hosted at ThFree Co, Russian Federation.

git.blender.org/blender.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
path: root/intern/cycles/kernel/device/optix
diff options
context:
space:
mode:
Diffstat (limited to 'intern/cycles/kernel/device/optix')
-rw-r--r--intern/cycles/kernel/device/optix/globals.h16
-rw-r--r--intern/cycles/kernel/device/optix/kernel.cu48
-rw-r--r--intern/cycles/kernel/device/optix/kernel_shader_raytrace.cu8
3 files changed, 36 insertions, 36 deletions
diff --git a/intern/cycles/kernel/device/optix/globals.h b/intern/cycles/kernel/device/optix/globals.h
index bb752c531f0..7af2e421378 100644
--- a/intern/cycles/kernel/device/optix/globals.h
+++ b/intern/cycles/kernel/device/optix/globals.h
@@ -28,21 +28,21 @@ struct KernelParamsOptiX {
/* Global scene data and textures */
KernelData data;
-#define KERNEL_TEX(type, name) const type *name;
-#include "kernel/textures.h"
+#define KERNEL_DATA_ARRAY(type, name) const type *name;
+#include "kernel/data_arrays.h"
/* Integrator state */
- IntegratorStateGPU __integrator_state;
+ IntegratorStateGPU integrator_state;
};
#ifdef __NVCC__
-extern "C" static __constant__ KernelParamsOptiX __params;
+extern "C" static __constant__ KernelParamsOptiX kernel_params;
#endif
/* Abstraction macros */
-#define kernel_data __params.data
-#define kernel_tex_array(t) __params.t
-#define kernel_tex_fetch(t, index) __params.t[(index)]
-#define kernel_integrator_state __params.__integrator_state
+#define kernel_data kernel_params.data
+#define kernel_data_array(name) kernel_params.name
+#define kernel_data_fetch(name, index) kernel_params.name[(index)]
+#define kernel_integrator_state kernel_params.integrator_state
CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu
index 9843b2e99be..949bf41d171 100644
--- a/intern/cycles/kernel/device/optix/kernel.cu
+++ b/intern/cycles/kernel/device/optix/kernel.cu
@@ -51,15 +51,15 @@ ccl_device_forceinline int get_object_id()
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest()
{
const int global_index = optixGetLaunchIndex().x;
- const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
+ const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
global_index;
- integrator_intersect_closest(nullptr, path_index, __params.render_buffer);
+ integrator_intersect_closest(nullptr, path_index, kernel_params.render_buffer);
}
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_shadow()
{
const int global_index = optixGetLaunchIndex().x;
- const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
+ const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
global_index;
integrator_intersect_shadow(nullptr, path_index);
}
@@ -67,7 +67,7 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_shadow()
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_subsurface()
{
const int global_index = optixGetLaunchIndex().x;
- const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
+ const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
global_index;
integrator_intersect_subsurface(nullptr, path_index);
}
@@ -75,7 +75,7 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_subsurfac
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_volume_stack()
{
const int global_index = optixGetLaunchIndex().x;
- const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
+ const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
global_index;
integrator_intersect_volume_stack(nullptr, path_index);
}
@@ -151,17 +151,17 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
isect->t = optixGetRayTmax();
isect->prim = prim;
isect->object = get_object_id();
- isect->type = kernel_tex_fetch(__objects, isect->object).primitive_type;
+ isect->type = kernel_data_fetch(objects, isect->object).primitive_type;
const float2 barycentrics = optixGetTriangleBarycentrics();
isect->u = 1.0f - barycentrics.y - barycentrics.x;
isect->v = barycentrics.x;
/* Record geometric normal. */
- const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
- const float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0);
- const float3 tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1);
- const float3 tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
+ const uint tri_vindex = kernel_data_fetch(tri_vindex, prim).w;
+ const float3 tri_a = kernel_data_fetch(tri_verts, tri_vindex + 0);
+ const float3 tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1);
+ const float3 tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2);
local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
/* Continue tracing (without this the trace call would return after the first hit). */
@@ -176,7 +176,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
const uint object = get_object_id();
# ifdef __VISIBILITY_FLAG__
const uint visibility = optixGetPayload_4();
- if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
+ if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
return optixIgnoreIntersection();
}
# endif
@@ -192,14 +192,14 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
const float2 barycentrics = optixGetTriangleBarycentrics();
u = 1.0f - barycentrics.y - barycentrics.x;
v = barycentrics.x;
- type = kernel_tex_fetch(__objects, object).primitive_type;
+ type = kernel_data_fetch(objects, object).primitive_type;
}
# ifdef __HAIR__
else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) {
u = __uint_as_float(optixGetAttribute_0());
v = __uint_as_float(optixGetAttribute_1());
- const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
+ const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
type = segment.type;
prim = segment.prim;
@@ -212,7 +212,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
}
# endif
else {
- type = kernel_tex_fetch(__objects, object).primitive_type;
+ type = kernel_data_fetch(objects, object).primitive_type;
u = 0.0f;
v = 0.0f;
}
@@ -307,12 +307,12 @@ extern "C" __global__ void __anyhit__kernel_optix_volume_test()
const uint object = get_object_id();
#ifdef __VISIBILITY_FLAG__
const uint visibility = optixGetPayload_4();
- if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
+ if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
return optixIgnoreIntersection();
}
#endif
- if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
+ if ((kernel_data_fetch(object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
return optixIgnoreIntersection();
}
@@ -340,7 +340,7 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
const uint object = get_object_id();
const uint visibility = optixGetPayload_4();
#ifdef __VISIBILITY_FLAG__
- if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
+ if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
return optixIgnoreIntersection();
}
#endif
@@ -377,10 +377,10 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x));
optixSetPayload_2(__float_as_uint(barycentrics.x));
optixSetPayload_3(prim);
- optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type);
+ optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type);
}
else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) {
- const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
+ const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */
optixSetPayload_2(optixGetAttribute_1());
optixSetPayload_3(segment.prim);
@@ -390,7 +390,7 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
optixSetPayload_1(0);
optixSetPayload_2(0);
optixSetPayload_3(prim);
- optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type);
+ optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type);
}
}
@@ -401,7 +401,7 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type)
# ifdef __VISIBILITY_FLAG__
const uint visibility = optixGetPayload_4();
- if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
+ if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
return;
}
# endif
@@ -436,7 +436,7 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type)
extern "C" __global__ void __intersection__curve_ribbon()
{
- const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, optixGetPrimitiveIndex());
+ const KernelCurveSegment segment = kernel_data_fetch(curve_segments, optixGetPrimitiveIndex());
const int prim = segment.prim;
const int type = segment.type;
if (type & PRIMITIVE_CURVE_RIBBON) {
@@ -451,11 +451,11 @@ extern "C" __global__ void __intersection__point()
{
const int prim = optixGetPrimitiveIndex();
const int object = get_object_id();
- const int type = kernel_tex_fetch(__objects, object).primitive_type;
+ const int type = kernel_data_fetch(objects, object).primitive_type;
# ifdef __VISIBILITY_FLAG__
const uint visibility = optixGetPayload_4();
- if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
+ if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
return;
}
# endif
diff --git a/intern/cycles/kernel/device/optix/kernel_shader_raytrace.cu b/intern/cycles/kernel/device/optix/kernel_shader_raytrace.cu
index 3bd57bc0f1a..41e6224f6da 100644
--- a/intern/cycles/kernel/device/optix/kernel_shader_raytrace.cu
+++ b/intern/cycles/kernel/device/optix/kernel_shader_raytrace.cu
@@ -11,15 +11,15 @@
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_surface_raytrace()
{
const int global_index = optixGetLaunchIndex().x;
- const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
+ const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
global_index;
- integrator_shade_surface_raytrace(nullptr, path_index, __params.render_buffer);
+ integrator_shade_surface_raytrace(nullptr, path_index, kernel_params.render_buffer);
}
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_surface_mnee()
{
const int global_index = optixGetLaunchIndex().x;
- const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
+ const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
global_index;
- integrator_shade_surface_mnee(nullptr, path_index, __params.render_buffer);
+ integrator_shade_surface_mnee(nullptr, path_index, kernel_params.render_buffer);
}
generated by cgit v1.2.3 (git 2.25.1) at 2024-06-19 22:08:55 +0300