31 lines
1.2 KiB
Plaintext
31 lines
1.2 KiB
Plaintext
// No host SDK, CRT, or device library dependency. Compiled only for gfx1201.
|
|
extern "C" __attribute__((global)) void d4r_ptx_pattern(
|
|
unsigned int* output, unsigned int count, unsigned int seed)
|
|
{
|
|
const unsigned int i = __builtin_amdgcn_workgroup_id_x() * 128u +
|
|
__builtin_amdgcn_workitem_id_x();
|
|
if (i < count)
|
|
output[i] = (i * 1664525u + seed) ^ 0xa5a55a5au;
|
|
}
|
|
|
|
extern "C" __attribute__((global)) void d4r_probe_pattern(
|
|
unsigned int* output, unsigned int count, unsigned int seed)
|
|
{
|
|
const unsigned int i = __builtin_amdgcn_workgroup_id_x() * 128u +
|
|
__builtin_amdgcn_workitem_id_x();
|
|
if (i < count)
|
|
output[i] = (i * 1664525u + seed) ^ 0xa5a55a5au;
|
|
}
|
|
|
|
extern "C" __attribute__((global)) void d4r_probe_interop(
|
|
unsigned int* shared, unsigned int count, unsigned int seed, unsigned int* errors)
|
|
{
|
|
const unsigned int i = __builtin_amdgcn_workgroup_id_x() * 128u +
|
|
__builtin_amdgcn_workitem_id_x();
|
|
if (i < count) {
|
|
const unsigned int expected = (i * 1664525u + seed) ^ 0xa5a55a5au;
|
|
if (shared[i] != expected) __atomic_fetch_add(errors, 1u, __ATOMIC_RELAXED);
|
|
shared[i] ^= 0x3579bdf1u;
|
|
}
|
|
}
|