gpu: generic: sycl: lnorm and bnorm Intel GPU precision issues

kala855 · kala855 · commit db382e06b8ff · 2024-10-09T14:46:53.000+01:00
diff --git a/src/gpu/generic/sycl/ref_batch_normalization.cpp b/src/gpu/generic/sycl/ref_batch_normalization.cpp
@@ -84,9 +84,13 @@ status_t ref_batch_normalization_fwd_t::init(impl::engine_t *engine) {
                 = ::sycl::get_kernel_id<batch_normalization_fwd_kernel_vec_t>();
         CHECK(create_kernel(engine, kid, &kernel_));
     } else {
+        // Enabling the IEEE div compliant implementation
+        setenv("SYCL_PROGRAM_COMPILE_OPTIONS",
+                "-cl-fp32-correctly-rounded-divide-sqrt", 1);
         const auto kid = ::sycl::get_kernel_id<
                 batch_normalization_fwd_kernel_vec_t1>();
         CHECK(create_kernel(engine, kid, &kernel_));
+        unsetenv("SYCL_PROGRAM_COMPILE_OPTIONS");
     }
     return status::success;
 }
diff --git a/src/gpu/generic/sycl/ref_layer_normalizations.cpp b/src/gpu/generic/sycl/ref_layer_normalizations.cpp
@@ -82,9 +82,13 @@ status_t ref_layer_normalization_fwd_t::init(impl::engine_t *engine) {
                 = ::sycl::get_kernel_id<layer_normalization_fwd_kernel_vec_t>();
         CHECK(create_kernel(engine, kid, &kernel_));
     } else {
+        // Enabling the IEEE div compliant implementation
+        setenv("SYCL_PROGRAM_COMPILE_OPTIONS",
+                "-cl-fp32-correctly-rounded-divide-sqrt", 1);
         const auto kid = ::sycl::get_kernel_id<
                 layer_normalization_fwd_kernel_vec1_t>();
         CHECK(create_kernel(engine, kid, &kernel_));
+        unsetenv("SYCL_PROGRAM_COMPILE_OPTIONS");
     }
     return status::success;
 }

Original file line number	Diff line number	Diff line change
`@@ -84,9 +84,13 @@ status_t ref_batch_normalization_fwd_t::init(impl::engine_t *engine) {`
`84`	`84`	`= ::sycl::get_kernel_id<batch_normalization_fwd_kernel_vec_t>();`
`85`	`85`	`CHECK(create_kernel(engine, kid, &kernel_));`
`86`	`86`	`} else {`
	`87`	`+ // Enabling the IEEE div compliant implementation`
	`88`	`+ setenv("SYCL_PROGRAM_COMPILE_OPTIONS",`
	`89`	`+ "-cl-fp32-correctly-rounded-divide-sqrt", 1);`
`87`	`90`	`const auto kid = ::sycl::get_kernel_id<`
`88`	`91`	`batch_normalization_fwd_kernel_vec_t1>();`
`89`	`92`	`CHECK(create_kernel(engine, kid, &kernel_));`
	`93`	`+ unsetenv("SYCL_PROGRAM_COMPILE_OPTIONS");`
`90`	`94`	`}`
`91`	`95`	`return status::success;`
`92`	`96`	`}`
Original file line number	Diff line number	Diff line change
`@@ -82,9 +82,13 @@ status_t ref_layer_normalization_fwd_t::init(impl::engine_t *engine) {`
`82`	`82`	`= ::sycl::get_kernel_id<layer_normalization_fwd_kernel_vec_t>();`
`83`	`83`	`CHECK(create_kernel(engine, kid, &kernel_));`
`84`	`84`	`} else {`
	`85`	`+ // Enabling the IEEE div compliant implementation`
	`86`	`+ setenv("SYCL_PROGRAM_COMPILE_OPTIONS",`
	`87`	`+ "-cl-fp32-correctly-rounded-divide-sqrt", 1);`
`85`	`88`	`const auto kid = ::sycl::get_kernel_id<`
`86`	`89`	`layer_normalization_fwd_kernel_vec1_t>();`
`87`	`90`	`CHECK(create_kernel(engine, kid, &kernel_));`
	`91`	`+ unsetenv("SYCL_PROGRAM_COMPILE_OPTIONS");`
`88`	`92`	`}`
`89`	`93`	`return status::success;`
`90`	`94`	`}`