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
diff options
context:
space:
mode:
authorDalai Felinto <dfelinto@gmail.com>2016-01-15 18:00:56 +0300
committerDalai Felinto <dfelinto@gmail.com>2016-01-15 18:00:56 +0300
commit9a76354585e2cd2011267e79bd99ca59a06588f8 (patch)
treee775e7c44dc210ef9978b483930ade6a9b4d6fc5 /intern/cycles/kernel
parent9137a4401440d3f3206e989f49f3539079d685b8 (diff)
Cycles-Bake: Custom Baking passes
The combined pass is built with the contributions the user finds fit. It is useful for lightmap baking, as well as non-view dependent effects baking. The manual will be updated once we get closer to the 2.77 release. Meanwhile the new page can be found here: http://dalaifelinto.com/blender-manual/render/cycles/baking.html Reviewers: sergey, brecht Differential Revision: https://developer.blender.org/D1674
Diffstat (limited to 'intern/cycles/kernel')
-rw-r--r--intern/cycles/kernel/kernel_bake.h202
-rw-r--r--intern/cycles/kernel/kernel_types.h45
-rw-r--r--intern/cycles/kernel/kernels/cpu/kernel_cpu.h1
-rw-r--r--intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h2
-rw-r--r--intern/cycles/kernel/kernels/cuda/kernel.cu4
-rw-r--r--intern/cycles/kernel/kernels/opencl/kernel.cl6
6 files changed, 158 insertions, 102 deletions
diff --git a/intern/cycles/kernel/kernel_bake.h b/intern/cycles/kernel/kernel_bake.h
index 2a2220ceb99..31e58de0b48 100644
--- a/intern/cycles/kernel/kernel_bake.h
+++ b/intern/cycles/kernel/kernel_bake.h
@@ -19,7 +19,7 @@ CCL_NAMESPACE_BEGIN
#undef USE_BAKE_JITTER
ccl_device void compute_light_pass(KernelGlobals *kg, ShaderData *sd, PathRadiance *L, RNG rng,
- const bool is_combined, const bool is_ao, const bool is_sss, int sample)
+ const bool is_ao, const bool is_sss, int sample)
{
/* initialize master radiance accumulator */
kernel_assert(kernel_data.film.use_light_pass);
@@ -56,13 +56,13 @@ ccl_device void compute_light_pass(KernelGlobals *kg, ShaderData *sd, PathRadian
#endif
/* sample ambient occlusion */
- if(is_combined || is_ao) {
+ if(is_ao) {
kernel_path_ao(kg, sd, &L_sample, &state, &rng, throughput);
}
#ifdef __SUBSURFACE__
/* sample subsurface scattering */
- if((is_combined || is_sss_sample) && (sd->flag & SD_BSSRDF)) {
+ if(is_sss_sample && (sd->flag & SD_BSSRDF)) {
/* when mixing BSSRDF and BSDF closures we should skip BSDF lighting if scattering was successful */
SubsurfaceIndirectRays ss_indirect;
kernel_path_subsurface_init_indirect(&ss_indirect);
@@ -124,13 +124,13 @@ ccl_device void compute_light_pass(KernelGlobals *kg, ShaderData *sd, PathRadian
/* branched path tracer */
/* sample ambient occlusion */
- if(is_combined || is_ao) {
+ if(is_ao) {
kernel_branched_path_ao(kg, sd, &L_sample, &state, &rng, throughput);
}
#ifdef __SUBSURFACE__
/* sample subsurface scattering */
- if((is_combined || is_sss_sample) && (sd->flag & SD_BSSRDF)) {
+ if(is_sss_sample && (sd->flag & SD_BSSRDF)) {
/* when mixing BSSRDF and BSDF closures we should skip BSDF lighting if scattering was successful */
kernel_branched_path_subsurface_scatter(kg, sd, &L_sample, &state, &rng, &ray, throughput);
}
@@ -175,21 +175,27 @@ ccl_device bool is_aa_pass(ShaderEvalType type)
}
}
-ccl_device bool is_light_pass(ShaderEvalType type)
+/* Keep it synced with BakeManager::is_light_pass. */
+ccl_device bool is_light_pass(ShaderEvalType type, const int pass_filter)
{
switch(type) {
case SHADER_EVAL_AO:
- case SHADER_EVAL_COMBINED:
case SHADER_EVAL_SHADOW:
- case SHADER_EVAL_DIFFUSE_DIRECT:
- case SHADER_EVAL_GLOSSY_DIRECT:
- case SHADER_EVAL_TRANSMISSION_DIRECT:
- case SHADER_EVAL_SUBSURFACE_DIRECT:
- case SHADER_EVAL_DIFFUSE_INDIRECT:
- case SHADER_EVAL_GLOSSY_INDIRECT:
- case SHADER_EVAL_TRANSMISSION_INDIRECT:
- case SHADER_EVAL_SUBSURFACE_INDIRECT:
return true;
+ case SHADER_EVAL_DIFFUSE:
+ case SHADER_EVAL_GLOSSY:
+ case SHADER_EVAL_TRANSMISSION:
+ return ((pass_filter & BAKE_FILTER_DIRECT) != 0) ||
+ ((pass_filter & BAKE_FILTER_INDIRECT) != 0);
+ case SHADER_EVAL_COMBINED:
+ return ((pass_filter & BAKE_FILTER_AO) != 0) ||
+ ((pass_filter & BAKE_FILTER_EMISSION) != 0) ||
+ ((((pass_filter & BAKE_FILTER_DIRECT) != 0) ||
+ ((pass_filter & BAKE_FILTER_INDIRECT) != 0)) &&
+ (((pass_filter & BAKE_FILTER_DIFFUSE) != 0) ||
+ ((pass_filter & BAKE_FILTER_GLOSSY) != 0) ||
+ ((pass_filter & BAKE_FILTER_TRANSMISSION) != 0) ||
+ ((pass_filter & BAKE_FILTER_SUBSURFACE) != 0)));
default:
return false;
}
@@ -208,15 +214,52 @@ ccl_device_inline float bake_clamp_mirror_repeat(float u)
return (((int)fu) & 1)? 1.0f - u: u;
}
+ccl_device float3 kernel_bake_evaluate_direct_indirect(KernelGlobals *kg, ShaderData *sd, PathState *state,
+ float3 (*shader_bsdf)(KernelGlobals *kg, ShaderData *sd),
+ float3 direct, float3 indirect, const int pass_filter)
+{
+ float3 color;
+ const bool is_color = (pass_filter & BAKE_FILTER_COLOR) != 0;
+ const bool is_direct = (pass_filter & BAKE_FILTER_DIRECT) != 0;
+ const bool is_indirect = (pass_filter & BAKE_FILTER_INDIRECT) != 0;
+ float3 out = make_float3(0.0f, 0.0f, 0.0f);
+
+ if(is_color) {
+ if(is_direct || is_indirect) {
+ /* Leave direct and diffuse channel colored. */
+ color = make_float3(1.0f, 1.0f, 1.0f);
+ }
+ else {
+ /* surface color of the pass only */
+ shader_eval_surface(kg, sd, state, 0.0f, 0, SHADER_CONTEXT_MAIN);
+ return shader_bsdf(kg, sd);
+ }
+ }
+ else {
+ shader_eval_surface(kg, sd, state, 0.0f, 0, SHADER_CONTEXT_MAIN);
+ color = shader_bsdf(kg, sd);
+ }
+
+ if(is_direct) {
+ out += safe_divide_color(direct, color);
+ }
+
+ if(is_indirect) {
+ out += safe_divide_color(indirect, color);
+ }
+
+ return out;
+}
+
ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input, ccl_global float4 *output,
- ShaderEvalType type, int i, int offset, int sample)
+ ShaderEvalType type, int pass_filter, int i, int offset, int sample)
{
ShaderData sd;
PathState state = {0};
uint4 in = input[i * 2];
uint4 diff = input[i * 2 + 1];
- float3 out;
+ float3 out = make_float3(0.0f, 0.0f, 0.0f);
int object = in.x;
int prim = in.y;
@@ -279,13 +322,23 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
sd.dv.dy = dvdy;
/* light passes */
- if(is_light_pass(type)) {
- compute_light_pass(kg, &sd, &L, rng,
- (type == SHADER_EVAL_COMBINED),
- (type == SHADER_EVAL_AO),
- (type == SHADER_EVAL_SUBSURFACE_DIRECT ||
- type == SHADER_EVAL_SUBSURFACE_INDIRECT),
- sample);
+ if(is_light_pass(type, pass_filter)) {
+ bool is_ao, is_sss;
+
+ if (type == SHADER_EVAL_COMBINED) {
+ is_ao = (pass_filter & BAKE_FILTER_AO) != 0;
+ is_sss = ((pass_filter & BAKE_FILTER_SUBSURFACE) != 0) &&
+ (((pass_filter & BAKE_FILTER_DIRECT) != 0) ||
+ ((pass_filter & BAKE_FILTER_INDIRECT) != 0));
+ }
+ else {
+ is_ao = (type == SHADER_EVAL_AO);
+ is_sss = (type == SHADER_EVAL_SUBSURFACE) &&
+ (((pass_filter & BAKE_FILTER_DIRECT) != 0) ||
+ ((pass_filter & BAKE_FILTER_INDIRECT) != 0));
+ }
+
+ compute_light_pass(kg, &sd, &L, rng, is_ao, is_sss, sample);
}
switch(type) {
@@ -305,32 +358,6 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
out = primitive_uv(kg, &sd);
break;
}
- case SHADER_EVAL_DIFFUSE_COLOR:
- {
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = shader_bsdf_diffuse(kg, &sd);
- break;
- }
- case SHADER_EVAL_GLOSSY_COLOR:
- {
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = shader_bsdf_glossy(kg, &sd);
- break;
- }
- case SHADER_EVAL_TRANSMISSION_COLOR:
- {
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = shader_bsdf_transmission(kg, &sd);
- break;
- }
- case SHADER_EVAL_SUBSURFACE_COLOR:
- {
-#ifdef __SUBSURFACE__
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = shader_bsdf_subsurface(kg, &sd);
-#endif
- break;
- }
case SHADER_EVAL_EMISSION:
{
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_EMISSION);
@@ -347,7 +374,34 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
}
case SHADER_EVAL_COMBINED:
{
- out = path_radiance_clamp_and_sum(kg, &L);
+ if((pass_filter & BAKE_FILTER_COMBINED) == BAKE_FILTER_COMBINED) {
+ out = path_radiance_clamp_and_sum(kg, &L);
+ break;
+ }
+
+ if((pass_filter & BAKE_FILTER_DIFFUSE_DIRECT) == BAKE_FILTER_DIFFUSE_DIRECT)
+ out += L.direct_diffuse;
+ if((pass_filter & BAKE_FILTER_DIFFUSE_INDIRECT) == BAKE_FILTER_DIFFUSE_INDIRECT)
+ out += L.indirect_diffuse;
+
+ if((pass_filter & BAKE_FILTER_GLOSSY_DIRECT) == BAKE_FILTER_GLOSSY_DIRECT)
+ out += L.direct_glossy;
+ if((pass_filter & BAKE_FILTER_GLOSSY_INDIRECT) == BAKE_FILTER_GLOSSY_INDIRECT)
+ out += L.indirect_glossy;
+
+ if((pass_filter & BAKE_FILTER_TRANSMISSION_DIRECT) == BAKE_FILTER_TRANSMISSION_DIRECT)
+ out += L.direct_transmission;
+ if((pass_filter & BAKE_FILTER_TRANSMISSION_INDIRECT) == BAKE_FILTER_TRANSMISSION_INDIRECT)
+ out += L.indirect_transmission;
+
+ if((pass_filter & BAKE_FILTER_SUBSURFACE_DIRECT) == BAKE_FILTER_SUBSURFACE_DIRECT)
+ out += L.direct_subsurface;
+ if((pass_filter & BAKE_FILTER_SUBSURFACE_INDIRECT) == BAKE_FILTER_SUBSURFACE_INDIRECT)
+ out += L.indirect_subsurface;
+
+ if((pass_filter & BAKE_FILTER_EMISSION) != 0)
+ out += L.emission;
+
break;
}
case SHADER_EVAL_SHADOW:
@@ -355,55 +409,25 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
out = make_float3(L.shadow.x, L.shadow.y, L.shadow.z);
break;
}
- case SHADER_EVAL_DIFFUSE_DIRECT:
- {
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = safe_divide_color(L.direct_diffuse, shader_bsdf_diffuse(kg, &sd));
- break;
- }
- case SHADER_EVAL_GLOSSY_DIRECT:
- {
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = safe_divide_color(L.direct_glossy, shader_bsdf_glossy(kg, &sd));
- break;
- }
- case SHADER_EVAL_TRANSMISSION_DIRECT:
- {
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = safe_divide_color(L.direct_transmission, shader_bsdf_transmission(kg, &sd));
- break;
- }
- case SHADER_EVAL_SUBSURFACE_DIRECT:
- {
-#ifdef __SUBSURFACE__
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = safe_divide_color(L.direct_subsurface, shader_bsdf_subsurface(kg, &sd));
-#endif
- break;
- }
- case SHADER_EVAL_DIFFUSE_INDIRECT:
+ case SHADER_EVAL_DIFFUSE:
{
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = safe_divide_color(L.indirect_diffuse, shader_bsdf_diffuse(kg, &sd));
+ out = kernel_bake_evaluate_direct_indirect(kg, &sd, &state, &shader_bsdf_diffuse, L.direct_diffuse, L.indirect_diffuse, pass_filter);
break;
}
- case SHADER_EVAL_GLOSSY_INDIRECT:
+ case SHADER_EVAL_GLOSSY:
{
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = safe_divide_color(L.indirect_glossy, shader_bsdf_glossy(kg, &sd));
+ out = kernel_bake_evaluate_direct_indirect(kg, &sd, &state, &shader_bsdf_glossy, L.direct_glossy, L.indirect_glossy, pass_filter);
break;
}
- case SHADER_EVAL_TRANSMISSION_INDIRECT:
+ case SHADER_EVAL_TRANSMISSION:
{
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = safe_divide_color(L.indirect_transmission, shader_bsdf_transmission(kg, &sd));
+ out = kernel_bake_evaluate_direct_indirect(kg, &sd, &state, &shader_bsdf_transmission, L.direct_transmission, L.indirect_transmission, pass_filter);
break;
}
- case SHADER_EVAL_SUBSURFACE_INDIRECT:
+ case SHADER_EVAL_SUBSURFACE:
{
#ifdef __SUBSURFACE__
- shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
- out = safe_divide_color(L.indirect_subsurface, shader_bsdf_subsurface(kg, &sd));
+ out = kernel_bake_evaluate_direct_indirect(kg, &sd, &state, &shader_bsdf_subsurface, L.direct_subsurface, L.indirect_subsurface, pass_filter);
#endif
break;
}
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index 77fc16e3d35..da2416bc09b 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -218,14 +218,10 @@ typedef enum ShaderEvalType {
SHADER_EVAL_AO,
SHADER_EVAL_COMBINED,
SHADER_EVAL_SHADOW,
- SHADER_EVAL_DIFFUSE_DIRECT,
- SHADER_EVAL_GLOSSY_DIRECT,
- SHADER_EVAL_TRANSMISSION_DIRECT,
- SHADER_EVAL_SUBSURFACE_DIRECT,
- SHADER_EVAL_DIFFUSE_INDIRECT,
- SHADER_EVAL_GLOSSY_INDIRECT,
- SHADER_EVAL_TRANSMISSION_INDIRECT,
- SHADER_EVAL_SUBSURFACE_INDIRECT,
+ SHADER_EVAL_DIFFUSE,
+ SHADER_EVAL_GLOSSY,
+ SHADER_EVAL_TRANSMISSION,
+ SHADER_EVAL_SUBSURFACE,
/* extra */
SHADER_EVAL_ENVIRONMENT,
@@ -355,6 +351,39 @@ typedef enum PassType {
#define PASS_ALL (~0)
+typedef enum BakePassFilter {
+ BAKE_FILTER_NONE = 0,
+ BAKE_FILTER_DIRECT = (1 << 0),
+ BAKE_FILTER_INDIRECT = (1 << 1),
+ BAKE_FILTER_COLOR = (1 << 2),
+ BAKE_FILTER_DIFFUSE = (1 << 3),
+ BAKE_FILTER_GLOSSY = (1 << 4),
+ BAKE_FILTER_TRANSMISSION = (1 << 5),
+ BAKE_FILTER_SUBSURFACE = (1 << 6),
+ BAKE_FILTER_EMISSION = (1 << 7),
+ BAKE_FILTER_AO = (1 << 8),
+} BakePassFilter;
+
+typedef enum BakePassFilterCombos {
+ BAKE_FILTER_COMBINED = (
+ BAKE_FILTER_DIRECT |
+ BAKE_FILTER_INDIRECT |
+ BAKE_FILTER_DIFFUSE |
+ BAKE_FILTER_GLOSSY |
+ BAKE_FILTER_TRANSMISSION |
+ BAKE_FILTER_SUBSURFACE |
+ BAKE_FILTER_EMISSION |
+ BAKE_FILTER_AO),
+ BAKE_FILTER_DIFFUSE_DIRECT = (BAKE_FILTER_DIRECT | BAKE_FILTER_DIFFUSE),
+ BAKE_FILTER_GLOSSY_DIRECT = (BAKE_FILTER_DIRECT | BAKE_FILTER_GLOSSY),
+ BAKE_FILTER_TRANSMISSION_DIRECT = (BAKE_FILTER_DIRECT | BAKE_FILTER_TRANSMISSION),
+ BAKE_FILTER_SUBSURFACE_DIRECT = (BAKE_FILTER_DIRECT | BAKE_FILTER_SUBSURFACE),
+ BAKE_FILTER_DIFFUSE_INDIRECT = (BAKE_FILTER_INDIRECT | BAKE_FILTER_DIFFUSE),
+ BAKE_FILTER_GLOSSY_INDIRECT = (BAKE_FILTER_INDIRECT | BAKE_FILTER_GLOSSY),
+ BAKE_FILTER_TRANSMISSION_INDIRECT = (BAKE_FILTER_INDIRECT | BAKE_FILTER_TRANSMISSION),
+ BAKE_FILTER_SUBSURFACE_INDIRECT = (BAKE_FILTER_INDIRECT | BAKE_FILTER_SUBSURFACE),
+} BakePassFilterCombos;
+
#ifdef __PASSES__
typedef ccl_addr_space struct PathRadiance {
diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h
index 1ce1e41272b..1a07c705f1c 100644
--- a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h
+++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h
@@ -44,6 +44,7 @@ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg,
float4 *output,
float *output_luma,
int type,
+ int filter,
int i,
int offset,
int sample);
diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
index 0249610b381..1454f925ab9 100644
--- a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
+++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
@@ -101,6 +101,7 @@ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg,
float4 *output,
float *output_luma,
int type,
+ int filter,
int i,
int offset,
int sample)
@@ -111,6 +112,7 @@ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg,
input,
output,
(ShaderEvalType)type,
+ filter,
i,
offset,
sample);
diff --git a/intern/cycles/kernel/kernels/cuda/kernel.cu b/intern/cycles/kernel/kernels/cuda/kernel.cu
index e094612de01..c8940b981bb 100644
--- a/intern/cycles/kernel/kernels/cuda/kernel.cu
+++ b/intern/cycles/kernel/kernels/cuda/kernel.cu
@@ -183,12 +183,12 @@ kernel_cuda_shader(uint4 *input,
extern "C" __global__ void
CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
-kernel_cuda_bake(uint4 *input, float4 *output, int type, int sx, int sw, int offset, int sample)
+kernel_cuda_bake(uint4 *input, float4 *output, int type, int filter, int sx, int sw, int offset, int sample)
{
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
if(x < sx + sw)
- kernel_bake_evaluate(NULL, input, output, (ShaderEvalType)type, x, offset, sample);
+ kernel_bake_evaluate(NULL, input, output, (ShaderEvalType)type, filter, x, offset, sample);
}
#endif
diff --git a/intern/cycles/kernel/kernels/opencl/kernel.cl b/intern/cycles/kernel/kernels/opencl/kernel.cl
index 4c9f7ba1d7c..0a44a4d0301 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel.cl
@@ -99,7 +99,7 @@ __kernel void kernel_ocl_bake(
ccl_global type *name,
#include "../../kernel_textures.h"
- int type, int sx, int sw, int offset, int sample)
+ int type, int filter, int sx, int sw, int offset, int sample)
{
KernelGlobals kglobals, *kg = &kglobals;
@@ -115,7 +115,7 @@ __kernel void kernel_ocl_bake(
#ifdef __NO_BAKING__
output[x] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
#else
- kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, x, offset, sample);
+ kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, filter, x, offset, sample);
#endif
}
}
@@ -174,4 +174,4 @@ __kernel void kernel_ocl_convert_to_half_float(
kernel_film_convert_to_half_float(kg, rgba, buffer, sample_scale, x, y, offset, stride);
}
-#endif // __COMPILE_ONLY_MEGAKERNEL__ \ No newline at end of file
+#endif // __COMPILE_ONLY_MEGAKERNEL__