Attachment #853184 for bug #895286

Lines 99-106 The following table lists the available functions along with a brief description Link Here

(-)a/doc/OCKL.md (-2 / +2 lines)
99	\| `int __ockl_mul24_i32(int,int);` \| Multiply assuming operands fit in 24 bits \|	99	\| `int __ockl_mul24_i32(int,int);` \| Multiply assuming operands fit in 24 bits \|
100	\| `uint __ockl_mul24_u32(uint,uint);` \| \|	100	\| `uint __ockl_mul24_u32(uint,uint);` \| \|
101	\| - \| \|	101	\| - \| \|
102	\| `ulong __ockl_cyclectr_u64(void);` \| Current value of free running 64-bit clock counter \|	102	\| `ulong __ockl_memtime_u64(void);` \| Current value of free running 64-bit clock counter \|
103	\| `ulong __ockl_steadyctr_u64(void);` \| Current value of constant speed 64-bit clock counter \|	103	\| `ulong __ockl_memrealtime_u64(void);` \| Current value of constant speed 64-bit clock counter \|
104	\| - \| \|	104	\| - \| \|
105	\| `uint __ockl_activelane_u32(void);` \| Index of currently lane counting only active lanes in wavefront \|	105	\| `uint __ockl_activelane_u32(void);` \| Index of currently lane counting only active lanes in wavefront \|
106	\| - \| \|	106	\| - \| \|

Lines 143-151 DECL_OCKL_NULLARY_U32(activelane) Link Here

(-)a/ockl/inc/ockl.h (-3 lines)
143		143
144	DECL_OCKL_NULLARY_U64(memtime)	144	DECL_OCKL_NULLARY_U64(memtime)
145	DECL_OCKL_NULLARY_U64(memrealtime)	145	DECL_OCKL_NULLARY_U64(memrealtime)
146	DECL_OCKL_NULLARY_U64(cyclectr)
147	DECL_OCKL_NULLARY_U64(steadyctr)
148
149		146
150	extern half OCKL_MANGLE_T(wfred_add,f16)(half x);	147	extern half OCKL_MANGLE_T(wfred_add,f16)(half x);
151	extern float OCKL_MANGLE_T(wfred_add,f32)(float x);	148	extern float OCKL_MANGLE_T(wfred_add,f32)(float x);

Lines 201-206 get_heap_ptr(void) { Link Here

(-)a/ockl/src/dm.cl (-4 / +11 lines)
201	}	201	}
202	}	202	}
203		203
		204	// realtime
		205	__attribute__((target("s-memrealtime"))) static ulong
		206	realtime(void)
		207	{
		208	return __builtin_amdgcn_s_memrealtime();
		209	}
		210
