Update Linux Renderscript prebuilts for LLVM rebase to r235153

Change-Id: Ib3fcc7f969e42b0ffbc0422b19e79d9e253d8bf0
(cherry picked from commit 37fb08ad7037023182f86faf415c23e88b7cc62e)

35 files changed, 1069 insertions, 597 deletions

diff --git a/renderscript/clang-include/avx2intrin.h b/renderscript/clang-include/avx2intrin.h
index 7485bdd..949195b 100644
--- a/renderscript/clang-include/avx2intrin.h
+++ b/renderscript/clang-include/avx2intrin.h
@@ -160,7 +160,7 @@ _mm256_blendv_epi8(__m256i __V1, __m256i __V2, __m256i __M)
 #define _mm256_blend_epi16(V1, V2, M) __extension__ ({ \
   __m256i __V1 = (V1); \
   __m256i __V2 = (V2); \
-  (__m256d)__builtin_shufflevector((__v16hi)__V1, (__v16hi)__V2, \
+  (__m256i)__builtin_shufflevector((__v16hi)__V1, (__v16hi)__V2, \
                                    (((M) & 0x01) ? 16 : 0), \
                                    (((M) & 0x02) ? 17 : 1), \
                                    (((M) & 0x04) ? 18 : 2), \
diff --git a/renderscript/clang-include/avx512fintrin.h b/renderscript/clang-include/avx512fintrin.h
index c6d46cb..72af281 100644
--- a/renderscript/clang-include/avx512fintrin.h
+++ b/renderscript/clang-include/avx512fintrin.h
@@ -162,6 +162,168 @@ _mm512_castps512_ps128(__m512 __a)
   return __builtin_shufflevector(__a, __a, 0, 1, 2, 3);
 }
 
+/* Bitwise operators */
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_and_epi32(__m512i __a, __m512i __b)
+{
+  return __a & __b;
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_mask_and_epi32(__m512i __src, __mmask16 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_pandd512_mask((__v16si) __a,
+              (__v16si) __b,
+              (__v16si) __src,
+              (__mmask16) __k);
+}
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_maskz_and_epi32(__mmask16 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_pandd512_mask((__v16si) __a,
+              (__v16si) __b,
+              (__v16si)
+              _mm512_setzero_si512 (),
+              (__mmask16) __k);
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_and_epi64(__m512i __a, __m512i __b)
+{
+  return __a & __b;
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_mask_and_epi64(__m512i __src, __mmask8 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_pandq512_mask ((__v8di) __a,
+              (__v8di) __b,
+              (__v8di) __src,
+              (__mmask8) __k);
+}
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_maskz_and_epi64(__mmask8 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_pandq512_mask ((__v8di) __a,
+              (__v8di) __b,
+              (__v8di)
+              _mm512_setzero_si512 (),
+              (__mmask8) __k);
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_or_epi32(__m512i __a, __m512i __b)
+{
+  return __a | __b;
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_mask_or_epi32(__m512i __src, __mmask16 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_pord512_mask((__v16si) __a,
+              (__v16si) __b,
+              (__v16si) __src,
+              (__mmask16) __k);
+}
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_maskz_or_epi32(__mmask16 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_pord512_mask((__v16si) __a,
+              (__v16si) __b,
+              (__v16si)
+              _mm512_setzero_si512 (),
+              (__mmask16) __k);
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_or_epi64(__m512i __a, __m512i __b)
+{
+  return __a | __b;
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_mask_or_epi64(__m512i __src, __mmask8 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_porq512_mask ((__v8di) __a,
+              (__v8di) __b,
+              (__v8di) __src,
+              (__mmask8) __k);
+}
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_maskz_or_epi64(__mmask8 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_porq512_mask ((__v8di) __a,
+              (__v8di) __b,
+              (__v8di)
+              _mm512_setzero_si512 (),
+              (__mmask8) __k);
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_xor_epi32(__m512i __a, __m512i __b)
+{
+  return __a ^ __b;
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_mask_xor_epi32(__m512i __src, __mmask16 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_pxord512_mask((__v16si) __a,
+              (__v16si) __b,
+              (__v16si) __src,
+              (__mmask16) __k);
+}
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_maskz_xor_epi32(__mmask16 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_pxord512_mask((__v16si) __a,
+              (__v16si) __b,
+              (__v16si)
+              _mm512_setzero_si512 (),
+              (__mmask16) __k);
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_xor_epi64(__m512i __a, __m512i __b)
+{
+  return __a ^ __b;
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_mask_xor_epi64(__m512i __src, __mmask8 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_pxorq512_mask ((__v8di) __a,
+              (__v8di) __b,
+              (__v8di) __src,
+              (__mmask8) __k);
+}
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_maskz_xor_epi64(__mmask8 __k, __m512i __a, __m512i __b)
+{
+  return (__m512i) __builtin_ia32_pxorq512_mask ((__v8di) __a,
+              (__v8di) __b,
+              (__v8di)
+              _mm512_setzero_si512 (),
+              (__mmask8) __k);
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_and_si512(__m512i __a, __m512i __b)
+{
+  return __a & __b;
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_or_si512(__m512i __a, __m512i __b)
+{
+  return __a | __b;
+}
+
+static __inline__ __m512i __attribute__((__always_inline__, __nodebug__))
+_mm512_xor_si512(__m512i __a, __m512i __b)
+{
+  return __a ^ __b;
+}
 /* Arithmetic */
 
 static __inline __m512d __attribute__((__always_inline__, __nodebug__))
diff --git a/renderscript/clang-include/htmintrin.h b/renderscript/clang-include/htmintrin.h
index 4598ee0..0088c7c 100644
--- a/renderscript/clang-include/htmintrin.h
+++ b/renderscript/clang-include/htmintrin.h
@@ -128,4 +128,99 @@ typedef uintptr_t tfhar_t;
 
 #endif /* __powerpc */
 
+#ifdef __s390__
+
+/* Condition codes generated by tbegin  */
+#define _HTM_TBEGIN_STARTED       0
+#define _HTM_TBEGIN_INDETERMINATE 1
+#define _HTM_TBEGIN_TRANSIENT     2
+#define _HTM_TBEGIN_PERSISTENT    3
+
+/* The abort codes below this threshold are reserved for machine use.  */
+#define _HTM_FIRST_USER_ABORT_CODE 256
+
+/* The transaction diagnostic block is it is defined in the Principles
+   of Operation chapter 5-91.  */
+
+struct __htm_tdb {
+  unsigned char format;                /*   0 */
+  unsigned char flags;
+  unsigned char reserved1[4];
+  unsigned short nesting_depth;
+  unsigned long long abort_code;       /*   8 */
+  unsigned long long conflict_token;   /*  16 */
+  unsigned long long atia;             /*  24 */
+  unsigned char eaid;                  /*  32 */
+  unsigned char dxc;
+  unsigned char reserved2[2];
+  unsigned int program_int_id;
+  unsigned long long exception_id;     /*  40 */
+  unsigned long long bea;              /*  48 */
+  unsigned char reserved3[72];         /*  56 */
+  unsigned long long gprs[16];         /* 128 */
+} __attribute__((__packed__, __aligned__ (8)));
+
+
+/* Helper intrinsics to retry tbegin in case of transient failure.  */
+
+static __inline int __attribute__((__always_inline__, __nodebug__))
+__builtin_tbegin_retry_null (int retry)
+{
+  int cc, i = 0;
+
+  while ((cc = __builtin_tbegin(0)) == _HTM_TBEGIN_TRANSIENT
+         && i++ < retry)
+    __builtin_tx_assist(i);
+
+  return cc;
+}
+
+static __inline int __attribute__((__always_inline__, __nodebug__))
+__builtin_tbegin_retry_tdb (void *tdb, int retry)
+{
+  int cc, i = 0;
+
+  while ((cc = __builtin_tbegin(tdb)) == _HTM_TBEGIN_TRANSIENT
+         && i++ < retry)
+    __builtin_tx_assist(i);
+
+  return cc;
+}
+
+#define __builtin_tbegin_retry(tdb, retry) \
+  (__builtin_constant_p(tdb == 0) && tdb == 0 ? \
+   __builtin_tbegin_retry_null(retry) : \
+   __builtin_tbegin_retry_tdb(tdb, retry))
+
+static __inline int __attribute__((__always_inline__, __nodebug__))
+__builtin_tbegin_retry_nofloat_null (int retry)
+{
+  int cc, i = 0;
+
+  while ((cc = __builtin_tbegin_nofloat(0)) == _HTM_TBEGIN_TRANSIENT
+         && i++ < retry)
+    __builtin_tx_assist(i);
+
+  return cc;
+}
+
+static __inline int __attribute__((__always_inline__, __nodebug__))
+__builtin_tbegin_retry_nofloat_tdb (void *tdb, int retry)
+{
+  int cc, i = 0;
+
+  while ((cc = __builtin_tbegin_nofloat(tdb)) == _HTM_TBEGIN_TRANSIENT
+         && i++ < retry)
+    __builtin_tx_assist(i);
+
+  return cc;
+}
+
+#define __builtin_tbegin_retry_nofloat(tdb, retry) \
+  (__builtin_constant_p(tdb == 0) && tdb == 0 ? \
+   __builtin_tbegin_retry_nofloat_null(retry) : \
+   __builtin_tbegin_retry_nofloat_tdb(tdb, retry))
+
+#endif /* __s390__ */
+
 #endif /* __HTMINTRIN_H */
diff --git a/renderscript/clang-include/htmxlintrin.h b/renderscript/clang-include/htmxlintrin.h
index 8791afe..30f524d 100644
--- a/renderscript/clang-include/htmxlintrin.h
+++ b/renderscript/clang-include/htmxlintrin.h
@@ -212,4 +212,152 @@ __TM_failure_code(void* const TM_buff)
 
 #endif /* __powerpc__ */
 
+#ifdef __s390__
+
+#include <stdint.h>
+
+/* These intrinsics are being made available for compatibility with
+   the IBM XL compiler.  For documentation please see the "z/OS XL
+   C/C++ Programming Guide" publically available on the web.  */
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_simple_begin ()
+{
+  return __builtin_tbegin_nofloat (0);
+}
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_begin (void* const tdb)
+{
+  return __builtin_tbegin_nofloat (tdb);
+}
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_end ()
+{
+  return __builtin_tend ();
+}
+
+static __inline void __attribute__((__always_inline__))
+__TM_abort ()
+{
+  return __builtin_tabort (_HTM_FIRST_USER_ABORT_CODE);
+}
+
+static __inline void __attribute__((__always_inline__, __nodebug__))
+__TM_named_abort (unsigned char const code)
+{
+  return __builtin_tabort ((int)_HTM_FIRST_USER_ABORT_CODE + code);
+}
+
+static __inline void __attribute__((__always_inline__, __nodebug__))
+__TM_non_transactional_store (void* const addr, long long const value)
+{
+  __builtin_non_tx_store ((uint64_t*)addr, (uint64_t)value);
+}
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_nesting_depth (void* const tdb_ptr)
+{
+  int depth = __builtin_tx_nesting_depth ();
+  struct __htm_tdb *tdb = (struct __htm_tdb*)tdb_ptr;
+
+  if (depth != 0)
+    return depth;
+
+  if (tdb->format != 1)
+    return 0;
+  return tdb->nesting_depth;
+}
+
+/* Transaction failure diagnostics */
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_is_user_abort (void* const tdb_ptr)
+{
+  struct __htm_tdb *tdb = (struct __htm_tdb*)tdb_ptr;
+
+  if (tdb->format != 1)
+    return 0;
+
+  return !!(tdb->abort_code >= _HTM_FIRST_USER_ABORT_CODE);
+}
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_is_named_user_abort (void* const tdb_ptr, unsigned char* code)
+{
+  struct __htm_tdb *tdb = (struct __htm_tdb*)tdb_ptr;
+
+  if (tdb->format != 1)
+    return 0;
+
+  if (tdb->abort_code >= _HTM_FIRST_USER_ABORT_CODE)
+    {
+      *code = tdb->abort_code - _HTM_FIRST_USER_ABORT_CODE;
+      return 1;
+    }
+  return 0;
+}
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_is_illegal (void* const tdb_ptr)
+{
+  struct __htm_tdb *tdb = (struct __htm_tdb*)tdb_ptr;
+
+  return (tdb->format == 1
+	  && (tdb->abort_code == 4 /* unfiltered program interruption */
+	      || tdb->abort_code == 11 /* restricted instruction */));
+}
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_is_footprint_exceeded (void* const tdb_ptr)
+{
+  struct __htm_tdb *tdb = (struct __htm_tdb*)tdb_ptr;
+
+  return (tdb->format == 1
+	  && (tdb->abort_code == 7 /* fetch overflow */
+	      || tdb->abort_code == 8 /* store overflow */));
+}
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_is_nested_too_deep (void* const tdb_ptr)
+{
+  struct __htm_tdb *tdb = (struct __htm_tdb*)tdb_ptr;
+
+  return tdb->format == 1 && tdb->abort_code == 13; /* depth exceeded */
+}
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_is_conflict (void* const tdb_ptr)
+{
+  struct __htm_tdb *tdb = (struct __htm_tdb*)tdb_ptr;
+
+  return (tdb->format == 1
+	  && (tdb->abort_code == 9 /* fetch conflict */
+	      || tdb->abort_code == 10 /* store conflict */));
+}
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_is_failure_persistent (long const result)
+{
+  return result == _HTM_TBEGIN_PERSISTENT;
+}
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_failure_address (void* const tdb_ptr)
+{
+  struct __htm_tdb *tdb = (struct __htm_tdb*)tdb_ptr;
+  return tdb->atia;
+}
+
+static __inline long __attribute__((__always_inline__, __nodebug__))
+__TM_failure_code (void* const tdb_ptr)
+{
+  struct __htm_tdb *tdb = (struct __htm_tdb*)tdb_ptr;
+
+  return tdb->abort_code;
+}
+
+#endif /* __s390__ */
+
 #endif /* __HTMXLINTRIN_H  */
diff --git a/renderscript/clang-include/s390intrin.h b/renderscript/clang-include/s390intrin.h
new file mode 100644
index 0000000..b209895
--- /dev/null
+++ b/renderscript/clang-include/s390intrin.h
@@ -0,0 +1,35 @@
+/*===---- s390intrin.h - SystemZ intrinsics --------------------------------===
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#ifndef __S390INTRIN_H
+#define __S390INTRIN_H
+
+#ifndef __s390__
+#error "<s390intrin.h> is for s390 only"
+#endif
+
+#ifdef __HTM__
+#include <htmintrin.h>
+#endif
+
+#endif /* __S390INTRIN_H*/
diff --git a/renderscript/include/rs_allocation_data.rsh b/renderscript/include/rs_allocation_data.rsh
index 928b649..d8785a2 100644
--- a/renderscript/include/rs_allocation_data.rsh
+++ b/renderscript/include/rs_allocation_data.rsh
@@ -113,16 +113,16 @@ extern void __attribute__((overloadable))
  * values incrementing other dimensions.  Use multiple calls to rsGetElementAt()
  * instead.
  *
- * For example, when calling rsAllocationVLoadX_int4(a, 20, 30),
- * an int4 composed of a[20, 30], a[21, 30], a[22, 30], and a[23, 30] is returned.
+ * For example, when calling rsAllocationVLoadX_int4(a, 20, 30), an int4 composed
+ * of a[20, 30], a[21, 30], a[22, 30], and a[23, 30] is returned.
  *
- * When retrieving from a three dimensional allocations, use the x, y, z
- * variant.   Similarly, use the x, y variant for two dimensional
- * allocations and x for the mono dimensional allocations.
+ * When retrieving from a three dimensional allocations, use the x, y, z variant.
+ * Similarly, use the x, y variant for two dimensional allocations and x for the
+ * mono dimensional allocations.
  *
- * For efficiency, this function does not validate the inputs.  Trying to
- * wrap the X index, exceeding the size of the allocation, or using indexes
- * incompatible with the dimensionality of the allocation yields undefined results.
+ * For efficiency, this function does not validate the inputs.  Trying to wrap
+ * the X index, exceeding the size of the allocation, or using indices incompatible
+ * with the dimensionality of the allocation yields undefined results.
  *
  * See also rsAllocationVStoreX().
  *
@@ -585,24 +585,23 @@ extern ulong4 __attribute__((overloadable))
 /*
  * rsAllocationVStoreX: Store a vector into an allocation of scalars
  *
- * This function stores the entries of a vector into successive cells of an
- * allocation.  It assumes that the allocation contains scalars.
+ * This function stores the entries of a vector into successive cells of an allocation.
+ * It assumes that the allocation contains scalars.
  *
- * The "X" in the name indicates that successive values are stored by
- * increasing the X index.  There are currently no functions to store successive
- * values incrementing other dimensions.  Use multiple calls to rsSetElementAt()
- * instead.
+ * The "X" in the name indicates that successive values are stored by increasing
+ * the X index.  There are currently no functions to store successive values
+ * incrementing other dimensions.  Use multiple calls to rsSetElementAt() instead.
  *
- * For example, when calling rsAllocationVStoreX_int3(a, v, 20, 30),
- * v.x is stored at a[20, 30], v.y at a[21, 30], and v.z at a[22, 30].
+ * For example, when calling rsAllocationVStoreX_int3(a, v, 20, 30), v.x is stored
+ * at a[20, 30], v.y at a[21, 30], and v.z at a[22, 30].
  *
- * When storing into a three dimensional allocations, use the x, y, z
- * variant.   Similarly, use the x, y variant for two dimensional
- * allocations and x for the mono dimensional allocations.
+ * When storing into a three dimensional allocations, use the x, y, z variant.
+ * Similarly, use the x, y variant for two dimensional allocations and x for the
+ * mono dimensional allocations.
  *
- * For efficiency, this function does not validate the inputs.  Trying to
- * wrap the X index, exceeding the size of the allocation, or using indexes
- * incompatible with the dimensionality of the allocation yiels undefined results.
+ * For efficiency, this function does not validate the inputs.  Trying to wrap the
+ * X index, exceeding the size of the allocation, or using indexes incompatible
+ * with the dimensionality of the allocation yiels undefined results.
  *
  * See also rsAllocationVLoadX().
  *
@@ -1068,14 +1067,13 @@ extern void __attribute__((overloadable))
  *
  * This function extracts a single cell from an allocation.
  *
- * When retrieving from a three dimensional allocations, use the x, y, z
- * variant.   Similarly, use the x, y variant for two dimensional
- * allocations and x for the mono dimensional allocations.
+ * When retrieving from a three dimensional allocations, use the x, y, z variant.
+ * Similarly, use the x, y variant for two dimensional allocations and x for the
+ * mono dimensional allocations.
  *
- * This function has two styles.  One returns the address of the value using a
- * void*, the other returns the actual value, e.g. rsGetElementAt() vs.
- * rsGetElementAt_int4().  For primitive types, always use the latter as it is
- * more efficient.
+ * This function has two styles.  One returns the address of the value using a void*,
+ * the other returns the actual value, e.g. rsGetElementAt() vs. rsGetElementAt_int4().
+ * For primitive types, always use the latter as it is more efficient.
  */
 extern const void* __attribute__((overloadable))
     rsGetElementAt(rs_allocation a, uint32_t x);
@@ -2615,13 +2613,13 @@ extern float4 __attribute__((overloadable))
  *
  * This function stores a value into a single cell of an allocation.
  *
- * When storing into a three dimensional allocations, use the x, y, z
- * variant.   Similarly, use the x, y variant for two dimensional
- * allocations and x for the mono dimensional allocations.
+ * When storing into a three dimensional allocations, use the x, y, z variant.
+ * Similarly, use the x, y variant for two dimensional allocations and x for
+ * the mono dimensional allocations.
  *
- * This function has two styles.  One passes the value to be stored using
- * a void*, the other has the actual value as an argument, e.g. rsSetElementAt()
- * vs. rsSetElementAt_int4().  For primitive types, always use the latter as it is
+ * This function has two styles.  One passes the value to be stored using a void*,
+ * the other has the actual value as an argument, e.g. rsSetElementAt() vs.
+ * rsSetElementAt_int4().  For primitive types, always use the latter as it is
  * more efficient.
  *
  * See also rsGetElementAt().
diff --git a/renderscript/include/rs_atomic.rsh b/renderscript/include/rs_atomic.rsh
index cc2b8d5..0ab26ca 100644
--- a/renderscript/include/rs_atomic.rsh
+++ b/renderscript/include/rs_atomic.rsh
@@ -89,9 +89,9 @@ extern int32_t __attribute__((overloadable))
  * by rsAtomicCas() is compareValue.
  *
  * Parameters:
- *   addr: The address of the value to compare and replace if the test passes.
- *   compareValue: The value to test *addr against.
- *   newValue: The value to write if the test passes.
+ *   addr: Address of the value to compare and replace if the test passes.
+ *   compareValue: Value to test *addr against.
+ *   newValue: Value to write if the test passes.
  *
  * Returns: Value of *addr prior to the operation.
  */
diff --git a/renderscript/include/rs_convert.rsh b/renderscript/include/rs_convert.rsh
index 7abd557..7cf6c50 100644
--- a/renderscript/include/rs_convert.rsh
+++ b/renderscript/include/rs_convert.rsh
@@ -19,8 +19,8 @@
 /*
  * rs_convert.rsh: Conversion Functions
  *
- * The functions below convert from a numerical vector type to another,
- * of from one color representation to another.
+ * The functions below convert from a numerical vector type to another, of from one color
+ * representation to another.
  */
 
 #ifndef RENDERSCRIPT_RS_CONVERT_RSH
@@ -29,10 +29,9 @@
 /*
  * convert: Convert numerical vectors
  *
- * Converts a vector from one numerical type to another.  The conversion are
- * done entry per entry.
+ * Converts a vector from one numerical type to another.  The conversion are done entry per entry.
  *
- * E.g calling a = convert_short3(b) is equivalent to doing
+ * E.g calling a = convert_short3(b); is equivalent to doing
  * a.x = (short)b.x; a.y = (short)b.y; a.z = (short)b.z;.
  *
  * Converting floating point values to integer types truncates.
@@ -1250,13 +1249,12 @@ extern ulong4 __attribute__((const, overloadable))
 /*
  * rsPackColorTo8888: Create a uchar4 RGBA from floats
  *
- * Packs three or four floating point RGBA values into a uchar4.  The RGBA values should
- * be between 0.0 and 1.0 inclusive.  Values outside of this range are clamped to
- * this range.  However numbers greater than INT_MAX or less than INT_MIN can result
- * in undefined behavior.
+ * Packs three or four floating point RGBA values into a uchar4.  The RGBA values should be
+ * between 0.0 and 1.0 inclusive.  Values outside of this range are clamped to this range.
+ * However numbers greater than INT_MAX or less than INT_MIN can result in undefined behavior.
  *
- * If the alpha component is not specified, it is assumed to be 1.0, i.e. the
- * result will have an alpha set to 255.
+ * If the alpha component is not specified, it is assumed to be 1.0, i.e. the result will
+ * have an alpha set to 255.
  *
  * Parameters:
  *   r: Red component.
@@ -1280,8 +1278,7 @@ extern uchar4 __attribute__((const, overloadable))
 /*
  * rsUnpackColor8888: Create a float4 RGBA from uchar4
  *
- * Unpacks a uchar4 color to float4.  The resulting floats will be between 0.0 and
- * 1.0 inclusive.
+ * Unpacks a uchar4 color to float4.  The resulting floats will be between 0.0 and 1.0 inclusive.
  */
 extern float4 __attribute__((const))
     rsUnpackColor8888(uchar4 c);
@@ -1294,9 +1291,9 @@ extern float4 __attribute__((const))
  * We currently don't provide a function to do the reverse conversion.
  *
  * Parameters:
- *   y: Luminance component
- *   u: U chrominance component
- *   v: V chrominance component
+ *   y: Luminance component.
+ *   u: U chrominance component.
+ *   v: V chrominance component.
  */
 extern float4 __attribute__((const, overloadable))
     rsYuvToRGBA_float4(uchar y, uchar u, uchar v);
diff --git a/renderscript/include/rs_core.rsh b/renderscript/include/rs_core.rsh
index be9f7fa..128c689 100644
--- a/renderscript/include/rs_core.rsh
+++ b/renderscript/include/rs_core.rsh
@@ -19,19 +19,19 @@
 /*
  * rs_core.rsh: Overview
  *
- * RenderScript is a high-performance runtime that provides
- * compute operations at the native level. RenderScript code is compiled on devices
- * at runtime to allow platform-independence as well.
- * This reference documentation describes the RenderScript runtime APIs, which you
- * can utilize to write RenderScript code in C99. The RenderScript compute header
- * files are automatically included for you.
+ * RenderScript is a high-performance runtime that provides compute operations at the native level.
+ * RenderScript code is compiled on devices at runtime to allow platform-independence as well.
  *
- * To use RenderScript, you need to utilize the RenderScript runtime APIs documented here
- * as well as the Android framework APIs for RenderScript.
- * For documentation on the Android framework APIs, see the android.renderscript package reference.
- * For more information on how to develop with RenderScript and how the runtime and
- * Android framework APIs interact, see the RenderScript developer guide
- * and the RenderScript samples.
+ * This reference documentation describes the RenderScript runtime APIs, which you can utilize
+ * to write RenderScript code in C99. The RenderScript compute header files are automatically
+ * included for you.
+ *
+ * To use RenderScript, you need to utilize the RenderScript runtime APIs documented here as well
+ * as the Android framework APIs for RenderScript.  For documentation on the Android framework
+ * APIs, see the android.renderscript package reference.
+ *
+ * For more information on how to develop with RenderScript and how the runtime and Android
+ * framework APIs interact, see the RenderScript developer guide and the RenderScript samples.
  */
 
 #ifndef RENDERSCRIPT_RS_CORE_RSH
diff --git a/renderscript/include/rs_for_each.rsh b/renderscript/include/rs_for_each.rsh
index f08a4e3..1bbcf9f 100644
--- a/renderscript/include/rs_for_each.rsh
+++ b/renderscript/include/rs_for_each.rsh
@@ -57,8 +57,8 @@ typedef enum rs_for_each_strategy {
  * A kernel may be executed in parallel over multiple threads.  Each thread will have its
  * own context.
  *
- * You can access the context by adding a rs_kernel_context argument to your
- * kernel function.  See rsGetDimX() and rsGetArray0() for examples.
+ * You can access the context by adding a rs_kernel_context argument to your kernel
+ * function.  See rsGetDimX() and rsGetArray0() for examples.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 23))
 typedef const struct rs_kernel_context_t * rs_kernel_context;
@@ -85,6 +85,12 @@ typedef struct rs_script_call {
     uint32_t zEnd; // Ending index (exclusive) in the Z dimension.
     uint32_t arrayStart; // Starting index in the Array0 dimension.
     uint32_t arrayEnd; // Ending index (exclusive) in the Array0 dimension.
+    uint32_t array1Start; // Starting index in the Array1 dimension.
+    uint32_t array1End; // Ending index (exclusive) in the Array1 dimension.
+    uint32_t array2Start; // Starting index in the Array2 dimension.
+    uint32_t array2End; // Ending index (exclusive) in the Array2 dimension.
+    uint32_t array3Start; // Starting index in the Array3 dimension.
+    uint32_t array3End; // Ending index (exclusive) in the Array3 dimension.
 } rs_script_call_t;
 
 /*
@@ -115,7 +121,7 @@ typedef struct rs_script_call {
  *   output: Allocation to write date into.
  *   usrData: User defined data to pass to the script.  May be NULL.
  *   sc: Extra control information used to select a sub-region of the allocation to be processed or suggest a walking strategy.  May be NULL.
- *   usrDataLen: The size of the userData structure.  This will be used to perform a shallow copy of the data if necessary.
+ *   usrDataLen: Size of the userData structure.  This will be used to perform a shallow copy of the data if necessary.
  */
 #if !defined(RS_VERSION) || (RS_VERSION <= 13)
 extern void __attribute__((overloadable))
@@ -148,15 +154,15 @@ extern void __attribute__((overloadable))
 /*
  * rsGetArray0: Index in the Array0 dimension for the specified context
  *
- * Returns the index in the Array0 dimension of the cell being processed,
- * as specified by the supplied context.
+ * Returns the index in the Array0 dimension of the cell being processed, as specified
+ * by the supplied context.
  *
- * This context is created when a kernel is launched and updated at each
- * iteration.  It contains common characteristics of the allocations being
- * iterated over and rarely used indexes, like the Array0 index.
+ * This context is created when a kernel is launched and updated at each iteration.
+ * It contains common characteristics of the allocations being iterated over and rarely
+ * used indexes, like the Array0 index.
  *
- * You can access the context by adding a rs_kernel_context argument to your
- * kernel function.  E.g.
+ * You can access the context by adding a rs_kernel_context argument to your kernel
+ * function.  E.g.
  * short RS_KERNEL myKernel(short value, uint32_t x, rs_kernel_context context) {
  *   // The current index in the common x, y, z, w dimensions are accessed by
  *   // adding these variables as arguments.  For the more rarely used indexes
@@ -175,9 +181,8 @@ extern uint32_t __attribute__((overloadable))
 /*
  * rsGetArray1: Index in the Array1 dimension for the specified context
  *
- * Returns the index in the Array1 dimension of the cell being processed,
- * as specified by the supplied context.  See rsGetArray0() for an explanation
- * of the context.
+ * Returns the index in the Array1 dimension of the cell being processed, as specified
+ * by the supplied context.  See rsGetArray0() for an explanation of the context.
  *
  * Returns 0 if the Array1 dimension is not present.
  */
@@ -203,9 +208,8 @@ extern uint32_t __attribute__((overloadable))
 /*
  * rsGetArray3: Index in the Array3 dimension for the specified context
  *
- * Returns the index in the Array3 dimension of the cell being processed,
- * as specified by the supplied context.  See rsGetArray0() for an explanation
- * of the context.
+ * Returns the index in the Array3 dimension of the cell being processed, as specified
+ * by the supplied context.  See rsGetArray0() for an explanation of the context.
  *
  * Returns 0 if the Array3 dimension is not present.
  */
@@ -269,9 +273,9 @@ extern uint32_t __attribute__((overloadable))
 /*
  * rsGetDimHasFaces: Presence of more than one face for the specified context
  *
- * If the context refers to a cubemap, this function returns true if there's
- * more than one face present.  In all other cases, it returns false.
- * See rsGetDimX() for an explanation of the context.
+ * If the context refers to a cubemap, this function returns true if there's more than
+ * one face present.  In all other cases, it returns false.  See rsGetDimX() for an
+ * explanation of the context.
  *
  * rsAllocationGetDimFaces() is similar but returns 0 or 1 instead of a bool.
  *
@@ -285,8 +289,9 @@ extern bool __attribute__((overloadable))
 /*
  * rsGetDimLod: Number of levels of detail for the specified context
  *
- * Returns the number of levels of detail for the specified context.
- * This is useful for mipmaps.  See rsGetDimX() for an explanation of the context.
+ * Returns the number of levels of detail for the specified context.  This is useful
+ * for mipmaps.  See rsGetDimX() for an explanation of the context.
+ *
  * Returns 0 if Level of Detail is not used.
  *
  * rsAllocationGetDimLOD() is similar but returns 0 or 1 instead the actual
@@ -351,8 +356,8 @@ extern uint32_t __attribute__((overloadable))
 /*
  * rsGetFace: Coordinate of the Face for the specified context
  *
- * Returns the face on which the cell being processed is found, as specified
- * by the supplied context.  See rsGetArray0() for an explanation of the context.
+ * Returns the face on which the cell being processed is found, as specified by the
+ * supplied context.  See rsGetArray0() for an explanation of the context.
  *
  * Returns RS_ALLOCATION_CUBEMAP_FACE_POSITIVE_X if the face dimension is not
  * present.
@@ -365,9 +370,9 @@ extern rs_allocation_cubemap_face __attribute__((overloadable))
 /*
  * rsGetLod: Index in the Levels of Detail dimension for the specified context
  *
- * Returns the index in the Levels of Detail dimension of the cell being
- * processed, as specified by the supplied context.  See rsGetArray0() for
- * an explanation of the context.
+ * Returns the index in the Levels of Detail dimension of the cell being processed,
+ * as specified by the supplied context.  See rsGetArray0() for an explanation of
+ * the context.
  *
  * Returns 0 if the Levels of Detail dimension is not present.
  */
diff --git a/renderscript/include/rs_graphics.rsh b/renderscript/include/rs_graphics.rsh
index c133aad..5439f06 100644
--- a/renderscript/include/rs_graphics.rsh
+++ b/renderscript/include/rs_graphics.rsh
@@ -214,12 +214,10 @@ typedef _RS_HANDLE rs_program_store;
 /*
  * rsClearObject: Release an object
  *
- * Tells the run time that this handle will no longer be used to access the
- * the related object.  If this was the last handle to that object, resource
- * recovery may happen.
+ * Tells the run time that this handle will no longer be used to access the the related
+ * object.  If this was the last handle to that object, resource recovery may happen.
  *
- * After calling this function, *dst will be set to an empty handle.  See
- * rsIsObject().
+ * After calling this function, *dst will be set to an empty handle.  See rsIsObject().
  */
 #ifndef __LP64__
 extern void __attribute__((overloadable))
@@ -259,9 +257,8 @@ extern void __attribute__((overloadable))
  * This function does not validate that the internal pointer used in the handle
  * points to an actual valid object; it only checks for null.
  *
- * This function can be used to check the element returned by
- * rsElementGetSubElement() or see if rsClearObject() has been called on a
- * handle.
+ * This function can be used to check the Element returned by rsElementGetSubElement()
+ * or see if rsClearObject() has been called on a handle.
  */
 #ifndef __LP64__
 extern bool __attribute__((overloadable))
diff --git a/renderscript/include/rs_math.rsh b/renderscript/include/rs_math.rsh
index 0f094f6..7ab6975 100644
--- a/renderscript/include/rs_math.rsh
+++ b/renderscript/include/rs_math.rsh
@@ -19,12 +19,10 @@
 /*
  * rs_math.rsh: Mathematical Constants and Functions
  *
- * The mathematical functions below can be applied to scalars and vectors.
- * When applied to vectors, a vector of the function applied to each entry
- * of the input is returned.
+ * The mathematical functions below can be applied to scalars and vectors.   When applied
+ * to vectors, the returned value is a vector of the function applied to each entry of the input.
  *
  * For example:
- *
  * float3 a, b;
  * // The following call sets
  * //   a.x to sin(b.x),
@@ -33,17 +31,22 @@
  * a = sin(b);
  *
  *
- * See "Vector math functions" for functions like distance() and length()
- * that interpret instead the input as a single vector in n-dimensional space.
+ * See Vector Math Functions for functions like distance() and length() that interpret
+ * instead the input as a single vector in n-dimensional space.
  *
- * The precision of the mathematical operations is affected by the pragmas
- * rs_fp_relaxed and rs_fp_full.
+ * The precision of the mathematical operations on 32 bit floats is affected by the pragmas
+ * rs_fp_relaxed and rs_fp_full.  Under rs_fp_relaxed, subnormal values may be flushed to zero and
+ * rounding may be done towards zero.  In comparison, rs_fp_full requires correct handling of
+ * subnormal values, i.e. smaller than 1.17549435e-38f.  rs_fp_rull also requires round to nearest
+ * with ties to even.
  *
- * Different precision/speed tradeoffs can be achieved by using three variants
- * of common math functions.  Functions with a name starting with
- * - native_ may have custom hardware implementations with weaker precision,
- * - half_ may perform internal computations using 16 bit floats, and
- * - fast_ are n-dimensional space computations that may use 16 bit floats.
+ * Different precision/speed tradeoffs can be achieved by using variants of the common math
+ * functions.  Functions with a name starting with
+ * - native_: May have custom hardware implementations with weaker precision.  Additionally,
+ *   subnormal values may be flushed to zero, rounding towards zero may be used, and NaN and
+ *   infinity input may not be handled correctly.
+ * - half_: May perform internal computations using 16 bit floats.  Additionally, subnormal
+ *   values may be flushed to zero, and rounding towards zero may be used.
  *
  */
 
@@ -338,8 +341,8 @@ extern float4 __attribute__((const, overloadable))
  * See also native_atan2().
  *
  * Parameters:
- *   numerator: The numerator
- *   denominator: The denominator.  Can be 0.
+ *   numerator: Numerator.
+ *   denominator: Denominator.  Can be 0.
  */
 extern float __attribute__((const, overloadable))
     atan2(float numerator, float denominator);
@@ -363,8 +366,8 @@ extern float4 __attribute__((const, overloadable))
  * See also native_atan2pi().
  *
  * Parameters:
- *   numerator: The numerator
- *   denominator: The denominator.  Can be 0.
+ *   numerator: Numerator.
+ *   denominator: Denominator.  Can be 0.
  */
 extern float __attribute__((const, overloadable))
     atan2pi(float numerator, float denominator);
@@ -1125,10 +1128,9 @@ extern float4 __attribute__((const, overloadable))
  *
  * Multiply and add.  Returns (multiplicand1 * multiplicand2) + offset.
  *
- * This function is similar to mad().  fma() retains full precision of the
- * multiplied result and rounds only after the addition.  mad() rounds after the
- * multiplication and the addition.  This extra precision is not guaranteed in
- * rs_fp_relaxed mode.
+ * This function is similar to mad().  fma() retains full precision of the multiplied result
+ * and rounds only after the addition.  mad() rounds after the multiplication and the addition.
+ * This extra precision is not guaranteed in rs_fp_relaxed mode.
  */
 extern float __attribute__((const, overloadable))
     fma(float multiplicand1, float multiplicand2, float offset);
@@ -1406,8 +1408,7 @@ extern float4 __attribute__((const, overloadable))
  *
  * For example, ilogb(8.5f) returns 3.
  *
- * Because of the difference in mantissa, this number is one less than
- * is returned by frexp().
+ * Because of the difference in mantissa, this number is one less than is returned by frexp().
  *
  * logb() is similar but returns a float.
  */
@@ -1432,8 +1433,8 @@ extern int4 __attribute__((const, overloadable))
  * See frexp() for the reverse operation.
  *
  * Parameters:
- *   mantissa: The mantissa
- *   exponent: The exponent, a single component or matching vector.
+ *   mantissa: Mantissa.
+ *   exponent: Exponent, a single component or matching vector.
  */
 extern float __attribute__((const, overloadable))
     ldexp(float mantissa, int exponent);
@@ -1575,8 +1576,7 @@ extern float4 __attribute__((const, overloadable))
  *
  * For example, logb(8.5f) returns 3.f.
  *
- * Because of the difference in mantissa, this number is one less than
- * is returned by frexp().
+ * Because of the difference in mantissa, this number is one less than is returned by frexp().
  *
  * ilogb() is similar but returns an integer.
  */
@@ -1597,10 +1597,9 @@ extern float4 __attribute__((const, overloadable))
  *
  * Multiply and add.  Returns (multiplicand1 * multiplicand2) + offset.
  *
- * This function is similar to fma().  fma() retains full precision of the
- * multiplied result and rounds only after the addition.  mad() rounds after the
- * multiplication and the addition.  In rs_fp_relaxed mode, mad() may not do the
- * rounding after multiplicaiton.
+ * This function is similar to fma().  fma() retains full precision of the multiplied result
+ * and rounds only after the addition.  mad() rounds after the multiplication and the addition.
+ * In rs_fp_relaxed mode, mad() may not do the rounding after multiplicaiton.
  */
 extern float __attribute__((const, overloadable))
     mad(float multiplicand1, float multiplicand2, float offset);
@@ -2453,7 +2452,8 @@ extern ulong4 __attribute__((const, overloadable))
  *
  * Returns start + ((stop - start) * fraction).
  *
- * This can be useful for mixing two values.  For example, to create a new color that is 40% color1 and 60% color2, use mix(color1, color2, 0.6f).
+ * This can be useful for mixing two values.  For example, to create a new color that is
+ * 40% color1 and 60% color2, use mix(color1, color2, 0.6f).
  */
 extern float __attribute__((const, overloadable))
     mix(float start, float stop, float fraction);
@@ -2481,13 +2481,14 @@ extern float4 __attribute__((const, overloadable))
  *
  * Returns the integral and fractional components of a number.
  *
- * Both components will have the same sign as x.  For example, for an input of -3.72f, iret will be set to -3.f and .72f will be returned.
+ * Both components will have the same sign as x.  For example, for an input of -3.72f,
+ * iret will be set to -3.f and .72f will be returned.
  *
  * Parameters:
- *   v: Source value
+ *   v: Source value.
  *   integral_part: *integral_part will be set to the integral portion of the number.
  *
- * Returns: The floating point portion of the value.
+ * Returns: Floating point portion of the value.
  */
 extern float __attribute__((overloadable))
     modf(float v, float* integral_part);
@@ -2517,8 +2518,7 @@ extern float __attribute__((const, overloadable))
  *
  * Returns the approximate inverse cosine, in radians.
  *
- * This function yields undefined results from input values less than -1 or greater
- * than 1.
+ * This function yields undefined results from input values less than -1 or greater than 1.
  *
  * See also acos().
  */
@@ -2576,8 +2576,7 @@ extern float4 __attribute__((const, overloadable))
  *
  * To get an inverse cosine measured in degrees, use acospi(a) * 180.f.
  *
- * This function yields undefined results from input values less than -1 or greater
- * than 1.
+ * This function yields undefined results from input values less than -1 or greater than 1.
  *
  * See also acospi().
  */
@@ -2606,8 +2605,7 @@ extern float4 __attribute__((const, overloadable))
  *
  * Returns the approximate inverse sine, in radians.
  *
- * This function yields undefined results from input values less than -1 or greater
- * than 1.
+ * This function yields undefined results from input values less than -1 or greater than 1.
  *
  * See also asin().
  */
@@ -2665,8 +2663,7 @@ extern float4 __attribute__((const, overloadable))
  *
  * To get an inverse sine measured in degrees, use asinpi(a) * 180.f.
  *
- * This function yields undefined results from input values less than -1 or greater
- * than 1.
+ * This function yields undefined results from input values less than -1 or greater than 1.
  *
  * See also asinpi().
  */
@@ -2725,8 +2722,8 @@ extern float4 __attribute__((const, overloadable))
  * See also atan2().
  *
  * Parameters:
- *   numerator: The numerator
- *   denominator: The denominator.  Can be 0.
+ *   numerator: Numerator.
+ *   denominator: Denominator.  Can be 0.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 21))
 extern float __attribute__((const, overloadable))
@@ -2751,15 +2748,16 @@ extern float4 __attribute__((const, overloadable))
 /*
  * native_atan2pi: Approximate inverse tangent of a ratio, divided by pi
  *
- * Returns the approximate inverse tangent of (numerator / denominator), in radians, divided by pi.
+ * Returns the approximate inverse tangent of (numerator / denominator),
+ * in radians, divided by pi.
  *
  * To get an inverse tangent measured in degrees, use atan2pi(n, d) * 180.f.
  *
  * See also atan2pi().
  *
  * Parameters:
- *   numerator: The numerator
- *   denominator: The denominator.  Can be 0.
+ *   numerator: Numerator.
+ *   denominator: Denominator.  Can be 0.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 21))
 extern float __attribute__((const, overloadable))
@@ -2977,7 +2975,8 @@ extern float4 __attribute__((const, overloadable))
  *
  * Fast approximate exp.
  *
- * It is valid for inputs from -86.f to 86.f.  The precision is no worse than what would be expected from using 16 bit floating point values.
+ * It is valid for inputs from -86.f to 86.f.  The precision is no worse than what would be
+ * expected from using 16 bit floating point values.
  *
  * See also exp().
  */
@@ -3006,7 +3005,8 @@ extern float4 __attribute__((const, overloadable))
  *
  * Fast approximate exp10.
  *
- * It is valid for inputs from -37.f to 37.f.  The precision is no worse than what would be expected from using 16 bit floating point values.
+ * It is valid for inputs from -37.f to 37.f.  The precision is no worse than what would be
+ * expected from using 16 bit floating point values.
  *
  * See also exp10().
  */
@@ -3035,7 +3035,8 @@ extern float4 __attribute__((const, overloadable))
  *
  * Fast approximate exp2.
  *
- * It is valid for inputs from -125.f to 125.f.  The precision is no worse than what would be expected from using 16 bit floating point values.
+ * It is valid for inputs from -125.f to 125.f.  The precision is no worse than what would be
+ * expected from using 16 bit floating point values.
  *
  * See also exp2().
  */
@@ -3374,10 +3375,10 @@ extern float4 __attribute__((const, overloadable))
  * See also sincos().
  *
  * Parameters:
- *   v: The incoming value in radians.
+ *   v: Incoming value in radians.
  *   cos: *cos will be set to the cosine value.
  *
- * Returns: sine
+ * Returns: Sine.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 21))
 extern float __attribute__((overloadable))
@@ -3568,9 +3569,8 @@ extern float4 __attribute__((const, overloadable))
  *
  * Returns the next representable floating point number from v towards target.
  *
- * In rs_fp_relaxed mode, a denormalized input value may not yield the next
- * denormalized  value, as support of denormalized values is optional in
- * relaxed mode.
+ * In rs_fp_relaxed mode, a denormalized input value may not yield the next denormalized
+ * value, as support of denormalized values is optional in relaxed mode.
  */
 extern float __attribute__((const, overloadable))
     nextafter(float v, float target);
@@ -3589,7 +3589,8 @@ extern float4 __attribute__((const, overloadable))
  *
  * Returns base raised to the power exponent, i.e. base ^ exponent.
  *
- * pown() and powr() are similar.  pown() takes an integer exponent. powr() assumes the base to be non-negative.
+ * pown() and powr() are similar.  pown() takes an integer exponent. powr() assumes the
+ * base to be non-negative.
  */
 extern float __attribute__((const, overloadable))
     pow(float base, float exponent);
@@ -3608,7 +3609,8 @@ extern float4 __attribute__((const, overloadable))
  *
  * Returns base raised to the power exponent, i.e. base ^ exponent.
  *
- * pow() and powr() are similar.  The both take a float exponent. powr() also assumes the base to be non-negative.
+ * pow() and powr() are similar.  The both take a float exponent. powr() also assumes the
+ * base to be non-negative.
  */
 extern float __attribute__((const, overloadable))
     pown(float base, int exponent);
@@ -3627,7 +3629,8 @@ extern float4 __attribute__((const, overloadable))
  *
  * Returns base raised to the power exponent, i.e. base ^ exponent.  base must be >= 0.
  *
- * pow() and pown() are similar.  They both make no assumptions about the base.  pow() takes a float exponent while pown() take an integer.
+ * pow() and pown() are similar.  They both make no assumptions about the base.
+ * pow() takes a float exponent while pown() take an integer.
  *
  * See also native_powr().
  */
@@ -3663,7 +3666,8 @@ extern float4 __attribute__((const, overloadable))
 /*
  * remainder: Remainder of a division
  *
- * Returns the remainder of (numerator / denominator), where the quotient is rounded towards the nearest integer.
+ * Returns the remainder of (numerator / denominator), where the quotient is rounded towards
+ * the nearest integer.
  *
  * The function fmod() is similar but rounds toward the closest interger.
  * For example, fmod(-3.8f, 2.f) returns -1.8f (-3.8f - -1.f * 2.f)
@@ -3688,16 +3692,20 @@ extern float4 __attribute__((const, overloadable))
  *
  * Only the sign and lowest three bits of the quotient are guaranteed to be accurate.
  *
- * This function is useful for implementing periodic functions.  The low three bits of the quotient gives the quadrant and the remainder the distance within the quadrant.  For example, an implementation of sin(x) could call remquo(x, PI / 2.f, &quadrant) to reduce very large value of x to something within a limited range.
+ * This function is useful for implementing periodic functions.  The low three bits of the
+ * quotient gives the quadrant and the remainder the distance within the quadrant.
+ * For example, an implementation of sin(x) could call remquo(x, PI / 2.f, &quadrant)
+ * to reduce very large value of x to something within a limited range.
  *
- * Example: remquo(-23.5f, 8.f, &quot) sets the lowest three bits of quot to 3 and the sign negative.  It returns 0.5f.
+ * Example: remquo(-23.5f, 8.f, &quot) sets the lowest three bits of quot to 3
+ * and the sign negative.  It returns 0.5f.
  *
  * Parameters:
- *   numerator: The numerator.
- *   denominator: The denominator.
+ *   numerator: Numerator.
+ *   denominator: Denominator.
  *   quotient: *quotient will be set to the integer quotient.
  *
- * Returns: The remainder, precise only for the low three bits.
+ * Returns: Remainder, precise only for the low three bits.
  */
 extern float __attribute__((overloadable))
     remquo(float numerator, float denominator, int* quotient);
@@ -3716,7 +3724,9 @@ extern float4 __attribute__((overloadable))
  *
  * Rounds to the nearest integral value.
  *
- * rint() rounds half values to even.  For example, rint(0.5f) returns 0.f and rint(1.5f) returns 2.f.  Similarly, rint(-0.5f) returns -0.f and rint(-1.5f) returns -2.f.
+ * rint() rounds half values to even.  For example, rint(0.5f) returns 0.f and
+ * rint(1.5f) returns 2.f.  Similarly, rint(-0.5f) returns -0.f and
+ * rint(-1.5f) returns -2.f.
  *
  * round() is similar but rounds away from zero.  trunc() truncates the decimal fraction.
  */
@@ -3756,7 +3766,9 @@ extern float4 __attribute__((const, overloadable))
  *
  * Round to the nearest integral value.
  *
- * round() rounds half values away from zero.  For example, round(0.5f) returns 1.f and round(1.5f) returns 2.f.  Similarly, round(-0.5f) returns -1.f and round(-1.5f) returns -2.f.
+ * round() rounds half values away from zero.  For example, round(0.5f) returns 1.f
+ * and round(1.5f) returns 2.f.  Similarly, round(-0.5f) returns -1.f
+ * and round(-1.5f) returns -2.f.
  *
  * rint() is similar but rounds half values toward even.  trunc() truncates the decimal fraction.
  */
@@ -3839,10 +3851,10 @@ extern float4 __attribute__((const, overloadable))
  * See also native_sincos().
  *
  * Parameters:
- *   v: The incoming value in radians
+ *   v: Incoming value in radians.
  *   cos: *cos will be set to the cosine value.
  *
- * Returns: sine of v
+ * Returns: Sine of v.
  */
 extern float __attribute__((overloadable))
     sincos(float v, float* cos);
@@ -3920,7 +3932,9 @@ extern float4 __attribute__((const, overloadable))
  *
  * Returns 0.f if v < edge, 1.f otherwise.
  *
- * This can be useful to create conditional computations without using loops and branching instructions.  For example, instead of computing (a[i] < b[i]) ? 0.f : atan2(a[i], b[i]) for the corresponding elements of a vector, you could instead use step(a, b) * atan2(a, b).
+ * This can be useful to create conditional computations without using loops and branching
+ * instructions.  For example, instead of computing (a[i] < b[i]) ? 0.f : atan2(a[i], b[i])
+ * for the corresponding elements of a vector, you could instead use step(a, b) * atan2(a, b).
  */
 extern float __attribute__((const, overloadable))
     step(float edge, float v);
@@ -4065,9 +4079,9 @@ extern float4 __attribute__((const, overloadable))
  * Clamp a value between low and high.
  *
  * Parameters:
- *   amount: The value to clamp
- *   low: Lower bound
- *   high: Upper bound
+ *   amount: Value to clamp.
+ *   low: Lower bound.
+ *   high: Upper bound.
  */
 extern char __attribute__((const, always_inline, overloadable))
     rsClamp(char amount, char low, char high);
diff --git a/renderscript/include/rs_matrix.rsh b/renderscript/include/rs_matrix.rsh
index bb7c6e8..aafc864 100644
--- a/renderscript/include/rs_matrix.rsh
+++ b/renderscript/include/rs_matrix.rsh
@@ -20,29 +20,29 @@
  * rs_matrix.rsh: Matrix Functions
  *
  * These functions let you manipulate square matrices of rank 2x2, 3x3, and 4x4.
- * They are particularly useful for graphical transformations and are
- * compatible with OpenGL.
+ * They are particularly useful for graphical transformations and are compatible
+ * with OpenGL.
  *
- * We use a zero-based index for rows and columns.  E.g. the last element of
- * a rs_matrix4x4 is found at (3, 3).
+ * We use a zero-based index for rows and columns.  E.g. the last element of a
+ * rs_matrix4x4 is found at (3, 3).
  *
- * RenderScript uses column-major matrices and column-based vectors.
- * Transforming a vector is done by postmultiplying the vector,
- * e.g. (matrix * vector), as provided by rsMatrixMultiply().
+ * RenderScript uses column-major matrices and column-based vectors.  Transforming
+ * a vector is done by postmultiplying the vector, e.g. (matrix * vector),
+ * as provided by rsMatrixMultiply().
  *
- * To create a transformation matrix that performs two transformations at
- * once, multiply the two source matrices, with the first transformation as the
- * right argument.  E.g. to create a transformation matrix that applies the
+ * To create a transformation matrix that performs two transformations at once,
+ * multiply the two source matrices, with the first transformation as the right
+ * argument.  E.g. to create a transformation matrix that applies the
  * transformation s1 followed by s2, call rsMatrixLoadMultiply(&combined, &s2, &s1).
  * This derives from s2 * (s1 * v), which is (s2 * s1) * v.
  *
  * We have two style of functions to create transformation matrices:
- * rsMatrixLoadTransformation and rsMatrixTransformation.  The
- * former style simply stores the transformation matrix in the first argument.
- * The latter modifies a pre-existing transformation matrix so that the new
- * transformation happens first.  E.g. if you call rsMatrixTranslate()
- * on a matrix that already does a scaling, the resulting matrix when applied
- * to a vector will first do the translation then the scaling.
+ * rsMatrixLoadTransformation and rsMatrixTransformation.  The former
+ * style simply stores the transformation matrix in the first argument.  The latter
+ * modifies a pre-existing transformation matrix so that the new transformation
+ * happens first.  E.g. if you call rsMatrixTranslate() on a matrix that already
+ * does a scaling, the resulting matrix when applied to a vector will first do the
+ * translation then the scaling.
  */
 
 #ifndef RENDERSCRIPT_RS_MATRIX_RSH
@@ -56,13 +56,13 @@
  * Computes 6 frustum planes from the view projection matrix
  *
  * Parameters:
- *   viewProj: matrix to extract planes from
- *   left: left plane
- *   right: right plane
- *   top: top plane
- *   bottom: bottom plane
- *   near: near plane
- *   far: far plane
+ *   viewProj: Matrix to extract planes from.
+ *   left: Left plane.
+ *   right: Right plane.
+ *   top: Top plane.
+ *   bottom: Bottom plane.
+ *   near: Near plane.
+ *   far: Far plane.
  */
 static inline void __attribute__((always_inline, overloadable))
     rsExtractFrustumPlanes(const rs_matrix4x4* viewProj, float4* left, float4* right, float4* top,
@@ -118,13 +118,13 @@ static inline void __attribute__((always_inline, overloadable))
  * Returns true if the sphere is within the 6 frustum planes.
  *
  * Parameters:
- *   sphere: float4 representing the sphere
- *   left: left plane
- *   right: right plane
- *   top: top plane
- *   bottom: bottom plane
- *   near: near plane
- *   far: far plane
+ *   sphere: float4 representing the sphere.
+ *   left: Left plane.
+ *   right: Right plane.
+ *   top: Top plane.
+ *   bottom: Bottom plane.
+ *   near: Near plane.
+ *   far: Far plane.
  */
 static inline bool __attribute__((always_inline, overloadable))
     rsIsSphereInFrustum(float4* sphere, float4* left, float4* right, float4* top, float4* bottom,
@@ -164,9 +164,9 @@ static inline bool __attribute__((always_inline, overloadable))
  * Warning: The order of the column and row parameters may be unexpected.
  *
  * Parameters:
- *   m: The matrix to extract the element from.
- *   col: The zero-based column of the element to be extracted.
- *   row: The zero-based row of the element to extracted.
+ *   m: Matrix to extract the element from.
+ *   col: Zero-based column of the element to be extracted.
+ *   row: Zero-based row of the element to extracted.
  */
 extern float __attribute__((overloadable))
     rsMatrixGet(const rs_matrix4x4* m, uint32_t col, uint32_t row);
@@ -183,7 +183,7 @@ extern float __attribute__((overloadable))
  * Returns true if the matrix was successfully inverted.
  *
  * Parameters:
- *   m: The matrix to invert.
+ *   m: Matrix to invert.
  */
 extern bool __attribute__((overloadable))
     rsMatrixInverse(rs_matrix4x4* m);
@@ -191,11 +191,11 @@ extern bool __attribute__((overloadable))
 /*
  * rsMatrixInverseTranspose: Inverts and transpose a matrix in place
  *
- * The matrix is first inverted then transposed.
- * Returns true if the matrix was successfully inverted.
+ * The matrix is first inverted then transposed. Returns true if the matrix was
+ * successfully inverted.
  *
  * Parameters:
- *   m: The matrix to modify.
+ *   m: Matrix to modify.
  */
 extern bool __attribute__((overloadable))
     rsMatrixInverseTranspose(rs_matrix4x4* m);
@@ -209,8 +209,8 @@ extern bool __attribute__((overloadable))
  * row 0, column 0 should be first, followed by the element at
  * row 0, column 1, etc.
  *
- * If loading from a matrix and the source is smaller than the destination, the rest of the
- * destination is filled with elements of the identity matrix.  E.g.
+ * If loading from a matrix and the source is smaller than the destination, the rest
+ * of the destination is filled with elements of the identity matrix.  E.g.
  * loading a rs_matrix2x2 into a rs_matrix4x4 will give:
  *
  * m00 m01 0.0 0.0
@@ -220,9 +220,9 @@ extern bool __attribute__((overloadable))
  *
  *
  * Parameters:
- *   destination: The matrix to set.
- *   array: The array of values to set the matrix to. These arrays should be 4, 9, or 16 floats long, depending on the matrix size.
- *   source: The source matrix.
+ *   destination: Matrix to set.
+ *   array: Array of values to set the matrix to. These arrays should be 4, 9, or 16 floats long, depending on the matrix size.
+ *   source: Source matrix.
  */
 extern void __attribute__((overloadable))
     rsMatrixLoad(rs_matrix4x4* destination, const float* array);
@@ -251,15 +251,14 @@ extern void __attribute__((overloadable))
 /*
  * rsMatrixLoadFrustum: Load a frustum projection matrix
  *
- * Constructs a frustum projection matrix, transforming the box
- * identified by the six clipping planes left, right, bottom, top,
- * near, far.
+ * Constructs a frustum projection matrix, transforming the box identified by
+ * the six clipping planes left, right, bottom, top, near, far.
  *
- * To apply this projection to a vector, multiply the vector by the
- * created matrix using rsMatrixMultiply().
+ * To apply this projection to a vector, multiply the vector by the created
+ * matrix using rsMatrixMultiply().
  *
  * Parameters:
- *   m: The matrix to set.
+ *   m: Matrix to set.
  */
 extern void __attribute__((overloadable))
     rsMatrixLoadFrustum(rs_matrix4x4* m, float left, float right, float bottom, float top,
@@ -271,7 +270,7 @@ extern void __attribute__((overloadable))
  * Set the elements of a matrix to the identity matrix.
  *
  * Parameters:
- *   m: The matrix to set.
+ *   m: Matrix to set.
  */
 extern void __attribute__((overloadable))
     rsMatrixLoadIdentity(rs_matrix4x4* m);
@@ -289,8 +288,7 @@ extern void __attribute__((overloadable))
  *
  * To combine two 4x4 transformaton matrices, multiply the second transformation matrix
  * by the first transformation matrix.  E.g. to create a transformation matrix that applies
- * the transformation s1 followed by s2, call
- * rsMatrixLoadMultiply(&combined, &s2, &s1).
+ * the transformation s1 followed by s2, call rsMatrixLoadMultiply(&combined, &s2, &s1).
  *
  * Warning: Prior to version 21, storing the result back into right matrix is not supported and
  * will result in undefined behavior.  Use rsMatrixMulitply instead.   E.g. instead of doing
@@ -298,9 +296,9 @@ extern void __attribute__((overloadable))
  * rsMatrixLoadMultiply (&m2l, &m2r, &m2l) works as expected.
  *
  * Parameters:
- *   m: The matrix to set.
- *   lhs: The left matrix of the product.
- *   rhs: The right matrix of the product.
+ *   m: Matrix to set.
+ *   lhs: Left matrix of the product.
+ *   rhs: Right matrix of the product.
  */
 extern void __attribute__((overloadable))
     rsMatrixLoadMultiply(rs_matrix4x4* m, const rs_matrix4x4* lhs, const rs_matrix4x4* rhs);
@@ -314,18 +312,17 @@ extern void __attribute__((overloadable))
 /*
  * rsMatrixLoadOrtho: Load an orthographic projection matrix
  *
- * Constructs an orthographic projection matrix, transforming the box
- * identified by the six clipping planes left, right, bottom, top,
- * near, far into a unit cube with a corner at
- * (-1, -1, -1) and the opposite at (1, 1, 1).
+ * Constructs an orthographic projection matrix, transforming the box identified by the
+ * six clipping planes left, right, bottom, top, near, far into a unit cube
+ * with a corner at (-1, -1, -1) and the opposite at (1, 1, 1).
  *
- * To apply this projection to a vector, multiply the vector by the
- * created matrix using rsMatrixMultiply().
+ * To apply this projection to a vector, multiply the vector by the created matrix
+ * using rsMatrixMultiply().
  *
  * See https://en.wikipedia.org/wiki/Orthographic_projection .
  *
  * Parameters:
- *   m: The matrix to set.
+ *   m: Matrix to set.
  */
 extern void __attribute__((overloadable))
     rsMatrixLoadOrtho(rs_matrix4x4* m, float left, float right, float bottom, float top, float near,
@@ -336,15 +333,15 @@ extern void __attribute__((overloadable))
  *
  * Constructs a perspective projection matrix, assuming a symmetrical field of view.
  *
- * To apply this projection to a vector, multiply the vector by the
- * created matrix using rsMatrixMultiply().
+ * To apply this projection to a vector, multiply the vector by the created matrix
+ * using rsMatrixMultiply().
  *
  * Parameters:
- *   m: The matrix to set.
+ *   m: Matrix to set.
  *   fovy: Field of view, in degrees along the Y axis.
  *   aspect: Ratio of x / y.
- *   near: The near clipping plane.
- *   far: The far clipping plane.
+ *   near: Near clipping plane.
+ *   far: Far clipping plane.
  */
 extern void __attribute__((overloadable))
     rsMatrixLoadPerspective(rs_matrix4x4* m, float fovy, float aspect, float near, float far);
@@ -352,20 +349,18 @@ extern void __attribute__((overloadable))
 /*
  * rsMatrixLoadRotate: Load a rotation matrix
  *
- * This function creates a rotation matrix.  The axis of rotation is the
- * (x, y, z) vector.
+ * This function creates a rotation matrix.  The axis of rotation is the (x, y, z) vector.
  *
- * To rotate a vector, multiply the vector by the created matrix
- * using rsMatrixMultiply().
+ * To rotate a vector, multiply the vector by the created matrix using rsMatrixMultiply().
  *
  * See http://en.wikipedia.org/wiki/Rotation_matrix .
  *
  * Parameters:
- *   m: The matrix to set.
+ *   m: Matrix to set.
  *   rot: How much rotation to do, in degrees.
- *   x: The x component of the vector that is the axis of rotation.
- *   y: The y component of the vector that is the axis of rotation.
- *   z: The z component of the vector that is the axis of rotation.
+ *   x: X component of the vector that is the axis of rotation.
+ *   y: Y component of the vector that is the axis of rotation.
+ *   z: Z component of the vector that is the axis of rotation.
  */
 extern void __attribute__((overloadable))
     rsMatrixLoadRotate(rs_matrix4x4* m, float rot, float x, float y, float z);
@@ -373,17 +368,16 @@ extern void __attribute__((overloadable))
 /*
  * rsMatrixLoadScale: Load a scaling matrix
  *
- * This function creates a scaling matrix, where each component of a
- * vector is multiplied by a number.  This number can be negative.
+ * This function creates a scaling matrix, where each component of a vector is multiplied
+ * by a number.  This number can be negative.
  *
- * To scale a vector, multiply the vector by the created matrix
- * using rsMatrixMultiply().
+ * To scale a vector, multiply the vector by the created matrix using rsMatrixMultiply().
  *
  * Parameters:
- *   m: The matrix to set.
- *   x: The multiple to scale the x components by.
- *   y: The multiple to scale the y components by.
- *   z: The multiple to scale the z components by.
+ *   m: Matrix to set.
+ *   x: Multiple to scale the x components by.
+ *   y: Multiple to scale the y components by.
+ *   z: Multiple to scale the z components by.
  */
 extern void __attribute__((overloadable))
     rsMatrixLoadScale(rs_matrix4x4* m, float x, float y, float z);
@@ -391,17 +385,17 @@ extern void __attribute__((overloadable))
 /*
  * rsMatrixLoadTranslate: Load a translation matrix
  *
- * This function creates a translation matrix, where a
- * number is added to each element of a vector.
+ * This function creates a translation matrix, where a number is added to each element of
+ * a vector.
  *
- * To translate a vector, multiply the vector by the created matrix
- * using rsMatrixMultiply().
+ * To translate a vector, multiply the vector by the created matrix using
+ * rsMatrixMultiply().
  *
  * Parameters:
- *   m: The matrix to set.
- *   x: The number to add to each x component.
- *   y: The number to add to each y component.
- *   z: The number to add to each z component.
+ *   m: Matrix to set.
+ *   x: Number to add to each x component.
+ *   y: Number to add to each y component.
+ *   z: Number to add to each z component.
  */
 extern void __attribute__((overloadable))
     rsMatrixLoadTranslate(rs_matrix4x4* m, float x, float y, float z);
@@ -427,8 +421,8 @@ extern void __attribute__((overloadable))
  * Starting with API 14, this function takes a const matrix as the first argument.
  *
  * Parameters:
- *   m: The left matrix of the product and the matrix to be set.
- *   rhs: The right matrix of the product.
+ *   m: Left matrix of the product and the matrix to be set.
+ *   rhs: Right matrix of the product.
  */
 extern void __attribute__((overloadable))
     rsMatrixMultiply(rs_matrix4x4* m, const rs_matrix4x4* rhs);
@@ -504,18 +498,18 @@ extern float2 __attribute__((overloadable))
  *
  * Multiply the matrix m with a rotation matrix.
  *
- * This function modifies a transformation matrix to first do a rotation.
- * The axis of rotation is the (x, y, z) vector.
+ * This function modifies a transformation matrix to first do a rotation.  The axis of
+ * rotation is the (x, y, z) vector.
  *
- * To apply this combined transformation to a vector, multiply
- * the vector by the created matrix using rsMatrixMultiply().
+ * To apply this combined transformation to a vector, multiply the vector by the created
+ * matrix using rsMatrixMultiply().
  *
  * Parameters:
- *   m: The matrix to modify.
+ *   m: Matrix to modify.
  *   rot: How much rotation to do, in degrees.
- *   x: The x component of the vector that is the axis of rotation.
- *   y: The y component of the vector that is the axis of rotation.
- *   z: The z component of the vector that is the axis of rotation.
+ *   x: X component of the vector that is the axis of rotation.
+ *   y: Y component of the vector that is the axis of rotation.
+ *   z: Z component of the vector that is the axis of rotation.
  */
 extern void __attribute__((overloadable))
     rsMatrixRotate(rs_matrix4x4* m, float rot, float x, float y, float z);
@@ -525,18 +519,17 @@ extern void __attribute__((overloadable))
  *
  * Multiply the matrix m with a scaling matrix.
  *
- * This function modifies a transformation matrix to first do a scaling.
- * When scaling, each component of a vector is multiplied by a number.
- * This number can be negative.
+ * This function modifies a transformation matrix to first do a scaling.   When scaling,
+ * each component of a vector is multiplied by a number.  This number can be negative.
  *
- * To apply this combined transformation to a vector, multiply
- * the vector by the created matrix using rsMatrixMultiply().
+ * To apply this combined transformation to a vector, multiply the vector by the created
+ * matrix using rsMatrixMultiply().
  *
  * Parameters:
- *   m: The matrix to modify.
- *   x: The multiple to scale the x components by.
- *   y: The multiple to scale the y components by.
- *   z: The multiple to scale the z components by.
+ *   m: Matrix to modify.
+ *   x: Multiple to scale the x components by.
+ *   y: Multiple to scale the y components by.
+ *   z: Multiple to scale the z components by.
  */
 extern void __attribute__((overloadable))
     rsMatrixScale(rs_matrix4x4* m, float x, float y, float z);
@@ -549,10 +542,10 @@ extern void __attribute__((overloadable))
  * Warning: The order of the column and row parameters may be unexpected.
  *
  * Parameters:
- *   m: The matrix that will be modified.
- *   col: The zero-based column of the element to be set.
- *   row: The zero-based row of the element to be set.
- *   v: The value to set.
+ *   m: Matrix that will be modified.
+ *   col: Zero-based column of the element to be set.
+ *   row: Zero-based row of the element to be set.
+ *   v: Value to set.
  */
 extern void __attribute__((overloadable))
     rsMatrixSet(rs_matrix4x4* m, uint32_t col, uint32_t row, float v);
@@ -568,18 +561,17 @@ extern void __attribute__((overloadable))
  *
  * Multiply the matrix m with a translation matrix.
  *
- * This function modifies a transformation matrix to first
- * do a translation.  When translating, a number is added
- * to each component of a vector.
+ * This function modifies a transformation matrix to first do a translation.  When
+ * translating, a number is added to each component of a vector.
  *
- * To apply this combined transformation to a vector, multiply
- * the vector by the created matrix using rsMatrixMultiply().
+ * To apply this combined transformation to a vector, multiply the vector by the
+ * created matrix using rsMatrixMultiply().
  *
  * Parameters:
- *   m: The matrix to modify.
- *   x: The number to add to each x component.
- *   y: The number to add to each y component.
- *   z: The number to add to each z component.
+ *   m: Matrix to modify.
+ *   x: Number to add to each x component.
+ *   y: Number to add to each y component.
+ *   z: Number to add to each z component.
  */
 extern void __attribute__((overloadable))
     rsMatrixTranslate(rs_matrix4x4* m, float x, float y, float z);
@@ -590,7 +582,7 @@ extern void __attribute__((overloadable))
  * Transpose the matrix m in place.
  *
  * Parameters:
- *   m: The matrix to transpose.
+ *   m: Matrix to transpose.
  */
 extern void __attribute__((overloadable))
     rsMatrixTranspose(rs_matrix4x4* m);
diff --git a/renderscript/include/rs_object_info.rsh b/renderscript/include/rs_object_info.rsh
index 6ea7ed6..4305f2a 100644
--- a/renderscript/include/rs_object_info.rsh
+++ b/renderscript/include/rs_object_info.rsh
@@ -19,26 +19,52 @@
 /*
  * rs_object_info.rsh: Object Characteristics Functions
  *
- * The functions below can be used to query the characteristics of an allocation,
- * element, or sampler object.  These objects are created from Java.
+ * The functions below can be used to query the characteristics of an Allocation, Element,
+ * or Sampler object.  These objects are created from Java.  You can't create them from a
+ * script.
  *
- * The term "element" is used a bit ambiguously in RenderScript, as both
- * the type of an item of an allocation and the instantiation of that type:
+ * Allocations:
+ *
+ * Allocations are the primary method used to pass data to and from RenderScript kernels.
+ *
+ * They are a structured collection of cells that can be used to store bitmaps, textures,
+ * arbitrary data points, etc.
+ *
+ * This collection of cells may have many dimensions (X, Y, Z, Array0, Array1, Array2, Array3),
+ * faces (for cubemaps), and level of details (for mipmapping).
+ *
+ * See the android.renderscript.Allocation for details on to create Allocations.
+ *
+ * Elements:
+ *
+ * The term "element" is used a bit ambiguously in RenderScript, as both type information
+ * for the cells of an Allocation and the instantiation of that type.  For example:
  * - rs_element is a handle to a type specification, and
- * - In functions like rsGetElementAt(), "element" means the instantiation
- *     of the type, i.e. an item of an allocation.
+ * - In functions like rsGetElementAt(), "element" means the instantiation of the type,
+ *     i.e. a cell of an Allocation.
  *
  * The functions below let you query the characteristics of the type specificiation.
  *
- * To create complex elements, use the Element.Builder Java class.
- * For common elements, in Java you can simply use one of the many predefined elements
- * like F32_2.  You can't create elements from a script.
+ * An Element can specify a simple data types as found in C, e.g. an integer, float, or
+ * boolean.  It can also specify a handle to a RenderScript object.  See rs_data_type for
+ * a list of basic types.
  *
- * An element can be a simple data type as found in C/C++, a handle type,
- * a structure, or a fixed size vector (of size 2, 3, or 4) of sub-elements.
+ * Elements can specify fixed size vector (of size 2, 3, or 4) versions of the basic types.
+ * Elements can be grouped together into complex Elements, creating the equivalent of
+ * C structure definitions.
  *
- * Elements can also have a kind, which is semantic information used mostly to
- * interpret pixel data.
+ * Elements can also have a kind, which is semantic information used to interpret pixel
+ * data.  See rs_data_kind.
+ *
+ * When creating Allocations of common elements, you can simply use one of the many predefined
+ * Elements like F32_2.
+ *
+ * To create complex Elements, use the Element.Builder Java class.
+ *
+ * Samplers:
+ *
+ * Samplers objects define how Allocations can be read as structure within a kernel.
+ * See android.renderscript.S.
  */
 
 #ifndef RENDERSCRIPT_RS_OBJECT_INFO_RSH
@@ -47,8 +73,8 @@
 /*
  * rsAllocationGetDimFaces: Presence of more than one face
  *
- * If the allocation is a cubemap, this function returns 1 if there's more than
- * one face present.  In all other cases, it returns 0.
+ * If the Allocation is a cubemap, this function returns 1 if there's more than one face
+ * present.  In all other cases, it returns 0.
  *
  * Use rsGetDimHasFaces() to get the dimension of a currently running kernel.
  *
@@ -60,8 +86,8 @@ extern uint32_t __attribute__((overloadable))
 /*
  * rsAllocationGetDimLOD: Presence of levels of detail
  *
- * Query an allocation for the presence of more than one Level Of Detail.
- * This is useful for mipmaps.
+ * Query an Allocation for the presence of more than one Level Of Detail.  This is useful
+ * for mipmaps.
  *
  * Use rsGetDimLod() to get the dimension of a currently running kernel.
  *
@@ -73,11 +99,11 @@ extern uint32_t __attribute__((overloadable))
 /*
  * rsAllocationGetDimX: Size of the X dimension
  *
- * Returns the size of the X dimension of the allocation.
+ * Returns the size of the X dimension of the Allocation.
  *
  * Use rsGetDimX() to get the dimension of a currently running kernel.
  *
- * Returns: The X dimension of the allocation.
+ * Returns: X dimension of the Allocation.
  */
 extern uint32_t __attribute__((overloadable))
     rsAllocationGetDimX(rs_allocation a);
@@ -85,12 +111,12 @@ extern uint32_t __attribute__((overloadable))
 /*
  * rsAllocationGetDimY: Size of the Y dimension
  *
- * Returns the size of the Y dimension of the allocation.
- * If the allocation has less than two dimensions, returns 0.
+ * Returns the size of the Y dimension of the Allocation.  If the Allocation has less
+ * than two dimensions, returns 0.
  *
  * Use rsGetDimY() to get the dimension of a currently running kernel.
  *
- * Returns: The Y dimension of the allocation.
+ * Returns: Y dimension of the Allocation.
  */
 extern uint32_t __attribute__((overloadable))
     rsAllocationGetDimY(rs_allocation a);
@@ -98,23 +124,26 @@ extern uint32_t __attribute__((overloadable))
 /*
  * rsAllocationGetDimZ: Size of the Z dimension
  *
- * Returns the size of the Z dimension of the allocation.
- * If the allocation has less than three dimensions, returns 0.
+ * Returns the size of the Z dimension of the Allocation.  If the Allocation has less
+ * than three dimensions, returns 0.
  *
  * Use rsGetDimZ() to get the dimension of a currently running kernel.
  *
- * Returns: The Z dimension of the allocation.
+ * Returns: Z dimension of the Allocation.
  */
 extern uint32_t __attribute__((overloadable))
     rsAllocationGetDimZ(rs_allocation a);
 
 /*
- * Get the element object describing the allocation's layout
+ * rsAllocationGetElement: Get the object that describes the cell of an Allocation
+ *
+ * Get the Element object describing the type, kind, and other characteristics of a cell
+ * of an Allocation.  See the rsElement* functions below.
  *
  * Parameters:
- *   a: allocation to get data from
+ *   a: Allocation to get data from.
  *
- * Returns: element describing allocation layout
+ * Returns: Element describing Allocation layout.
  */
 extern rs_element __attribute__((overloadable))
     rsAllocationGetElement(rs_allocation a);
@@ -122,12 +151,10 @@ extern rs_element __attribute__((overloadable))
 /*
  * rsClearObject: Release an object
  *
- * Tells the run time that this handle will no longer be used to access the
- * the related object.  If this was the last handle to that object, resource
- * recovery may happen.
+ * Tells the run time that this handle will no longer be used to access the the related
+ * object.  If this was the last handle to that object, resource recovery may happen.
  *
- * After calling this function, *dst will be set to an empty handle.  See
- * rsIsObject().
+ * After calling this function, *dst will be set to an empty handle.  See rsIsObject().
  */
 extern void __attribute__((overloadable))
     rsClearObject(rs_element* dst);
@@ -152,9 +179,8 @@ extern void __attribute__((overloadable))
  * This function does not validate that the internal pointer used in the handle
  * points to an actual valid object; it only checks for null.
  *
- * This function can be used to check the element returned by
- * rsElementGetSubElement() or see if rsClearObject() has been called on a
- * handle.
+ * This function can be used to check the Element returned by rsElementGetSubElement()
+ * or see if rsClearObject() has been called on a handle.
  */
 extern bool __attribute__((overloadable))
     rsIsObject(rs_element v);
@@ -172,9 +198,9 @@ extern bool __attribute__((overloadable))
     rsIsObject(rs_script v);
 
 /*
- * rsElementGetBytesSize: Return the size of an element
+ * rsElementGetBytesSize: Size of an Element
  *
- * Returns the size in bytes that an instantiation of this element will occupy.
+ * Returns the size in bytes that an instantiation of this Element will occupy.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern uint32_t __attribute__((overloadable))
@@ -182,9 +208,9 @@ extern uint32_t __attribute__((overloadable))
 #endif
 
 /*
- * rsElementGetDataKind: Return the kind of an element
+ * rsElementGetDataKind: Kind of an Element
  *
- * Returns the element's data kind.  This is used to interpret pixel data.
+ * Returns the Element's data kind.  This is used to interpret pixel data.
  *
  * See rs_data_kind.
  */
@@ -194,17 +220,18 @@ extern rs_data_kind __attribute__((overloadable))
 #endif
 
 /*
- * rsElementGetDataType: Return the data type of an element
- *
- * Returns the element's base data type.  This can be a type similar to C/C++ (e.g. RS_TYPE_UNSIGNED_8),
- * a handle (e.g. RS_TYPE_ALLOCATION and RS_TYPE_ELEMENT), or a more complex numerical type
- * (e.g.RS_TYPE_UNSIGNED_5_6_5 and RS_TYPE_MATRIX_4X4).
+ * rsElementGetDataType: Data type of an Element
  *
- * If the element describes a vector, this function returns the data type of one of its items.
+ * Returns the Element's base data type.  This can be a type similar to C/C++ (e.g.
+ * RS_TYPE_UNSIGNED_8), a handle (e.g. RS_TYPE_ALLOCATION and RS_TYPE_ELEMENT), or a
+ * more complex numerical type (e.g. RS_TYPE_UNSIGNED_5_6_5 and RS_TYPE_MATRIX_4X4).
+ * See rs_data_type.
  *
- * If the element describes a structure, RS_TYPE_NONE is returned.
+ * If the Element describes a vector, this function returns the data type of one of its items.
+ * Use rsElementGetVectorSize to get the size of the vector.
  *
- * See rs_data_type.
+ * If the Element describes a structure, RS_TYPE_NONE is returned.  Use the rsElementGetSub*
+ * functions to explore this complex Element.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern rs_data_type __attribute__((overloadable))
@@ -212,19 +239,19 @@ extern rs_data_type __attribute__((overloadable))
 #endif
 
 /*
- * rsElementGetSubElement: Return a sub element of a complex element
+ * rsElementGetSubElement: Sub-element of a complex Element
  *
- * For the element represents a structure, this function returns the sub-element at
- * the specified index.
+ * For Elements that represents a structure, this function returns the sub-element at the
+ * specified index.
  *
- * If the element is not a structure or the index is greater or equal to the number
- * of sub-elements, an invalid handle is returned.
+ * If the Element is not a structure or the index is greater or equal to the number of
+ * sub-elements, an invalid handle is returned.
  *
  * Parameters:
- *   e: Element to query
- *   index: Index of the sub-element to return
+ *   e: Element to query.
+ *   index: Index of the sub-element to return.
  *
- * Returns: Sub-element at the given index
+ * Returns: Sub-element at the given index.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern rs_element __attribute__((overloadable))
@@ -232,17 +259,17 @@ extern rs_element __attribute__((overloadable))
 #endif
 
 /*
- * rsElementGetSubElementArraySize: Return the array size of a sub element of a complex element
+ * rsElementGetSubElementArraySize: Array size of a sub-element of a complex Element
  *
- * For complex elements, some sub-elements could be statically
- * sized arrays. This function returns the array size of the
- * sub-element at the index.
+ * For complex Elements, sub-elements can be statically sized arrays.  This function
+ * returns the array size of the sub-element at the index.  This sub-element repetition
+ * is different than fixed size vectors.
  *
  * Parameters:
- *   e: Element to query
- *   index: Index of the sub-element
+ *   e: Element to query.
+ *   index: Index of the sub-element.
  *
- * Returns: Array size of the sub-element at the given index
+ * Returns: Array size of the sub-element.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern uint32_t __attribute__((overloadable))
@@ -250,17 +277,16 @@ extern uint32_t __attribute__((overloadable))
 #endif
 
 /*
- * rsElementGetSubElementCount: Return the number of sub-elements
+ * rsElementGetSubElementCount: Number of sub-elements
  *
- * Elements could be simple, such as an int or a float, or a
- * structure with multiple sub-elements, such as a collection of
- * floats, float2, float4.  This function returns zero for simple
- * elements or the number of sub-elements otherwise.
+ * Elements can be simple, such as an int or a float, or a structure with multiple
+ * sub-elements.  This function returns zero for simple Elements and the number of
+ * sub-elements for complex Elements.
  *
  * Parameters:
- *   e: Element to get data from
+ *   e: Element to get data from.
  *
- * Returns: Number of sub-elements in this element
+ * Returns: Number of sub-elements.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern uint32_t __attribute__((overloadable))
@@ -268,18 +294,18 @@ extern uint32_t __attribute__((overloadable))
 #endif
 
 /*
- * rsElementGetSubElementName: Return the name of a sub-element
+ * rsElementGetSubElementName: Name of a sub-element
  *
- * For complex elements, this function returns the name of the sub-element
- * at the specified index.
+ * For complex Elements, this function returns the name of the sub-element at the
+ * specified index.
  *
  * Parameters:
- *   e: Element to get data from
- *   index: Index of the sub-element
- *   name: Array to store the name into
- *   nameLength: Length of the provided name array
+ *   e: Element to get data from.
+ *   index: Index of the sub-element.
+ *   name: Address of the array to store the name into.
+ *   nameLength: Length of the provided name array.
  *
- * Returns: Number of characters actually written, excluding the null terminator
+ * Returns: Number of characters copied, excluding the null terminator.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern uint32_t __attribute__((overloadable))
@@ -287,16 +313,16 @@ extern uint32_t __attribute__((overloadable))
 #endif
 
 /*
- * rsElementGetSubElementNameLength: Return the length of the name of a sub-element
+ * rsElementGetSubElementNameLength: Length of the name of a sub-element
  *
- * For complex elements, this function will return the length of
- * sub-element name at index
+ * For complex Elements, this function returns the length of the name of the sub-element
+ * at the specified index.
  *
  * Parameters:
- *   e: Element to get data from
- *   index: Index of the sub-element to return
+ *   e: Element to get data from.
+ *   index: Index of the sub-element.
  *
- * Returns: Length of the sub-element name including the null terminator (size of buffer needed to write the name)
+ * Returns: Length of the sub-element name including the null terminator.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern uint32_t __attribute__((overloadable))
@@ -304,14 +330,19 @@ extern uint32_t __attribute__((overloadable))
 #endif
 
 /*
- * This function specifies the location of a sub-element within
- * the element
+ * rsElementGetSubElementOffsetBytes: Offset of the instantiated sub-element
+ *
+ * This function returns the relative position of the instantiation of the specified
+ * sub-element within the instantiation of the Element.
+ *
+ * For example, if the Element describes a 32 bit float followed by a 32 bit integer,
+ * the offset return for the first will be 0 and the second 4.
  *
  * Parameters:
- *   e: Element to get data from
- *   index: Index of the sub-element
+ *   e: Element to get data from.
+ *   index: Index of the sub-element.
  *
- * Returns: Offset in bytes of sub-element in this element at given index
+ * Returns: Offset in bytes.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern uint32_t __attribute__((overloadable))
@@ -319,12 +350,15 @@ extern uint32_t __attribute__((overloadable))
 #endif
 
 /*
- * Returns the element's vector size
+ * rsElementGetVectorSize: Vector size of the Element
+ *
+ * Returns the Element's vector size.  If the Element does not represent a vector,
+ * 1 is returned.
  *
  * Parameters:
- *   e: Element to get data from
+ *   e: Element to get data from.
  *
- * Returns: Length of the element vector (for float2, float3, etc.)
+ * Returns: Length of the element vector.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern uint32_t __attribute__((overloadable))
@@ -332,24 +366,22 @@ extern uint32_t __attribute__((overloadable))
 #endif
 
 /*
- * rsGetAllocation: Returns the Allocation for a given pointer
+ * rsGetAllocation: Return the Allocation for a given pointer
  *
  * DEPRECATED.  Do not use.
  *
- * Returns the Allocation for a given pointer.  The pointer should point within
- * a valid allocation.  The results are undefined if the pointer is not from a
- * valid allocation.
+ * Returns the Allocation for a given pointer.  The pointer should point within a valid
+ * allocation.  The results are undefined if the pointer is not from a valid Allocation.
  */
 extern rs_allocation __attribute__((overloadable))
     rsGetAllocation(const void* p);
 
 /*
- *  Get sampler anisotropy
+ * rsSamplerGetAnisotropy: Anisotropy of the Sampler
  *
- * Parameters:
- *   s: sampler to query
+ * Get the Sampler's anisotropy.
  *
- * Returns: anisotropy
+ * See android.renderscript.S.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern float __attribute__((overloadable))
@@ -357,12 +389,11 @@ extern float __attribute__((overloadable))
 #endif
 
 /*
- * Get sampler magnification value
+ * rsSamplerGetMagnification: Sampler magnification value
  *
- * Parameters:
- *   s: sampler to query
+ * Get the Sampler's magnification value.
  *
- * Returns: magnification value
+ * See android.renderscript.S.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern rs_sampler_value __attribute__((overloadable))
@@ -370,12 +401,11 @@ extern rs_sampler_value __attribute__((overloadable))
 #endif
 
 /*
- * Get sampler minification value
+ * rsSamplerGetMinification: Sampler minification value
  *
- * Parameters:
- *   s: sampler to query
+ * Get the Sampler's minification value.
  *
- * Returns: minification value
+ * See android.renderscript.S.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern rs_sampler_value __attribute__((overloadable))
@@ -383,12 +413,11 @@ extern rs_sampler_value __attribute__((overloadable))
 #endif
 
 /*
- * Get sampler wrap S value
+ * rsSamplerGetWrapS: Sampler wrap S value
  *
- * Parameters:
- *   s: sampler to query
+ * Get the Sampler's wrap S value.
  *
- * Returns: wrap S value
+ * See android.renderscript.S.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern rs_sampler_value __attribute__((overloadable))
@@ -396,12 +425,11 @@ extern rs_sampler_value __attribute__((overloadable))
 #endif
 
 /*
- * Get sampler wrap T value
+ * rsSamplerGetWrapT: Sampler wrap T value
  *
- * Parameters:
- *   s: sampler to query
+ * Get the sampler's wrap T value.
  *
- * Returns: wrap T value
+ * See android.renderscript.S.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 extern rs_sampler_value __attribute__((overloadable))
diff --git a/renderscript/include/rs_object_types.rsh b/renderscript/include/rs_object_types.rsh
index 0fc5e4d..9b11b32 100644
--- a/renderscript/include/rs_object_types.rsh
+++ b/renderscript/include/rs_object_types.rsh
@@ -19,9 +19,8 @@
 /*
  * rs_object_types.rsh: Object Types
  *
- * The types below are used to manipulate RenderScript objects like allocations,
- * samplers, elements, and scripts.  Most of these object are created using the Java
- * RenderScript APIs.
+ * The types below are used to manipulate RenderScript objects like allocations, samplers,
+ * elements, and scripts.  Most of these object are created using the Java RenderScript APIs.
  */
 
 #ifndef RENDERSCRIPT_RS_OBJECT_TYPES_RSH
@@ -48,46 +47,52 @@ struct {\
 /*
  * rs_element: Handle to an element
  *
- * Opaque handle to a RenderScript element.
- * See: android.renderscript.Element
+ * An opaque handle to a RenderScript element.
+ *
+ * See android.renderscript.Element.
  */
 typedef _RS_HANDLE rs_element;
 
 /*
  * rs_type: Handle to a Type
  *
- * Opaque handle to a RenderScript type.
- * See: android.renderscript.Type
+ * An opaque handle to a RenderScript type.
+ *
+ * See android.renderscript.Type.
  */
 typedef _RS_HANDLE rs_type;
 
 /*
  * rs_allocation: Handle to an allocation
  *
- * Opaque handle to a RenderScript allocation.
- * See: android.renderscript.Allocation
+ * An opaque handle to a RenderScript allocation.
+ *
+ * See android.renderscript.Allocation.
  */
 typedef _RS_HANDLE rs_allocation;
 
 /*
  * rs_sampler: Handle to a Sampler
  *
- * Opaque handle to a RenderScript sampler object.
- * See: android.renderscript.Sampler
+ * An opaque handle to a RenderScript sampler object.
+ *
+ * See android.renderscript.Sampler.
  */
 typedef _RS_HANDLE rs_sampler;
 
 /*
  * rs_script: Handle to a Script
  *
- * Opaque handle to a RenderScript script object.
- * See: android.renderscript.ScriptC
+ * An opaque handle to a RenderScript script object.
+ *
+ * See android.renderscript.ScriptC.
  */
 typedef _RS_HANDLE rs_script;
 
 /*
  * rs_allocation_cubemap_face: Enum for selecting cube map faces
  *
+ * An enum used to specify one the six faces of a cubemap.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 14))
 typedef enum {
@@ -101,7 +106,7 @@ typedef enum {
 #endif
 
 /*
- * rs_allocation_usage_type: Bitfield to specify the usage types for an allocation
+ * rs_allocation_usage_type: Bitfield to specify how an allocation is used
  *
  * These values are ORed together to specify which usages or memory spaces are
  * relevant to an allocation or an operation on an allocation.
@@ -120,54 +125,45 @@ typedef enum {
 #endif
 
 /*
- * rs_data_type: Element data types
- *
- * DataType represents the basic type information for a basic element.  The
- * naming convention follows.  For numeric types it is FLOAT,
- * SIGNED, or UNSIGNED followed by the _BITS where BITS is the
- * size of the data.  BOOLEAN is a true / false (1,0)
- * represented in an 8 bit container.  The UNSIGNED variants
- * with multiple bit definitions are for packed graphical data
- * formats and represent vectors with per vector member sizes
- * which are treated as a single unit for packing and alignment
- * purposes.
- *
- * MATRIX the three matrix types contain FLOAT_32 elements and are treated
- * as 32 bits for alignment purposes.
- *
- * RS_* objects.  32 bit opaque handles.
+ * rs_data_type: Element basic data type
+ *
+ * rs_data_type is used to encode the type information of a basic element.
+ *
+ * RS_TYPE_UNSIGNED_5_6_5, RS_TYPE_UNSIGNED_5_5_5_1, RS_TYPE_UNSIGNED_4_4_4_4 are for packed
+ * graphical data formats and represent vectors with per vector member sizes which are treated
+ * as a single unit for packing and alignment purposes.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 typedef enum {
-    RS_TYPE_NONE = 0,
-    RS_TYPE_FLOAT_32 = 2,
-    RS_TYPE_FLOAT_64 = 3,
-    RS_TYPE_SIGNED_8 = 4,
-    RS_TYPE_SIGNED_16 = 5,
-    RS_TYPE_SIGNED_32 = 6,
-    RS_TYPE_SIGNED_64 = 7,
-    RS_TYPE_UNSIGNED_8 = 8,
-    RS_TYPE_UNSIGNED_16 = 9,
-    RS_TYPE_UNSIGNED_32 = 10,
-    RS_TYPE_UNSIGNED_64 = 11,
-    RS_TYPE_BOOLEAN = 12,
-    RS_TYPE_UNSIGNED_5_6_5 = 13,
-    RS_TYPE_UNSIGNED_5_5_5_1 = 14,
-    RS_TYPE_UNSIGNED_4_4_4_4 = 15,
-    RS_TYPE_MATRIX_4X4 = 16,
-    RS_TYPE_MATRIX_3X3 = 17,
-    RS_TYPE_MATRIX_2X2 = 18,
-    RS_TYPE_ELEMENT = 1000,
-    RS_TYPE_TYPE = 1001,
-    RS_TYPE_ALLOCATION = 1002,
-    RS_TYPE_SAMPLER = 1003,
-    RS_TYPE_SCRIPT = 1004,
-    RS_TYPE_MESH = 1005,
-    RS_TYPE_PROGRAM_FRAGMENT = 1006,
-    RS_TYPE_PROGRAM_VERTEX = 1007,
-    RS_TYPE_PROGRAM_RASTER = 1008,
-    RS_TYPE_PROGRAM_STORE = 1009,
-    RS_TYPE_FONT = 1010,
+    RS_TYPE_NONE = 0, // Element is a complex type, i.e. a struct.
+    RS_TYPE_FLOAT_32 = 2, // A 32 bit float point value.
+    RS_TYPE_FLOAT_64 = 3, // A 64 bit floating point value.
+    RS_TYPE_SIGNED_8 = 4, // An 8 bit signed integer.
+    RS_TYPE_SIGNED_16 = 5, // A 16 bit signed integer.
+    RS_TYPE_SIGNED_32 = 6, // A 32 bit signed integer.
+    RS_TYPE_SIGNED_64 = 7, // A 64 bit signed integer.
+    RS_TYPE_UNSIGNED_8 = 8, // An 8 bit unsigned integer.
+    RS_TYPE_UNSIGNED_16 = 9, // A 16 bit unsigned integer.
+    RS_TYPE_UNSIGNED_32 = 10, // A 32 bit unsigned integer.
+    RS_TYPE_UNSIGNED_64 = 11, // A 64 bit unsigned integer.
+    RS_TYPE_BOOLEAN = 12, // 0 or 1 (false or true) stored in an 8 bit container.
+    RS_TYPE_UNSIGNED_5_6_5 = 13, // A 16 bit unsigned integer packing graphical data in 5, 6, and 5 bit sections.
+    RS_TYPE_UNSIGNED_5_5_5_1 = 14, // A 16 bit unsigned integer packing graphical data in 5, 5, 5, and 1 bit sections.
+    RS_TYPE_UNSIGNED_4_4_4_4 = 15, // A 16 bit unsigned integer packing graphical data in 4, 4, 4, and 4 bit sections.
+    RS_TYPE_MATRIX_4X4 = 16, // A 4x4 matrix of 32 bit floats, aligned on a 32 bit boundary.
+    RS_TYPE_MATRIX_3X3 = 17, // A 3x3 matrix of 32 bit floats, aligned on a 32 bit boundary.
+    RS_TYPE_MATRIX_2X2 = 18, // A 2x2 matrix of 32 bit floats, aligned on a 32 bit boundary.
+    RS_TYPE_ELEMENT = 1000, // A handle to an Element.
+    RS_TYPE_TYPE = 1001, // A handle to a Type.
+    RS_TYPE_ALLOCATION = 1002, // A handle to an Allocation.
+    RS_TYPE_SAMPLER = 1003, // A handle to a Sampler.
+    RS_TYPE_SCRIPT = 1004, // A handle to a Script.
+    RS_TYPE_MESH = 1005, // Deprecated.
+    RS_TYPE_PROGRAM_FRAGMENT = 1006, // Deprecated.
+    RS_TYPE_PROGRAM_VERTEX = 1007, // Deprecated.
+    RS_TYPE_PROGRAM_RASTER = 1008, // Deprecated.
+    RS_TYPE_PROGRAM_STORE = 1009, // Deprecated.
+    RS_TYPE_FONT = 1010, // Deprecated.
     RS_TYPE_INVALID = 10000
 } rs_data_type;
 #endif
@@ -175,21 +171,26 @@ typedef enum {
 /*
  * rs_data_kind: Element data kind
  *
- * The special interpretation of the data if required.  This is primarly
- * useful for graphical data.  USER indicates no special interpretation is
- * expected.  PIXEL is used in conjunction with the standard data types for
- * representing texture formats.
+ * This enumeration is primarly useful for graphical data.  It provides additional information to
+ * help interpret the rs_data_type.
+ *
+ * RS_KIND_USER indicates no special interpretation is expected.
+ *
+ * The RS_KIND_PIXEL_* values are used in conjunction with the standard data types for representing
+ * texture formats.
+ *
+ * See the Element.createPixel() method.
  */
 #if (defined(RS_VERSION) && (RS_VERSION >= 16))
 typedef enum {
-    RS_KIND_USER         = 0,
-    RS_KIND_PIXEL_L      = 7,
-    RS_KIND_PIXEL_A      = 8,
-    RS_KIND_PIXEL_LA     = 9,
-    RS_KIND_PIXEL_RGB    = 10,
-    RS_KIND_PIXEL_RGBA   = 11,
-    RS_KIND_PIXEL_DEPTH  = 12,
-    RS_KIND_PIXEL_YUV    = 13,
+    RS_KIND_USER         = 0, // No special interpretation.
+    RS_KIND_PIXEL_L      = 7, // Luminance.
+    RS_KIND_PIXEL_A      = 8, // Alpha.
+    RS_KIND_PIXEL_LA     = 9, // Luminance and Alpha.
+    RS_KIND_PIXEL_RGB    = 10, // Red, Green, Blue.
+    RS_KIND_PIXEL_RGBA   = 11, // Red, Green, Blue, and Alpha.
+    RS_KIND_PIXEL_DEPTH  = 12, // Depth for a depth texture.
+    RS_KIND_PIXEL_YUV    = 13, // Luminance and chrominance.
     RS_KIND_INVALID      = 100
 } rs_data_kind;
 #endif
diff --git a/renderscript/include/rs_quaternion.rsh b/renderscript/include/rs_quaternion.rsh
index 041bdf5..524eeb0 100644
--- a/renderscript/include/rs_quaternion.rsh
+++ b/renderscript/include/rs_quaternion.rsh
@@ -19,6 +19,7 @@
 /*
  * rs_quaternion.rsh: Quaternion Functions
  *
+ * The following functions manipulate quaternions.
  */
 
 #ifndef RENDERSCRIPT_RS_QUATERNION_RSH
@@ -129,7 +130,7 @@ static inline void __attribute__((overloadable))
 }
 
 /*
- * rsQuaternionSet: Create a quarternion
+ * rsQuaternionSet: Create a quaternion
  *
  * Creates a quaternion from its four components or from another quaternion.
  *
@@ -210,8 +211,8 @@ static inline void __attribute__((overloadable))
  *
  * Parameters:
  *   q: Destination quaternion.
- *   scalar: Scalar to multiply the quarternion by.
- *   rhs: Quarternion to multiply the destination quaternion by.
+ *   scalar: Scalar to multiply the quaternion by.
+ *   rhs: Quaternion to multiply the destination quaternion by.
  */
 static inline void __attribute__((overloadable))
     rsQuaternionMultiply(rs_quaternion* q, float scalar) {
diff --git a/renderscript/include/rs_time.rsh b/renderscript/include/rs_time.rsh
index 8ce120f..05987f8 100644
--- a/renderscript/include/rs_time.rsh
+++ b/renderscript/include/rs_time.rsh
@@ -19,9 +19,8 @@
 /*
  * rs_time.rsh: Time Functions and Types
  *
- * The functions below can be used to tell the current clock time and the
- * current system up time.  It's not recommended to call these functions
- * inside of a kernel.
+ * The functions below can be used to tell the current clock time and the current
+ * system up time.  It is not recommended to call these functions inside of a kernel.
  */
 
 #ifndef RENDERSCRIPT_RS_TIME_RSH
@@ -61,8 +60,7 @@ typedef struct {
 /*
  * rsGetDt: Elapsed time since last call
  *
- * Returns the time in seconds since this function was last called in this
- * script.
+ * Returns the time in seconds since this function was last called in this script.
  *
  * Returns: Time in seconds.
  */
@@ -72,11 +70,11 @@ extern float __attribute__((overloadable))
 /*
  * rsLocaltime: Convert to local time
  *
- * Converts the time specified by timer into a rs_tm structure that provides year, month, hour, etc.
- * This value is stored at *local.
+ * Converts the time specified by timer into a rs_tm structure that provides year, month,
+ * hour, etc.  This value is stored at *local.
  *
- * This functions returns the same pointer that is passed as first argument.
- * If the local parameter is NULL, this function does nothing and returns NULL.
+ * This functions returns the same pointer that is passed as first argument.  If the
+ * local parameter is NULL, this function does nothing and returns NULL.
  *
  * Parameters:
  *   local: Pointer to time structure where the local time will be stored.
@@ -90,8 +88,7 @@ extern rs_tm* __attribute__((overloadable))
 /*
  * rsTime: Seconds since January 1, 1970
  *
- * Returns the number of seconds since the Epoch (00:00:00 UTC, January 1,
- * 1970).
+ * Returns the number of seconds since the Epoch (00:00:00 UTC, January 1, 1970).
  *
  * If timer is non-NULL, the result is also stored in the memory pointed to by
  * this variable.
@@ -119,8 +116,7 @@ extern int64_t __attribute__((overloadable))
  *
  * Returns the current system clock (uptime) in nanoseconds.
  *
- * The granularity of the values return by this call may be much
- * larger than a nanosecond.
+ * The granularity of the values return by this call may be much larger than a nanosecond.
  *
  * Returns: Uptime in nanoseconds.
  */
diff --git a/renderscript/include/rs_value_types.rsh b/renderscript/include/rs_value_types.rsh
index 13c0500..9dc7259 100644
--- a/renderscript/include/rs_value_types.rsh
+++ b/renderscript/include/rs_value_types.rsh
@@ -23,21 +23,11 @@
  *
  * RenderScript supports the following scalar numerical types:
  *
- * Integers:
- * - 8 bit: char, int8_t
- * - 16 bit: short, int16_t
- * - 32 bit: int, int32_t
- * - 64 bit: long, long long, int64_t
- *
- * Unsigned integers:
- * - 8 bit: uchar, uint8_t
- * - 16 bit: ushort, uint16_t
- * - 32 bit: uint, uint32_t
- * - 64 bit: ulong, uint64_t
- *
- * Floating point:
- * - 32 bit: float
- * - 64 bit: double
+ *                    8 bits           16 bits            32 bits          64 bits
+ * Integer:           char, int8_t     short, int16_t     int32_t          long, long long, int64_t
+ * Unsigned integer:  uchar, uint8_t   ushort, uint16_t   uint, uint32_t   ulong, uint64_t
+ * Floating point:                                        float            double
+ *
  *
  * Vectors:
  *
@@ -61,30 +51,31 @@
  *   myVar.z == myVar.b == myVar.s2 == myVar.S2
  *   myVar.w == myVar.a == myVar.s3 == myVar.S3
  *
- * Multiple entries of a vector can be accessed at once by using an identifier
- * that is the concatenation of multiple letters or indices.  The resulting vector
- * has a size equal to the number of entries named.
+ * Multiple entries of a vector can be accessed at once by using an identifier that is
+ * the concatenation of multiple letters or indices.  The resulting vector has a size
+ * equal to the number of entries named.
  *
  * With the example above, the middle two entries can be accessed using
  * myVar.yz, myVar.gb, myVar.s12, and myVar.S12.
  *
- * The entries don't have to be contiguous or in increasing order.
- * Entries can even be repeated, as long as we're not trying to assign
- * to it.  You also can't mix the naming styles.
+ * The entries don't have to be contiguous or in increasing order.  Entries can even be
+ * repeated, as long as we're not trying to assign to it.  You also can't mix the naming
+ * styles.
  *
- * Here are examples of what can or can't be done: 
+ * Here are examples of what can or can't be done:
  * float4 v4;
  * float3 v3;
  * float2 v2;
  * v2 = v4.xx; // Valid
  * v3 = v4.zxw; // Valid
  * v3 = v4.bba; // Valid
- * v3 = v4.s034; // Valid
+ * v3 = v4.s032; // Valid
  * v3.s120 = v4.S233; // Valid
  * v4.yz = v3.rg; // Valid
  * v4.yzx = v3.rg; // Invalid: mismatched sizes
  * v4.yzz = v3; // Invalid: z appears twice in an assignment
  * v3 = v3.xas0; // Invalid: can't mix xyzw with rgba nor s0...
+ * v3 = v4.s034; // Invalid: the digit can only be 0, 1, 2, or 3
  *
  *
  * Matrices and Quaternions:
@@ -93,7 +84,8 @@
  * The types are named rs_matrix2x2, rs_matrix3x3, and rs_matrix4x4.  See
  * Matrix Functions for the list of operations.
  *
- * Quaternions are also supported via rs_quaternion.  See Quaterion Functions. for the list of operations.
+ * Quaternions are also supported via rs_quaternion.  See Quaterion Functions for the list
+ * of operations.
  */
 
 #ifndef RENDERSCRIPT_RS_VALUE_TYPES_RSH
@@ -224,240 +216,243 @@ typedef int32_t ssize_t;
 /*
  * float2: Two 32 bit floats
  *
- * Vector version of the basic float type.
- * Provides two float fields packed into a single 64 bit field with 64 bit alignment.
+ * A vector of two floats.  These two floats are packed into a single 64 bit field
+ * with a 64 bit alignment.
+ *
+ * A vector of two floats.  These two floats are packed into a single 64 bit field
+ * with a 64 bit alignment.
  */
 typedef float __attribute__((ext_vector_type(2))) float2;
 
 /*
  * float3: Three 32 bit floats
  *
- * Vector version of the basic float type.
- * Provides three float fields packed into a single 128 bit field with 128 bit alignment.
+ * A vector of three floats.  These three floats are packed into a single 128 bit field
+ * with a 128 bit alignment.
  */
 typedef float __attribute__((ext_vector_type(3))) float3;
 
 /*
  * float4: Four 32 bit floats
  *
- * Vector version of the basic float type.
- * Provides four float fields packed into a single 128 bit field with 128 bit alignment.
+ * A vector of four floats type.  These four floats are packed into a single 128 bit field
+ * with a 128 bit alignment.
  */
 typedef float __attribute__((ext_vector_type(4))) float4;
 
 /*
  * double2: Two 64 bit floats
  *
- * Vector version of the basic double type. Provides two double fields packed
- * into a single 128 bit field with 128 bit alignment.
+ * A vector of two doubles.  These two double fields packed into a single 128 bit field
+ * with a 128 bit alignment.
  */
 typedef double __attribute__((ext_vector_type(2))) double2;
 
 /*
  * double3: Three 64 bit floats
  *
- * Vector version of the basic double type. Provides three double fields packed
- * into a single 256 bit field with 256 bit alignment.
+ * A vector of three doubles.  These three double fields packed into a single 256 bit field
+ * with a 256 bit alignment.
  */
 typedef double __attribute__((ext_vector_type(3))) double3;
 
 /*
  * double4: Four 64 bit floats
  *
- * Vector version of the basic double type. Provides four double fields packed
- * into a single 256 bit field with 256 bit alignment.
+ * A vector of four doubles.  These four double fields packed into a single 256 bit field
+ * with a 256 bit alignment.
  */
 typedef double __attribute__((ext_vector_type(4))) double4;
 
 /*
  * uchar2: Two 8 bit unsigned integers
  *
- * Vector version of the basic uchar type. Provides two uchar fields packed
- * into a single 16 bit field with 16 bit alignment.
+ * A vector of two uchars.  These two uchar fields packed into a single 16 bit field
+ * with a 16 bit alignment.
  */
 typedef uchar __attribute__((ext_vector_type(2))) uchar2;
 
 /*
  * uchar3: Three 8 bit unsigned integers
  *
- * Vector version of the basic uchar type. Provides three uchar fields packed
- * into a single 32 bit field with 32 bit alignment.
+ * A vector of three uchars.  These three uchar fields packed into a single 32 bit field
+ * with a 32 bit alignment.
  */
 typedef uchar __attribute__((ext_vector_type(3))) uchar3;
 
 /*
  * uchar4: Four 8 bit unsigned integers
  *
- * Vector version of the basic uchar type. Provides four uchar fields packed
- * into a single 32 bit field with 32 bit alignment.
+ * A vector of four uchars.  These four uchar fields packed into a single 32 bit field
+ * with a 32 bit alignment.
  */
 typedef uchar __attribute__((ext_vector_type(4))) uchar4;
 
 /*
  * ushort2: Two 16 bit unsigned integers
  *
- * Vector version of the basic ushort type. Provides two ushort fields packed
- * into a single 32 bit field with 32 bit alignment.
+ * A vector of two ushorts.  These two ushort fields packed into a single 32 bit field
+ * with a 32 bit alignment.
  */
 typedef ushort __attribute__((ext_vector_type(2))) ushort2;
 
 /*
  * ushort3: Three 16 bit unsigned integers
  *
- * Vector version of the basic ushort type. Provides three ushort fields packed
- * into a single 64 bit field with 64 bit alignment.
+ * A vector of three ushorts.  These three ushort fields packed into a single 64 bit field
+ * with a 64 bit alignment.
  */
 typedef ushort __attribute__((ext_vector_type(3))) ushort3;
 
 /*
  * ushort4: Four 16 bit unsigned integers
  *
- * Vector version of the basic ushort type. Provides four ushort fields packed
- * into a single 64 bit field with 64 bit alignment.
+ * A vector of four ushorts.  These four ushort fields packed into a single 64 bit field
+ * with a 64 bit alignment.
  */
 typedef ushort __attribute__((ext_vector_type(4))) ushort4;
 
 /*
  * uint2: Two 32 bit unsigned integers
  *
- * Vector version of the basic uint type. Provides two uint fields packed into a
- * single 64 bit field with 64 bit alignment.
+ * A vector of two uints.  These two uints are packed into a single 64 bit field
+ * with a 64 bit alignment.
  */
 typedef uint __attribute__((ext_vector_type(2))) uint2;
 
 /*
  * uint3: Three 32 bit unsigned integers
  *
- * Vector version of the basic uint type. Provides three uint fields packed into
- * a single 128 bit field with 128 bit alignment.
+ * A vector of three uints.  These three uints are packed into a single 128 bit field
+ * with a 128 bit alignment.
  */
 typedef uint __attribute__((ext_vector_type(3))) uint3;
 
 /*
  * uint4: Four 32 bit unsigned integers
  *
- * Vector version of the basic uint type. Provides four uint fields packed into
- * a single 128 bit field with 128 bit alignment.
+ * A vector of four uints.  These four uints are packed into a single 128 bit field
+ * with a 128 bit alignment.
  */
 typedef uint __attribute__((ext_vector_type(4))) uint4;
 
 /*
  * ulong2: Two 64 bit unsigned integers
  *
- * Vector version of the basic ulong type. Provides two ulong fields packed into
- * a single 128 bit field with 128 bit alignment.
+ * A vector of two ulongs.  These two ulongs are packed into a single 128 bit field
+ * with a 128 bit alignment.
  */
 typedef ulong __attribute__((ext_vector_type(2))) ulong2;
 
 /*
  * ulong3: Three 64 bit unsigned integers
  *
- * Vector version of the basic ulong type. Provides three ulong fields packed
- * into a single 256 bit field with 256 bit alignment.
+ * A vector of three ulongs.  These three ulong fields packed into a single 256 bit field
+ * with a 256 bit alignment.
  */
 typedef ulong __attribute__((ext_vector_type(3))) ulong3;
 
 /*
  * ulong4: Four 64 bit unsigned integers
  *
- * Vector version of the basic ulong type. Provides four ulong fields packed
- * into a single 256 bit field with 256 bit alignment.
+ * A vector of four ulongs.  These four ulong fields packed into a single 256 bit field
+ * with a 256 bit alignment.
  */
 typedef ulong __attribute__((ext_vector_type(4))) ulong4;
 
 /*
  * char2: Two 8 bit signed integers
  *
- * Vector version of the basic char type. Provides two char fields packed into a
- * single 16 bit field with 16 bit alignment.
+ * A vector of two chars.  These two chars are packed into a single 16 bit field
+ * with a 16 bit alignment.
  */
 typedef char __attribute__((ext_vector_type(2))) char2;
 
 /*
  * char3: Three 8 bit signed integers
  *
- * Vector version of the basic char type. Provides three char fields packed into
- * a single 32 bit field with 32 bit alignment.
+ * A vector of three chars.  These three chars are packed into a single 32 bit field
+ * with a 32 bit alignment.
  */
 typedef char __attribute__((ext_vector_type(3))) char3;
 
 /*
  * char4: Four 8 bit signed integers
  *
- * Vector version of the basic char type. Provides four char fields packed into
- * a single 32 bit field with 32 bit alignment.
+ * A vector of four chars.  These four chars are packed into a single 32 bit field
+ * with a 32 bit alignment.
  */
 typedef char __attribute__((ext_vector_type(4))) char4;
 
 /*
  * short2: Two 16 bit signed integers
  *
- * Vector version of the basic short type. Provides two short fields packed into
- * a single 32 bit field with 32 bit alignment.
+ * A vector of two shorts.  These two shorts are packed into a single 32 bit field
+ * with a 32 bit alignment.
  */
 typedef short __attribute__((ext_vector_type(2))) short2;
 
 /*
  * short3: Three 16 bit signed integers
  *
- * Vector version of the basic short type. Provides three short fields packed
- * into a single 64 bit field with 64 bit alignment.
+ * A vector of three shorts.  These three short fields packed into a single 64 bit field
+ * with a 64 bit alignment.
  */
 typedef short __attribute__((ext_vector_type(3))) short3;
 
 /*
  * short4: Four 16 bit signed integers
  *
- * Vector version of the basic short type. Provides four short fields packed
- * into a single 64 bit field with 64 bit alignment.
+ * A vector of four shorts.  These four short fields packed into a single 64 bit field
+ * with a 64 bit alignment.
  */
 typedef short __attribute__((ext_vector_type(4))) short4;
 
 /*
  * int2: Two 32 bit signed integers
  *
- * Vector version of the basic int type. Provides two int fields packed into a
- * single 64 bit field with 64 bit alignment.
+ * A vector of two ints.  These two ints are packed into a single 64 bit field
+ * with a 64 bit alignment.
  */
 typedef int __attribute__((ext_vector_type(2))) int2;
 
 /*
  * int3: Three 32 bit signed integers
  *
- * Vector version of the basic int type. Provides three int fields packed into a
- * single 128 bit field with 128 bit alignment.
+ * A vector of three ints.  These three ints are packed into a single 128 bit field
+ * with a 128 bit alignment.
  */
 typedef int __attribute__((ext_vector_type(3))) int3;
 
 /*
  * int4: Four 32 bit signed integers
  *
- * Vector version of the basic int type. Provides two four fields packed into a
- * single 128 bit field with 128 bit alignment.
+ * A vector of four ints.  These two fours are packed into a single 128 bit field
+ * with a 128 bit alignment.
  */
 typedef int __attribute__((ext_vector_type(4))) int4;
 
 /*
  * long2: Two 64 bit signed integers
  *
- * Vector version of the basic long type. Provides two long fields packed into a
- * single 128 bit field with 128 bit alignment.
+ * A vector of two longs.  These two longs are packed into a single 128 bit field
+ * with a 128 bit alignment.
  */
 typedef long __attribute__((ext_vector_type(2))) long2;
 
 /*
  * long3: Three 64 bit signed integers
  *
- * Vector version of the basic long type. Provides three long fields packed into
- * a single 256 bit field with 256 bit alignment.
+ * A vector of three longs.  These three longs are packed into a single 256 bit field
+ * with a 256 bit alignment.
  */
 typedef long __attribute__((ext_vector_type(3))) long3;
 
 /*
  * long4: Four 64 bit signed integers
  *
- * Vector version of the basic long type. Provides four long fields packed into
- * a single 256 bit field with 256 bit alignment.
+ * A vector of four longs.  These four longs are packed into a single 256 bit field
+ * with a 256 bit alignment.
  */
 typedef long __attribute__((ext_vector_type(4))) long4;
 
diff --git a/renderscript/include/rs_vector_math.rsh b/renderscript/include/rs_vector_math.rsh
index 8ad2cbf..a2757e9 100644
--- a/renderscript/include/rs_vector_math.rsh
+++ b/renderscript/include/rs_vector_math.rsh
@@ -19,16 +19,19 @@
 /*
  * rs_vector_math.rsh: Vector Math Functions
  *
- * These functions interpret the input arguments as representation of vectors in n-dimensional space.
+ * These functions interpret the input arguments as representation of vectors in
+ * n-dimensional space.
  *
- * The precision of the mathematical operations is affected by the pragmas
- * rs_fp_relaxed and rs_fp_full.
+ * The precision of the mathematical operations on 32 bit floats is affected by the pragmas
+ * rs_fp_relaxed and rs_fp_full.  See Mathematical Constants and Functions for details.
  *
- * Different precision/speed tradeoffs can be achieved by using three variants
- * of common math functions.  Functions with a name starting with
- * - native_ may have custom hardware implementations with weaker precision,
- * - half_ may perform internal computations using 16 bit floats, and
- * - fast_ are n-dimensional space computations that may use 16 bit floats.
+ * Different precision/speed tradeoffs can be achieved by using variants of the common math
+ * functions.  Functions with a name starting with
+ * - native_: May have custom hardware implementations with weaker precision.  Additionally,
+ *   subnormal values may be flushed to zero, rounding towards zero may be used, and NaN and
+ *   infinity input may not be handled correctly.
+ * - fast_: May perform internal computations using 16 bit floats.  Additionally, subnormal
+ *   values may be flushed to zero, and rounding towards zero may be used.
  *
  */
 
@@ -87,7 +90,8 @@ extern float __attribute__((const, overloadable))
  *
  * Computes the approximate distance between two points.
  *
- * The precision is what would be expected from doing the computation using 16 bit floating point values.
+ * The precision is what would be expected from doing the computation using 16 bit floating
+ * point values.
  *
  * See also distance(), native_distance().
  */
@@ -116,7 +120,8 @@ extern float __attribute__((const, overloadable))
  *
  * Computes the approximate length of a vector.
  *
- * The precision is what would be expected from doing the computation using 16 bit floating point values.
+ * The precision is what would be expected from doing the computation using 16 bit floating
+ * point values.
  *
  * See also length(), native_length().
  */
@@ -145,9 +150,11 @@ extern float __attribute__((const, overloadable))
  *
  * Approximately normalizes a vector.
  *
- * For vectors of size 1, returns -1.f for negative values, 0.f for null values, and 1.f for positive values.
+ * For vectors of size 1, returns -1.f for negative values, 0.f for null values, and 1.f for
+ * positive values.
  *
- * The precision is what would be expected from doing the computation using 16 bit floating point values.
+ * The precision is what would be expected from doing the computation using 16 bit floating
+ * point values.
  *
  * See also normalize(), native_normalize().
  */
@@ -276,7 +283,8 @@ extern float4 __attribute__((const, overloadable))
  *
  * Normalize a vector.
  *
- * For vectors of size 1, returns -1.f for negative values, 0.f for null values, and 1.f for positive values.
+ * For vectors of size 1, returns -1.f for negative values, 0.f for null values, and 1.f for
+ * positive values.
  *
  * See also fast_normalize(), native_normalize().
  */
diff --git a/renderscript/lib/arm/libRSSupport.so b/renderscript/lib/arm/libRSSupport.so
index afefc27..975dc54 100755
--- a/renderscript/lib/arm/libRSSupport.so
+++ b/renderscript/lib/arm/libRSSupport.so
diff --git a/renderscript/lib/arm/libRSSupportIO.so b/renderscript/lib/arm/libRSSupportIO.so
index d2f53fb..c6700be 100755
--- a/renderscript/lib/arm/libRSSupportIO.so
+++ b/renderscript/lib/arm/libRSSupportIO.so
diff --git a/renderscript/lib/arm/libc.so b/renderscript/lib/arm/libc.so
index bcbb1e7..76c73ee 100755
--- a/renderscript/lib/arm/libc.so
+++ b/renderscript/lib/arm/libc.so
diff --git a/renderscript/lib/arm/libm.so b/renderscript/lib/arm/libm.so
index 2e12f88..855e964 100755
--- a/renderscript/lib/arm/libm.so
+++ b/renderscript/lib/arm/libm.so
diff --git a/renderscript/lib/arm/librsjni.so b/renderscript/lib/arm/librsjni.so
index f0ef9b5..baf730b 100755
--- a/renderscript/lib/arm/librsjni.so
+++ b/renderscript/lib/arm/librsjni.so
diff --git a/renderscript/lib/javalib.jar b/renderscript/lib/javalib.jar
index d856e94..0058d42 100644
--- a/renderscript/lib/javalib.jar
+++ b/renderscript/lib/javalib.jar
diff --git a/renderscript/lib/mips/libRSSupport.so b/renderscript/lib/mips/libRSSupport.so
index edb68d9..9b360a3 100755
--- a/renderscript/lib/mips/libRSSupport.so
+++ b/renderscript/lib/mips/libRSSupport.so
diff --git a/renderscript/lib/mips/libRSSupportIO.so b/renderscript/lib/mips/libRSSupportIO.so
index 292fee0..6ccbff3 100755
--- a/renderscript/lib/mips/libRSSupportIO.so
+++ b/renderscript/lib/mips/libRSSupportIO.so
diff --git a/renderscript/lib/mips/libc.so b/renderscript/lib/mips/libc.so
index 89cfc0c..3f37e8b 100755
--- a/renderscript/lib/mips/libc.so
+++ b/renderscript/lib/mips/libc.so
diff --git a/renderscript/lib/mips/libm.so b/renderscript/lib/mips/libm.so
index 29055fb..d275296 100755
--- a/renderscript/lib/mips/libm.so
+++ b/renderscript/lib/mips/libm.so
diff --git a/renderscript/lib/mips/librsjni.so b/renderscript/lib/mips/librsjni.so
index 4d7ed75..bf1cd74 100755
--- a/renderscript/lib/mips/librsjni.so
+++ b/renderscript/lib/mips/librsjni.so
diff --git a/renderscript/lib/x86/libRSSupport.so b/renderscript/lib/x86/libRSSupport.so
index 60288d6..3a07846 100755
--- a/renderscript/lib/x86/libRSSupport.so
+++ b/renderscript/lib/x86/libRSSupport.so
diff --git a/renderscript/lib/x86/libRSSupportIO.so b/renderscript/lib/x86/libRSSupportIO.so
index afd8e79..5f5e49a 100755
--- a/renderscript/lib/x86/libRSSupportIO.so
+++ b/renderscript/lib/x86/libRSSupportIO.so
diff --git a/renderscript/lib/x86/libc.so b/renderscript/lib/x86/libc.so
index 5d785fc..80886c2 100755
--- a/renderscript/lib/x86/libc.so
+++ b/renderscript/lib/x86/libc.so
diff --git a/renderscript/lib/x86/libm.so b/renderscript/lib/x86/libm.so
index ce846a2..fbd135f 100755
--- a/renderscript/lib/x86/libm.so
+++ b/renderscript/lib/x86/libm.so
diff --git a/renderscript/lib/x86/librsjni.so b/renderscript/lib/x86/librsjni.so
index e6c1fd9..54428b8 100755
--- a/renderscript/lib/x86/librsjni.so
+++ b/renderscript/lib/x86/librsjni.so