yum-archive/TaSTT-Whisper

High-performance GPGPU inference of OpenAI's Whisper automatic speech recognition (ASR) model

git clone https://git.yummers.dev/yum-archive/TaSTT-Whisper

KonstantinSource codes8c4603c

master
branchesmasterall refs →
3.7 KiB139 linesraw
1groupshared float sharedAccumulators[ 32 ];
2
3// Compute horisontal sum of the numbers. The result is only correct on the thread #0 of the group.
4void horizontalSum( const uint thread, inout float sum )
5{
6	sharedAccumulators[ thread ] = sum;
7	for( uint i = 16; i > 1; i /= 2 )
8	{
9		GroupMemoryBarrierWithGroupSync();
10		if( thread < i )
11		{
12			sum += sharedAccumulators[ thread + i ];
13			sharedAccumulators[ thread ] = sum;
14		}
15	}
16	GroupMemoryBarrierWithGroupSync();
17	if( 0 == thread )
18		sum += sharedAccumulators[ 1 ];
19}
20
21// Compute horisontal sum of the numbers, and broadcast to all threads of the group.
22void horizontalSumBroadcast( const uint thread, inout float sum )
23{
24	horizontalSum( thread, sum );
25	if( 0 == thread )
26		sharedAccumulators[ 0 ] = sum;
27	GroupMemoryBarrierWithGroupSync();
28	sum = sharedAccumulators[ 0 ];
29}
30
31// Compute horisontal sum of the numbers, in the order equal to the CPU-running dot product implementation.
32// The result is only correct on the thread #0 of the group.
33void horizontalSumCompat( const uint thread, inout float sum )
34{
35	sharedAccumulators[ thread ] = sum;
36	GroupMemoryBarrierWithGroupSync();
37
38	if( 0 == ( thread & 8 ) )
39	{
40		// This runs on threads [ 0 .. 7 ] and [ 16 .. 23 ]
41		// sum01 = _mm256_add_ps( sum0, sum1 );
42		// sum23 = _mm256_add_ps( sum2, sum3 );
43		sum += sharedAccumulators[ thread + 8 ];
44		sharedAccumulators[ thread ] = sum;
45	}
46
47	GroupMemoryBarrierWithGroupSync();
48	if( thread < 8 )
49	{
50		// This runs on threads [ 0 .. 7 ]
51		// sum0123 = _mm256_add_ps( sum01, sum23 );
52		sum += sharedAccumulators[ thread + 16 ];
53		sharedAccumulators[ thread ] = sum;
54	}
55
56	GroupMemoryBarrierWithGroupSync();
57	if( thread < 4 )
58	{
59		// const __m128 r4 = _mm_add_ps( _mm256_castps256_ps128( sum0123 ), _mm256_extractf128_ps( sum0123, 1 ) );
60		sum += sharedAccumulators[ thread + 4 ];
61		sharedAccumulators[ thread ] = sum;
62	}
63
64	GroupMemoryBarrierWithGroupSync();
65	if( thread < 2 )
66	{
67		// const __m128 r2 = _mm_add_ps( r4, _mm_movehl_ps( r4, r4 ) );
68		sum += sharedAccumulators[ thread + 2 ];
69		sharedAccumulators[ thread ] = sum;
70	}
71
72	GroupMemoryBarrierWithGroupSync();
73	if( 0 == thread )
74	{
75		// const __m128 r1 = _mm_add_ss( r2, _mm_movehdup_ps( r2 ) );
76		sum += sharedAccumulators[ 1 ];
77	}
78}
79
80// Compute horisontal sum of the numbers, in yet another creative summation order recently implemented in the upstream
81void horizontalSumCompatNew( const uint thread, inout float sum )
82{
83	// GGML_F32x8_REDUCE
84	sharedAccumulators[ thread ] = sum;
85	GroupMemoryBarrierWithGroupSync();
86
87	if( 0 == ( thread & 8 ) )
88	{
89		// Runs on threads [ 0 .. 7 ] and [ 16 .. 23 ]
90		sum += sharedAccumulators[ thread | 8 ];
91		sharedAccumulators[ thread ] = sum;
92	}
93	GroupMemoryBarrierWithGroupSync();
94
95	if( thread < 8 )
96	{
97		// Runs on threads [ 0 .. 7 ]
98		sum += sharedAccumulators[ thread | 0x10 ];
99		sharedAccumulators[ thread ] = sum;
100	}
101	GroupMemoryBarrierWithGroupSync();
102
103	if( thread < 4 )
104	{
105		// Runs on threads [ 0 .. 3 ]
106		sum += sharedAccumulators[ thread | 4 ];
107		sharedAccumulators[ thread ] = sum;
108	}
109	GroupMemoryBarrierWithGroupSync();
110
111	if( thread < 4 && 0 == ( thread & 1 ) )
112	{
113		// Runs on threads [ 0, 2 ]
114		sum += sharedAccumulators[ thread | 1 ];
115		sharedAccumulators[ thread ] = sum;
116	}
117	GroupMemoryBarrierWithGroupSync();
118
119	if( 0 == thread )
120		sum += sharedAccumulators[ 2 ];
121}
122
123
124// Compute horizontal maximum of the numbers, and broadcast to all threads of the group.
125void horizontalMaxBroadcast( const uint thread, inout float ax )
126{
127	sharedAccumulators[ thread ] = ax;
128	for( uint i = 16; i > 0; i /= 2 )
129	{
130		GroupMemoryBarrierWithGroupSync();
131		if( thread < i )
132		{
133			ax = max( ax, sharedAccumulators[ thread + i ] );
134			sharedAccumulators[ thread ] = ax;
135		}
136	}
137	GroupMemoryBarrierWithGroupSync();
138	ax = sharedAccumulators[ 0 ];
139}