From 8562b5a04f2fd2e39b7ddc0578abc2dfe377f774 Mon Sep 17 00:00:00 2001
From: Diogo <dgcruz983@gmail.com>
Date: Thu, 20 Jul 2023 17:20:05 -0400
Subject: [PATCH] fixes error when trying to convert float4 -> half4 (#1300)

---
 tinygrad/runtime/ops_cuda.py | 1 +
 1 file changed, 1 insertion(+)

diff --git a/tinygrad/runtime/ops_cuda.py b/tinygrad/runtime/ops_cuda.py
index ceaeb03f65..4dcd5831c7 100644
--- a/tinygrad/runtime/ops_cuda.py
+++ b/tinygrad/runtime/ops_cuda.py
@@ -88,6 +88,7 @@ class CUDACodegen(CStyleCodegen):
       #include <cuda_fp16.h>
       struct __align__(8) half4 {
         half2 x, y;
+        __device__ __forceinline__ explicit half4(const float4& a): x(make_half2(__float2half(a.x), __float2half(a.y))), y(make_half2(__float2half(a.z),__float2half(a.w))) {}
         __device__ __forceinline__ explicit operator float4() const {return make_float4(__half2float(x.x), __half2float(x.y), __half2float(y.x), __half2float(y.y)); }
       };
     """)