[Bf-blender-cvs] [d99c513] master: Cycles OpenCL: use AoS rather than SoA storage for ShaderData.

[Brecht Van Lommel]

Commit: d99c513afaf76324406c582304a2d7bf85fbc3a0
Author: Brecht Van Lommel
Date:   Wed Aug 24 01:44:48 2016 +0200
Branches: master
https://developer.blender.org/rBd99c513afaf76324406c582304a2d7bf85fbc3a0

Cycles OpenCL: use AoS rather than SoA storage for ShaderData.

Fixes part of the performance regressions compared to 2.76b, see T49046, T48876.

===================================================================

M	intern/cycles/kernel/kernel_types.h

===================================================================

diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index e299406..3152ec6 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -745,7 +745,7 @@ enum ShaderDataFlag {
 
 #ifdef __SPLIT_KERNEL__
 #  define SD_THREAD (get_global_id(1) * get_global_size(0) + get_global_id(0))
-#  if defined(__SPLIT_KERNEL_AOS__)
+#  if !defined(__SPLIT_KERNEL_SOA__)
      /* ShaderData is stored as an Array-of-Structures */
 #    define ccl_soa_member(type, name) type soa_##name
 #    define ccl_fetch(s, t) (s[SD_THREAD].soa_##t)