# How to use half2 in customer kernels?

**URL:** <https://discuss.pytorch.org/t/how-to-use-half2-in-customer-kernels/99008>\
**Category:** mixed-precision\
**Created:** [October 11, 2020, 9:12pm UTC](https://discuss.pytorch.org/t/how-to-use-half2-in-customer-kernels/99008 "2020-10-11T21:12:18Z")\
**Posts on this page:** 3\
**Page:** 1

<div class="post-metadata">

**Author:** ![hughperkins](https://discuss.pytorch.org/user_avatar/discuss.pytorch.org/hughperkins/32/915_2.png) [@hughperkins](https://discuss.pytorch.org/u/hughperkins)\
**Post date:** [October 11, 2020, 9:12pm UTC](https://discuss.pytorch.org/t/how-to-use-half2-in-customer-kernels/99008/1 "2020-10-11T21:12:18Z")

</div>

It looks like custom kernels use at::Half instead of half? And use of half itself doesnt seem to work? eg casts between half and float fail? I’m guessing these cast failures are something to do with the definition of **CUDA\_NO\_HALF\_OPERATORS** and **CUDA\_NO\_HALF2\_OPERATORS** in the compilation step?

Then, I figured that to use half2, I would use eg at::Half2, but this seems not to exist?

- why do we use at::Half instead of half?
- why are half casts etc removed?
- what should we do if we want to use half2? (or, is there some way to get the compiler to handle half2 vectorization for us, just as g++ handles SSE vecotrization automatically for us in cpu code?)

---

<div class="post-metadata">

**Author:** ![ptrblck](https://discuss.pytorch.org/user_avatar/discuss.pytorch.org/ptrblck/32/1823_2.png) [@ptrblck](https://discuss.pytorch.org/u/ptrblck)\
**Post date:** [October 12, 2020, 12:48am UTC](https://discuss.pytorch.org/t/how-to-use-half2-in-customer-kernels/99008/2 "2020-10-12T00:48:24Z")

</div>

Some questions are already answered in the [cross post](https://discuss.pytorch.org/t/cuda-no-half2-operators-for-cuda-9-2/18365/4).

> [@hughperkins](#):
>
> what should we do if we want to use half2?

You should be able to use `__half2` in your CUDA code as seen e.g. [here](https://github.com/NVIDIA/apex/blob/325f5a0bec542701edba1628ad34f3b2ea47c556/csrc/layer_norm_cuda_kernel.cu#L174).

---

<div class="post-metadata">

**Author:** ![hughperkins](https://discuss.pytorch.org/user_avatar/discuss.pytorch.org/hughperkins/32/915_2.png) [@hughperkins](https://discuss.pytorch.org/u/hughperkins)\
**Post date:** [October 12, 2020, 1:54am UTC](https://discuss.pytorch.org/t/how-to-use-half2-in-customer-kernels/99008/3 "2020-10-12T01:54:44Z")

</div>

sure, but there are no operators. Like eg for sigmoid I have to write:

```auto
__h2div(one, (__hadd2(one, h2exp(__hneg2(x)))))

```

instead of simply:

```auto
one / (one + hexp(-x))

```