Why does working with half's brake async function?

[trsh]

thrust::async::for_each(thrust::device,
			thrust::make_counting_iterator<unsigned int>(0),
			thrust::make_counting_iterator<unsigned int>(1),
			[] __device__(const unsigned int idx) {
				half kaka = 1.f;
				while (true) {}
			}
		);

So this blocks the process on CPU and acts as it would be a sync for_each. If I remove the half or replace with float, it works as expected.

[jrhemstad]

Looking at your code and comparing the generated PTX, it looks like when using a type other than __half the compiler just generates an empty kernel, whereas with __half it generates the infinite loop.

https://godbolt.org/z/TGY1qs8x8

An infinite loop results in undefined behavior, so the compiler is free to throw away the kernel entirely, which is what is happening here. I don't know why it doesn't do this for __half.

For reference, the thrust::async stuff is all deprecated.