204	// The actual number of blocks in a slab with blocks of kind k	211	// The actual number of blocks in a slab with blocks of kind k
205	static uint	212	static uint
206	num_blocks(kind_t k)	213	num_blocks(kind_t k)
Lines 466-472 new_slab_wait(__global heap_t *hp, kind_t k) Link Here
466	uint aid = __ockl_activelane_u32();	473	uint aid = __ockl_activelane_u32();
467	if (aid == 0) {	474	if (aid == 0) {
468	ulong expected = AL(&hp->salloc_time[k].value, memory_order_relaxed);	475	ulong expected = AL(&hp->salloc_time[k].value, memory_order_relaxed);
469	ulong now = __ockl_steadyctr_u64();	476	ulong now = realtime();
470	ulong dt = now - expected;	477	ulong dt = now - expected;
471	if (dt < SLAB_TICKS)	478	if (dt < SLAB_TICKS)
472	__ockl_rtcwait_u32(SLAB_TICKS - (uint)dt);	479	__ockl_rtcwait_u32(SLAB_TICKS - (uint)dt);
Lines 480-486 grow_recordable_wait(__global heap_t *hp, kind_t k) Link Here
480	uint aid = __ockl_activelane_u32();	487	uint aid = __ockl_activelane_u32();
481	if (aid == 0) {	488	if (aid == 0) {
482	ulong expected = AL(&hp->grow_time[k].value, memory_order_relaxed);	489	ulong expected = AL(&hp->grow_time[k].value, memory_order_relaxed);
483	ulong now = __ockl_steadyctr_u64();	490	ulong now = realtime();
484	ulong dt = now - expected;	491	ulong dt = now - expected;
485	if (dt < GROW_TICKS)	492	if (dt < GROW_TICKS)
486	__ockl_rtcwait_u32(GROW_TICKS - (uint)dt);	493	__ockl_rtcwait_u32(GROW_TICKS - (uint)dt);
Lines 540-546 try_grow_num_recordable_slabs(__global heap_t *hp, kind_t k) Link Here
540	uint ret = GROW_BUSY;	547	uint ret = GROW_BUSY;
541	if (aid == 0) {	548	if (aid == 0) {
542	ulong expected = AL(&hp->grow_time[k].value, memory_order_relaxed);	549	ulong expected = AL(&hp->grow_time[k].value, memory_order_relaxed);
543	ulong now = __ockl_steadyctr_u64();	550	ulong now = realtime();
544	if (now - expected >= GROW_TICKS &&	551	if (now - expected >= GROW_TICKS &&
545	ACE(&hp->grow_time[k].value, &expected, now, memory_order_relaxed))	552	ACE(&hp->grow_time[k].value, &expected, now, memory_order_relaxed))
546	ret = GROW_FAILURE;	553	ret = GROW_FAILURE;
Lines 687-693 try_allocate_new_slab(__global heap_t *hp, kind_t k) Link Here
687		694
688	if (aid == 0) {	695	if (aid == 0) {
689	ulong expected = AL(&hp->salloc_time[k].value, memory_order_relaxed);	696	ulong expected = AL(&hp->salloc_time[k].value, memory_order_relaxed);
690	ulong now = __ockl_steadyctr_u64();	697	ulong now = realtime();
691	if (now - expected >= SLAB_TICKS &&	698	if (now - expected >= SLAB_TICKS &&
692	ACE(&hp->salloc_time[k].value, &expected, now, memory_order_relaxed))	699	ACE(&hp->salloc_time[k].value, &expected, now, memory_order_relaxed))
693	ret = (__global sdata_t *)0;	700	ret = (__global sdata_t *)0;




 * License. See LICENSE.TXT for details.
 *===------------------------------------------------------------------------*/

#include "oclc.h"
#include "ockl.h"

__attribute__((target("s-memrealtime"))) static ulong
mem_realtime(void)
{
    return __builtin_amdgcn_s_memrealtime();
}

__attribute__((target("gfx11-insts"))) static ulong
msg_realtime(void)
{
    return __builtin_amdgcn_s_sendmsg_rtnl(0x83);
}

// Deprecated
__attribute__((target("s-memtime-inst"))) ulong
OCKL_MANGLE_U64(memtime)(void)
{
    return __builtin_amdgcn_s_memtime();
}

__attribute__((target("s-memrealtime"))) ulong
ulong
OCKL_MANGLE_U64(memrealtime)(void)
{
    return __builtin_amdgcn_s_memrealtime();
}

ulong
OCKL_MANGLE_U64(cyclectr)(void)
{
    return __builtin_readcyclecounter();
}

ulong
OCKL_MANGLE_U64(steadyctr)(void)
{
    if (__oclc_ISA_version >= 11000) {
        return msg_realtime();
    } else {
        return mem_realtime();
    }
}





#include "ockl.h"
#include "oclc.h"

__attribute__((target("s-memrealtime"))) void
OCKL_MANGLE_T(rtcwait,u32)(uint ticks)
{
    ulong now = __builtin_amdgcn_s_memrealtime();
    ulong end = now + __builtin_amdgcn_readfirstlane(ticks);

    if (__oclc_ISA_version >= 9000) {
        while (end > now + 1625) {
            __builtin_amdgcn_s_sleep(127);
            now = __builtin_amdgcn_s_memrealtime();
        }

        while (end > now + 806) {
            __builtin_amdgcn_s_sleep(63);
            now = __builtin_amdgcn_s_memrealtime();
        }

        while (end > now + 396) {
            __builtin_amdgcn_s_sleep(31);
            now = __builtin_amdgcn_s_memrealtime();
        }
    }

    while (end > now + 192) {
        __builtin_amdgcn_s_sleep(15);
        now = __builtin_amdgcn_s_memrealtime();
    }

    while (end > now + 89) {
        __builtin_amdgcn_s_sleep(7);
        now = __builtin_amdgcn_s_memrealtime();
    }

    while (end > now + 38) {
        __builtin_amdgcn_s_sleep(3);
        now = __builtin_amdgcn_s_memrealtime();
    }

    while (end > now) {
        __builtin_amdgcn_s_sleep(1);
        now = __builtin_amdgcn_s_memrealtime();
    }
}

- 

Return to bug 895286

Lines 5-52 Link Here

(-)a/ockl/src/mtime.cl (-33 / +2 lines)
5	* License. See LICENSE.TXT for details.	5	* License. See LICENSE.TXT for details.
6	===------------------------------------------------------------------------/	6	===------------------------------------------------------------------------/
7		7
8	#include "oclc.h"
9	#include "ockl.h"	8	#include "ockl.h"
10		9
11	__attribute__((target("s-memrealtime"))) static ulong
12	mem_realtime(void)
13	{
14	return __builtin_amdgcn_s_memrealtime();
15	}
16
17	__attribute__((target("gfx11-insts"))) static ulong
18	msg_realtime(void)
19	{
20	return __builtin_amdgcn_s_sendmsg_rtnl(0x83);
21	}
22
23	// Deprecated
24	__attribute__((target("s-memtime-inst"))) ulong	10	__attribute__((target("s-memtime-inst"))) ulong
25	OCKL_MANGLE_U64(memtime)(void)	11	OCKL_MANGLE_U64(memtime)(void)
26	{	12	{
27	return __builtin_amdgcn_s_memtime();	13	return __builtin_amdgcn_s_memtime();
28	}	14	}
29		15
30	// Deprecated	16	__attribute__((target("s-memrealtime"))) ulong
31	ulong
32	OCKL_MANGLE_U64(memrealtime)(void)	17	OCKL_MANGLE_U64(memrealtime)(void)
33	{	18	{
34	return mem_realtime();	19	return __builtin_amdgcn_s_memrealtime();
35	}
36
37	ulong
38	OCKL_MANGLE_U64(cyclectr)(void)
39	{
40	return __builtin_readcyclecounter();
41	}
42
43	ulong
44	OCKL_MANGLE_U64(steadyctr)(void)
45	{
46	if (__oclc_ISA_version >= 11000) {
47	return msg_realtime();
48	} else {
49	return mem_realtime();
50	}
51	}	20	}
52		21

Lines 10-56 Link Here

(-)a/ockl/src/wait.cl (-10 / +9 lines)
10	#include "ockl.h"	10	#include "ockl.h"
11	#include "oclc.h"	11	#include "oclc.h"
12		12
13	void	13	__attribute__((target("s-memrealtime"))) void
14	OCKL_MANGLE_T(rtcwait,u32)(uint ticks)	14	OCKL_MANGLE_T(rtcwait,u32)(uint ticks)
15	{	15	{
16	ulong now = __ockl_steadyctr_u64();	16	ulong now = __builtin_amdgcn_s_memrealtime();
17	ulong end = now + __builtin_amdgcn_readfirstlane(ticks);	17	ulong end = now + __builtin_amdgcn_readfirstlane(ticks);
18		18
19	if (__oclc_ISA_version >= 9000) {	19	if (__oclc_ISA_version >= 9000) {
20	while (end > now + 1625) {	20	while (end > now + 1625) {
21	__builtin_amdgcn_s_sleep(127);	21	__builtin_amdgcn_s_sleep(127);
22	now = __ockl_steadyctr_u64();	22	now = __builtin_amdgcn_s_memrealtime();
23	}	23	}
24		24
25	while (end > now + 806) {	25	while (end > now + 806) {
26	__builtin_amdgcn_s_sleep(63);	26	__builtin_amdgcn_s_sleep(63);
27	now = __ockl_steadyctr_u64();	27	now = __builtin_amdgcn_s_memrealtime();
28	}	28	}
29		29
30	while (end > now + 396) {	30	while (end > now + 396) {
31	__builtin_amdgcn_s_sleep(31);	31	__builtin_amdgcn_s_sleep(31);
32	now = __ockl_steadyctr_u64();	32	now = __builtin_amdgcn_s_memrealtime();
33	}	33	}
34	}	34	}
35		35
36	while (end > now + 192) {	36	while (end > now + 192) {
37	__builtin_amdgcn_s_sleep(15);	37	__builtin_amdgcn_s_sleep(15);
38	now = __ockl_steadyctr_u64();	38	now = __builtin_amdgcn_s_memrealtime();
39	}	39	}
40		40
41	while (end > now + 89) {	41	while (end > now + 89) {
42	__builtin_amdgcn_s_sleep(7);	42	__builtin_amdgcn_s_sleep(7);
43	now = __ockl_steadyctr_u64();	43	now = __builtin_amdgcn_s_memrealtime();
44	}	44	}
45		45
46	while (end > now + 38) {	46	while (end > now + 38) {
47	__builtin_amdgcn_s_sleep(3);	47	__builtin_amdgcn_s_sleep(3);
48	now = __ockl_steadyctr_u64();	48	now = __builtin_amdgcn_s_memrealtime();
49	}	49	}
50		50
51	while (end > now) {	51	while (end > now) {
52	__builtin_amdgcn_s_sleep(1);	52	__builtin_amdgcn_s_sleep(1);
53	now = __ockl_steadyctr_u64();	53	now = __builtin_amdgcn_s_memrealtime();
54	}	54	}
55	}	55	}
56		56
57	-