From b2b5756d86670b787dac440020922eeb5e8a1442 Mon Sep 17 00:00:00 2001
From: AlexeyAB <alexeyab84@gmail.com>
Date: Thu, 22 Feb 2018 20:52:43 +0000
Subject: [PATCH] Added __float2half_rn() and __half2float()

---
 src/convolutional_kernels.cu |    4 ++--
 1 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/src/convolutional_kernels.cu b/src/convolutional_kernels.cu
index 135a2ea..3918e5e 100644
--- a/src/convolutional_kernels.cu
+++ b/src/convolutional_kernels.cu
@@ -77,7 +77,7 @@
 __global__ void cuda_f32_to_f16(float* input_f32, size_t size, half *output_f16)
 {
 	int idx = blockIdx.x * blockDim.x + threadIdx.x;
-	if (idx < size) output_f16[idx] = input_f32[idx];
+	if (idx < size) output_f16[idx] = __float2half_rn(input_f32[idx]);
 }
 
 void cuda_convert_f32_to_f16(float* input_f32, size_t size, half *output_f16) {
@@ -87,7 +87,7 @@
 __global__ void cuda_f16_to_f32(half* input_f16, size_t size, float *output_f32)
 {
 	int idx = blockIdx.x * blockDim.x + threadIdx.x;
-	if (idx < size) output_f32[idx] = input_f16[idx];
+	if (idx < size) output_f32[idx] = __half2float(input_f16[idx]);
 }
 
 void cuda_convert_f16_to_f32(half* input_f16, size_t size, float *output_f32) {

--
Gitblit v1.10.0