diff --git a/tinygrad/runtime/ops_hip.py b/tinygrad/runtime/ops_hip.py index 1ea4bb6f76..0fc0331ce7 100644 --- a/tinygrad/runtime/ops_hip.py +++ b/tinygrad/runtime/ops_hip.py @@ -139,7 +139,7 @@ typedef float float8 __attribute__((ext_vector_type(8))); typedef _Float16 half16 __attribute__((ext_vector_type(16))); extern "C" __global__ """, launch_bounds=True, - smem_prefix = "__shared__ ", barrier = "__syncthreads();", float4 = "make_float4", uses_vload=True, uses_ptr_arithmetic=True, arg_int_prefix = "const int", + smem_prefix = "__shared__ ", smem_prefix_for_cast=False, barrier = "__syncthreads();", float4 = "make_float4", uses_vload=True, uses_ptr_arithmetic=True, arg_int_prefix = "const int", half_prekernel = "#include \nusing half4 = HIP_vector_type;" + """ __device__ float vload_half(size_t offset, const half *p) { return (float)*(p + offset); } __device__ float2 vload_half2(size_t offset, const half *p) { return make_float2((float)*(p + offset*2), (float)*(p + offset*2 + 1)); }