c11_atomics: unify host half representation and conversion with wrapper class

[Nuullll]

Introduce HostHalf wrapper class to eliminate explicit cl_half_from_float
and cl_half_to_float conversions throughout the test code. The wrapper
provides semantic value constructors/operators and automatic conversions,
simplifying half-precision arithmetic operations.

Key improvements:

• HostHalf class with operator overloading for arithmetic and comparisons
• Type traits is_host_atomic_fp_v and is_host_fp_v for generic FP handling
• Unified floating-point atomic operations (add/sub/min/max/exchange)
• Removed 300+ lines of half-specific conditional branches
• Consistent calculation for all FP types

[bashbaug]

Hi, I'm having a hard time connecting the dots together. Can you please provide a test command line that I can run that demonstrates the problem? Thanks!

[Nuullll]

Hi, I'm having a hard time connecting the dots together. Can you please provide a test command line that I can run that demonstrates the problem? Thanks!

Sorry for the late response - I was distracted by other work.

The issue can be reproduced by running:

./OpenCL-CTS/build/test_conformance/c11_atomics/test_c11_atomics svm_atomic_store

with Intel OpenCL CPU implementation (which has fp16 support enabled).

Root cause

In the OpenCL kernel:

#pragma OPENCL EXTENSION cl_khr_fp16 : enable
__kernel void test_atomic_kernel(uint threadCount, uint numDestItems, volatile __global atomic_half *destMemory, __global half *oldValues)
{
  uint  tid = get_global_id(0);
  atomic_store_explicit(&destMemory[tid], tid, memory_order_relaxed, memory_scope_all_devices);
}

When uint tid is stored as a half value, the OpenCL device properly converts the integer to IEEE 754 half-precision floating point format. However, in the host reference code, the original implementation simply cast cl_uint directly to cl_half:

host_atomic_store(&destMemory[tid], (HostDataType)tid, MemoryOrder());

cl_half is defined as uint16_t in OpenCL Header, a direct cast from cl_uint to cl_half just truncates the integer bits, rather than performing proper float-to-half conversion.

[bashbaug]

Ah, gotcha, thank you for the explanation!

FWIW, it was the static cast to a float in the snippet below that looked questionable (note: v is a uint)::

            // For half types, convert from float to proper half-precision bit
            // pattern
            return cl_half_from_float(static_cast<float>(v), gHalfRoundingMode);

But, it looks like this is matching what the kernel is doing (convert the uint to a half), so it's correct.

Interestingly, there does not appear to be a problem with our GPU device, which does not support SVM atomics but does support fp16 atomic load and store. I am running:

./test_conformance/c11_atomics/test_c11_atomics atomic_store

I'll add "focused review" and we'll see if we can get this merged next week.

[bashbaug]

Discussed in the November 4th teleconference. Will merge after @shajder 's review.

[Nuullll]

Interestingly, there does not appear to be a problem with our GPU device, which does not support SVM atomics but does support fp16 atomic load and store. I am running:

./test_conformance/c11_atomics/test_c11_atomics atomic_store

@bashbaug Thanks for pointing this out. atomic_store passes for both CPU and GPU, because by default it does not use HostFunction for verification:

OpenCL-CTS/test_conformance/c11_atomics/common.h

Lines 327 to 333 in e641de9

	virtual cl_uint MaxHostThreads()
	{
	if (UseSVM() || gHost)
	return MAX_HOST_THREADS;
	else
	return 0;
	}

OpenCL-CTS/test_conformance/c11_atomics/common.h

Line 1141 in e641de9

hostThreadCount = MaxHostThreads();

OpenCL-CTS/test_conformance/c11_atomics/common.h

Lines 1426 to 1428 in e641de9

	if (hostThreadCount > 0)
	ThreadPool_Do(HostThreadFunction, hostThreadCount,
	&hostThreadContexts[0]);

OpenCL-CTS/test_conformance/c11_atomics/common.h

Lines 146 to 155 in e641de9

	static cl_int HostThreadFunction(cl_uint job_id, cl_uint thread_id,
	void *userInfo)
	{
	THostThreadContext *threadContext =
	((THostThreadContext *)userInfo) + job_id;
	threadContext->test->HostFunction(
	threadContext->tid, threadContext->threadCount,
	threadContext->destMemory, threadContext->oldValues);
	return 0;
	}

OpenCL-CTS/test_conformance/c11_atomics/test_atomics.cpp

Lines 87 to 92 in e641de9

	virtual void HostFunction(cl_uint tid, cl_uint threadCount,
	volatile HostAtomicType *destMemory,
	HostDataType *oldValues)
	{
	host_atomic_store(&destMemory[tid], (HostDataType)tid, MemoryOrder());
	}

The same issue can be exposed on CPU if we force using host threads for verification:

./test_conformance/c11_atomics/test_c11_atomics atomic_store -host

[Nuullll]

I've refactored the host half handling with a new wrapper class HostHalf so that we don't need much extra caution for the half conversions everywhere.

[bashbaug]

@Nuullll would you mind resolving the merge conflicts? I took a look and they don't look too bad, but they aren't quite something I'm comfortable resolving via the web UI. Thanks!

[Nuullll]

This assertion indeed helped capture the kahan_sub implementation bug in #2368. See the fix here: https://github.com/KhronosGroup/OpenCL-CTS/pull/2503/files#diff-aac1c2081764e51ba19f3e93970d6de6c44ca0b530316e76acddfa4bb4a1829cR1533

[Nuullll]

(t-sum) == -y, so compensation == -2y here, causing sum to explode and overflow to infinity.

[shajder]

Thanks, good catch but this is still not quite correct. It should be:

double y = - num - compensation;
double t = sum + y;
compensation = (t - sum) - y;

[Nuullll]

@shajder @bashbaug I've resolved the conflict and fixed a minor issue of #2368 , please take another look. Thanks!

[shajder]

Thanks, good catch but this is still not quite correct. It should be:

double y = - num - compensation;
double t = sum + y;
compensation = (t - sum) - y;

[shajder]

@Nuullll We would like to merge this PR shortly. I’d appreciate it if you could apply any upcoming corrections in separate PRs. Thanks!