Fixed Shader Error

 Changes to be committed:
	modified:   src/vk/Mlp.def.cc
	modified:   src/vk/cl/Mlp.cl.c
	modified:   src/vk/cl/Mlp.h
	modified:   src/vk/cl/Slp.cl.c
	modified:   src/vk/cl/Slp.h
	modified:   test/vk/MlpTrainHalf-Primal.c
	modified:   test/vk/SlpPredict-Primal.c
	modified:   test/vk/SlpTrain-Primal.c

Author: Dalmurii

src/vk/Mlp.def.cc

@@ -371,10 +371,6 @@ ae2f_MAC(cllocfloat_t, ) _ae2fVK_AnnMlpMk_imp_V(
 							1, NULL
 							, &(v_mk).m_U1.m_clsrc
 							, ""
-							"-pod-pushconstant "
-							"-cl-fast-relaxed-math "
-							"-cl-single-precision-constant"
-
 							, ae2f_const_cast(
 								char**
 								, ae2f_reinterpret_cast(

src/vk/cl/Mlp.cl.c

@@ -1,6 +1,5 @@
 #pragma OPENCL EXTENSION cl-fast-relaxed-math : enable
-#pragma OPENCL EXTENSION -pod-pushconstant : enable
-#pragma OPENCL EXTENSION -cl-fast-relaxed-math : enable
+#pragma OPENCL EXTENSION pod-pushconstant : enable
 
 #define ae2f_NEED_CLASS 0
 
@@ -41,9 +40,9 @@
 #define llsz		(lsz - 1)
 
 /** @brief lsz * sizeof(uint32_t) */
-#define p_layerszlist	CAST(__global volatile uint32_t*, glob)
+#define p_layerszlist	CAST(__global uint32_t*, glob)
 /** @brief sizeof(ae2f_float_t) * lsz * pgsz */
-#define p_outstream	(CAST(__global volatile host_float_t*, p_layerszlist + lsz))
+#define p_outstream	(CAST(__global host_float_t*, p_layerszlist + lsz))
 #define p_inp		p_outstream
 
 /** @brief sizeof(ae2f_float_t) * pgsz_sqr * llsz */
@@ -112,7 +111,7 @@ typedef const struct sz2_t {
  * ae2f_float_t[Page]:	inp
  * ae2f_float_t[Page]:	out
  * */
-__kernel void kPredict(__global volatile void* glob, __local volatile ae2f_float_t* loc, sz2_t push) {
+__kernel void kPredict(__global void* glob, __local ae2f_float_t* loc, sz2_t push) {
 	const size_t
 		oidx = get_global_id(0)
 		, iidx = get_global_id(1)
@@ -138,7 +137,7 @@ __kernel void kPredict(__global volatile void* glob, __local volatile ae2f_float
  * ae2f_float_t[Page]:	inp
  * ae2f_float_t[Page]:	out
  * */
-__kernel void kPredictStream(__global volatile void* glob, __local volatile ae2f_float_t* loc, const sz2_t push) {
+__kernel void kPredictStream(__global void* glob, __local ae2f_float_t* loc, const sz2_t push) {
 	const size_t
 		oidx = get_global_id(0)
 		, iidx = get_global_id(1)
@@ -194,7 +193,7 @@ typedef char STATIC_ASSERT_LRLSZ_SZ[sizeof(lrlsz_t) ==  sizeof(lrlszel_t) * 4 ?
  * ae2f_float_t[lsz - 1][Page]: OutStream
  * ae2f_float_t[lsz - 1][Page]: DeltaStream
  * */
-__kernel void kFollow(__global volatile void* glob, __local volatile ae2f_float_t* loc, lrlsz_t lr) {
+__kernel void kFollow(__global void* glob, __local ae2f_float_t* loc, lrlsz_t lr) {
 	size_t		lidx = llsz - 1;
 	ae2f_float_t	v_tmp;
 
@@ -312,7 +311,7 @@ __kernel void kFollow(__global volatile void* glob, __local volatile ae2f_float_
  * ae2f_float_t[lsz - 1][Page]: OutStream
  * ae2f_float_t[lsz - 1][Page]: DeltaStream
  * */
-__kernel void kTrainAuto(__global volatile void* glob, __local volatile ae2f_float_t* loc, lrlsz_t lr) {
+__kernel void kTrainAuto(__global void* glob, __local ae2f_float_t* loc, lrlsz_t lr) {
 	size_t		lidx = 0;
 	const size_t
 		oidx = get_global_id(0)

src/vk/cl/Mlp.h

@@ -79,10 +79,7 @@ ae2f_MAC() clMlpGetHD1(
 
 	if((i_oidx) < (i_osz) && (i_iidx) < (i_isz)) {
 		unless((i_oidx)) (r_delta_then)[i_iidx] = 0;
-		atomic_add(
-				&(r_delta_then)[(i_iidx)]
-				, (i_weight)[(i_isz) * (i_oidx) + (i_iidx)] * (i_delta)[i_oidx]
-				);
+		atomic_add(&(r_delta_then)[(i_iidx)], (i_weight)[(i_isz) * (i_oidx) + (i_iidx)] * (i_delta)[i_oidx]);
 	}
 }
 

src/vk/cl/Slp.cl.c

@@ -1,6 +1,5 @@
 #pragma OPENCL EXTENSION cl-fast-relaxed-math : enable
-#pragma OPENCL EXTENSION -pod-pushconstant : enable
-#pragma OPENCL EXTENSION -cl-fast-relaxed-math : enable
+#pragma OPENCL EXTENSION pod-pushconstant : enable
 
 #define ae2f_NEED_CLASS 0
 

src/vk/cl/Slp.h

@@ -25,10 +25,8 @@ ae2f_MAC() clSlpPredict(
 	if((oidx) < (osz) && (iidx) < (isz)) {
 		unless((iidx)) (loc)[oidx] = 0;
 
-		atomic_add(
-				&(loc)[oidx]
-				, (p_weight)[(oidx) * (isz) + (iidx)] * (p_inp)[iidx]
-				);
+		atomic_add((loc)[oidx], (p_weight)[(oidx) * (isz) + (iidx)] * (p_inp)[iidx]);
+		/** (loc)[oidx] += (p_weight)[(oidx) * (isz) + (iidx)] * (p_inp)[iidx]; */
 
 		unless(iidx) {
 			(loc)[oidx] += (p_bias)[oidx];

test/vk/MlpTrainHalf-Primal.c

@@ -13,7 +13,7 @@ static void	LossDummy(ae2f_float* a, const ae2f_float* b, const ae2f_float* c, s
 	return;
 }
 
-static size_t	s_lpModelLen[] = {2, 10, 3, 12, 1};
+static size_t	s_lpModelLen[] = {2, 4, 3, 12, 1};
 
 #define	NULL_GENERIC(T)	((T*)NULL)
 

test/vk/SlpPredict-Primal.c

@@ -48,7 +48,7 @@ int main() {
 			, vkphydevmemprops
 			, NULL
 			, 
-			"#define CL_Q 1\n"
+			"#define CL_Q 0\n"
 			"#define LOSS_DERIV(r, y, y_desired, i, c)  \n"
 			"#define ACT(r, y, i, c) { (*(r) = (1.0 / (1.0 + exp(-y[i])))); } \n"
 			, "/** This is also a comment */\n"

test/vk/SlpTrain-Primal.c

@@ -33,14 +33,14 @@ int main() {
 	__ae2fVK_AnnSlpMk_imp(
 			mk,
 			0, 0, 0
-			, 400, 100, 0, 0
+			, 1000, 2, 0, 0
 			, ActDummy, ActDummy, LossDeriv
 			, 0.1, 0.1
 			, vkdev
 			, vkphydevmemprops
 			, NULL
 			, 
-			"#define CL_Q 1\n"
+			"#define CL_Q 0\n"
 			"#define ACT_DERIV(r, x, i, c) *(r) = (((x)[i] + 1e-7) * (1.0 - (x)[i] - 1e-7)); \n"
 			"#define ACT(r, x, i, c) *(r) = (1.0 / (1.0 + exp(-(x[i])))); \n"
 			"#define LOSS_DERIV(r, o, t, i, c) *(r) = ((o)[i] - (t)[i]) / (c);\n"