Kernels.hpp 10.3 KB
Newer Older
Gilbert Lee's avatar
Gilbert Lee committed
1
/*
gilbertlee-amd's avatar
gilbertlee-amd committed
2
Copyright (c) 2022-2023 Advanced Micro Devices, Inc. All rights reserved.
Gilbert Lee's avatar
Gilbert Lee committed
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24

Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/

#pragma once

gilbertlee-amd's avatar
gilbertlee-amd committed
25
#define PackedFloat_t   float4
26
#define MAX_BLOCKSIZE   512
gilbertlee-amd's avatar
gilbertlee-amd committed
27
28
29
#define FLOATS_PER_PACK (sizeof(PackedFloat_t) / sizeof(float))
#define MEMSET_CHAR     75
#define MEMSET_VAL      13323083.0f
Gilbert Lee's avatar
Gilbert Lee committed
30

gilbertlee-amd's avatar
gilbertlee-amd committed
31

32
33
34
35
#if defined(__NVCC__)
#define warpSize 32
#endif

gilbertlee-amd's avatar
gilbertlee-amd committed
36
37
38
39
#define MAX_WAVEGROUPS  MAX_BLOCKSIZE / warpSize
#define MAX_UNROLL      8
#define NUM_WAVEORDERS  6

gilbertlee-amd's avatar
gilbertlee-amd committed
40
41
42
43
// Each subExecutor is provided with subarrays to work on
#define MAX_SRCS 16
#define MAX_DSTS 16
struct SubExecParam
Gilbert Lee's avatar
Gilbert Lee committed
44
{
45
  // Inputs
gilbertlee-amd's avatar
gilbertlee-amd committed
46
47
48
49
50
  size_t    N;                                  // Number of floats this subExecutor works on
  int       numSrcs;                            // Number of source arrays
  int       numDsts;                            // Number of destination arrays
  float*    src[MAX_SRCS];                      // Source array pointers
  float*    dst[MAX_DSTS];                      // Destination array pointers
51
  int32_t   preferredXccId;                     // XCC ID to execute on
52

gilbertlee-amd's avatar
gilbertlee-amd committed
53
54
55
56
  // Prepared
  int       teamSize;                           // Index of this sub executor amongst team
  int       teamIdx;                            // Size of team this sub executor is part of

57
  // Outputs
gilbertlee-amd's avatar
gilbertlee-amd committed
58
59
  long long startCycle;                         // Start timestamp for in-kernel timing (GPU-GFX executor)
  long long stopCycle;                          // Stop  timestamp for in-kernel timing (GPU-GFX executor)
60
  uint32_t  hwId;                               // Hardware ID
gilbertlee-amd's avatar
gilbertlee-amd committed
61
  uint32_t  xccId;                              // XCC ID
gilbertlee-amd's avatar
gilbertlee-amd committed
62
};
Gilbert Lee's avatar
Gilbert Lee committed
63

64
65
66
67
68
69
70
71
72
// Macro for collecting HW_REG_HW_ID
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__NVCC__)
#define __trace_hwreg() \
  p.hwId = 0
#else
#define __trace_hwreg() \
  asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_HW_ID)" : "=s" (p.hwId));
#endif

gilbertlee-amd's avatar
gilbertlee-amd committed
73
74
// Macro for collecting HW_REG_XCC_ID
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
75
76
#define GetXccId(val) \
  asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s" (val));
gilbertlee-amd's avatar
gilbertlee-amd committed
77
#else
78
79
#define GetXccId(val) \
  val = 0
gilbertlee-amd's avatar
gilbertlee-amd committed
80
81
#endif

gilbertlee-amd's avatar
gilbertlee-amd committed
82
83
84
85
86
87
88
89
void CpuReduceKernel(SubExecParam const& p)
{
  int const& numSrcs = p.numSrcs;
  int const& numDsts = p.numDsts;

  if (numSrcs == 0)
  {
    for (int i = 0; i < numDsts; ++i)
90
      memset(p.dst[i], MEMSET_CHAR, p.N * sizeof(float));
gilbertlee-amd's avatar
gilbertlee-amd committed
91
92
93
94
95
96
  }
  else if (numSrcs == 1)
  {
    float const* __restrict__ src = p.src[0];
    for (int i = 0; i < numDsts; ++i)
    {
97
      memcpy(p.dst[i], src, p.N * sizeof(float));
gilbertlee-amd's avatar
gilbertlee-amd committed
98
99
100
101
102
103
104
105
106
107
108
109
110
    }
  }
  else
  {
    for (int j = 0; j < p.N; j++)
    {
      float sum = p.src[0][j];
      for (int i = 1; i < numSrcs; i++) sum += p.src[i][j];
      for (int i = 0; i < numDsts; i++) p.dst[i][j] = sum;
    }
  }
}

111
112
113
114
115
116
117
118
119
120
std::string PrepSrcValueString()
{
  return "Element i = ((i * 517) modulo 383 + 31) * (srcBufferIdx + 1)";
}

__host__ __device__ float PrepSrcValue(int srcBufferIdx, size_t idx)
{
  return (((idx % 383) * 517) % 383 + 31) * (srcBufferIdx + 1);
}

121
122
123
124
125
126
127
__global__ void CollectXccIdsKernel(int* xccIds)
{
  int xccId;
  GetXccId(xccId);
  xccIds[blockIdx.x] = xccId;
}

128
129
130
131
132
133
134
135
136
137
138
139
// GPU kernel to prepare src buffer data
__global__ void
PrepSrcDataKernel(float* ptr, size_t N, int srcBufferIdx)
{
  for (size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
       idx < N;
       idx += blockDim.x * gridDim.x)
  {
    ptr[idx] = PrepSrcValue(srcBufferIdx, idx);
  }
}

140
141
142
143
144
145
146
147
148
149
150
__device__ int64_t GetTimestamp()
{
#if defined(__NVCC__)
  int64_t result;
  asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(result));
  return result;
#else
  return wall_clock64();
#endif
}

gilbertlee-amd's avatar
gilbertlee-amd committed
151
152
153
154
155
// Helper function for memset
template <typename T> __device__ __forceinline__ T      MemsetVal();
template <>           __device__ __forceinline__ float  MemsetVal(){ return MEMSET_VAL; };
template <>           __device__ __forceinline__ float4 MemsetVal(){ return make_float4(MEMSET_VAL, MEMSET_VAL, MEMSET_VAL, MEMSET_VAL); }

gilbertlee-amd's avatar
gilbertlee-amd committed
156
157
158
template <int BLOCKSIZE, int UNROLL>
__global__ void __launch_bounds__(BLOCKSIZE)
  GpuReduceKernel(SubExecParam* params, int waveOrder)
gilbertlee-amd's avatar
gilbertlee-amd committed
159
{
160
  int64_t startCycle;
161
  if (threadIdx.x == 0) startCycle = GetTimestamp();
162
163
164

  SubExecParam& p = params[blockIdx.y];

gilbertlee-amd's avatar
gilbertlee-amd committed
165
  // (Experimental) Filter by XCC if desired
166
#if !defined(__NVCC__)
gilbertlee-amd's avatar
gilbertlee-amd committed
167
  int32_t xccId;
168
169
  GetXccId(xccId);
  if (p.preferredXccId != -1 && xccId != p.preferredXccId) return;
170
#endif
Gilbert Lee's avatar
Gilbert Lee committed
171

gilbertlee-amd's avatar
gilbertlee-amd committed
172
173
174
175
176
177
178
179
  // Collect data information
  int32_t const  numSrcs  = p.numSrcs;
  int32_t const  numDsts  = p.numDsts;
  float4  const* __restrict__ srcFloat4[MAX_SRCS];
  float4*        __restrict__ dstFloat4[MAX_DSTS];
  for (int i = 0; i < numSrcs; i++) srcFloat4[i] = (float4*)p.src[i];
  for (int i = 0; i < numDsts; i++) dstFloat4[i] = (float4*)p.dst[i];

Gilbert Lee's avatar
Gilbert Lee committed
180
  // Operate on wavefront granularity
gilbertlee-amd's avatar
gilbertlee-amd committed
181
182
183
184
185
  int32_t const nTeams   = p.teamSize;             // Number of threadblocks working together on this subarray
  int32_t const teamIdx  = p.teamIdx;              // Index of this threadblock within the team
  int32_t const nWaves   = BLOCKSIZE   / warpSize; // Number of wavefronts within this threadblock
  int32_t const waveIdx  = threadIdx.x / warpSize; // Index of this wavefront within the threadblock
  int32_t const tIdx     = threadIdx.x % warpSize; // Thread index within wavefront
Gilbert Lee's avatar
Gilbert Lee committed
186

gilbertlee-amd's avatar
gilbertlee-amd committed
187
  size_t  const numFloat4 = p.N / 4;
Gilbert Lee's avatar
Gilbert Lee committed
188

gilbertlee-amd's avatar
gilbertlee-amd committed
189
190
191
192
193
194
195
196
197
  int32_t teamStride, waveStride, unrlStride, teamStride2, waveStride2;
  switch (waveOrder)
  {
  case 0: /* U,W,C */ unrlStride = 1; waveStride = UNROLL; teamStride = UNROLL * nWaves;  teamStride2 = nWaves; waveStride2 = 1     ; break;
  case 1: /* U,C,W */ unrlStride = 1; teamStride = UNROLL; waveStride = UNROLL * nTeams;  teamStride2 = 1;      waveStride2 = nTeams; break;
  case 2: /* W,U,C */ waveStride = 1; unrlStride = nWaves; teamStride = nWaves * UNROLL;  teamStride2 = nWaves; waveStride2 = 1     ; break;
  case 3: /* W,C,U */ waveStride = 1; teamStride = nWaves; unrlStride = nWaves * nTeams;  teamStride2 = nWaves; waveStride2 = 1     ; break;
  case 4: /* C,U,W */ teamStride = 1; unrlStride = nTeams; waveStride = nTeams * UNROLL;  teamStride2 = 1;      waveStride2 = nTeams; break;
  case 5: /* C,W,U */ teamStride = 1; waveStride = nTeams; unrlStride = nTeams * nWaves;  teamStride2 = 1;      waveStride2 = nTeams; break;
Gilbert Lee's avatar
Gilbert Lee committed
198
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
199

gilbertlee-amd's avatar
gilbertlee-amd committed
200
201
202
  // First loop: Each wavefront in the team works on UNROLL float4s per thread
  size_t const loop1Stride = nTeams * nWaves * UNROLL * warpSize;
  size_t const loop1Limit  = numFloat4 / loop1Stride * loop1Stride;
Gilbert Lee's avatar
Gilbert Lee committed
203
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
204
205
    float4 val[UNROLL];
    if (numSrcs == 0)
Gilbert Lee's avatar
Gilbert Lee committed
206
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
207
208
209
      #pragma unroll
      for (int u = 0; u < UNROLL; u++)
        val[u] = MemsetVal<float4>();
Gilbert Lee's avatar
Gilbert Lee committed
210
211
    }

gilbertlee-amd's avatar
gilbertlee-amd committed
212
    for (size_t idx = (teamIdx * teamStride + waveIdx * waveStride) * warpSize + tIdx; idx < loop1Limit; idx += loop1Stride)
Gilbert Lee's avatar
Gilbert Lee committed
213
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
214
215
      // Read sources into memory and accumulate in registers
      if (numSrcs)
gilbertlee-amd's avatar
gilbertlee-amd committed
216
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
217
218
219
220
221
        for (int u = 0; u < UNROLL; u++)
          val[u] = srcFloat4[0][idx + u * unrlStride * warpSize];
        for (int s = 1; s < numSrcs; s++)
          for (int u = 0; u < UNROLL; u++)
            val[u] += srcFloat4[s][idx + u * unrlStride * warpSize];
gilbertlee-amd's avatar
gilbertlee-amd committed
222
      }
gilbertlee-amd's avatar
gilbertlee-amd committed
223
224
225

      // Write accumulation to all outputs
      for (int d = 0; d < numDsts; d++)
gilbertlee-amd's avatar
gilbertlee-amd committed
226
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
227
228
229
        #pragma unroll
        for (int u = 0; u < UNROLL; u++)
          dstFloat4[d][idx + u * unrlStride * warpSize] = val[u];
gilbertlee-amd's avatar
gilbertlee-amd committed
230
      }
Gilbert Lee's avatar
Gilbert Lee committed
231
232
233
    }
  }

gilbertlee-amd's avatar
gilbertlee-amd committed
234
  // Second loop: Deal with remaining float4s
gilbertlee-amd's avatar
gilbertlee-amd committed
235
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
236
    if (loop1Limit < numFloat4)
gilbertlee-amd's avatar
gilbertlee-amd committed
237
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
238
239
      float4 val;
      if (numSrcs == 0) val = MemsetVal<float4>();
gilbertlee-amd's avatar
gilbertlee-amd committed
240

gilbertlee-amd's avatar
gilbertlee-amd committed
241
242
      size_t const loop2Stride = nTeams * nWaves * warpSize;
      for (size_t idx = loop1Limit + (teamIdx * teamStride2 + waveIdx * waveStride2) * warpSize + tIdx; idx < numFloat4; idx += loop2Stride)
gilbertlee-amd's avatar
gilbertlee-amd committed
243
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
244
245
246
247
248
249
        if (numSrcs)
        {
          val = srcFloat4[0][idx];
          for (int s = 1; s < numSrcs; s++)
            val += srcFloat4[s][idx];
        }
gilbertlee-amd's avatar
gilbertlee-amd committed
250

gilbertlee-amd's avatar
gilbertlee-amd committed
251
252
        for (int d = 0; d < numDsts; d++)
          dstFloat4[d][idx] = val;
gilbertlee-amd's avatar
gilbertlee-amd committed
253
254
255
256
      }
    }
  }

gilbertlee-amd's avatar
gilbertlee-amd committed
257
  // Third loop; Deal with remaining floats
Gilbert Lee's avatar
Gilbert Lee committed
258
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
259
    if (numFloat4 * 4 < p.N)
gilbertlee-amd's avatar
gilbertlee-amd committed
260
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
261
262
      float val;
      if (numSrcs == 0) val = MemsetVal<float>();
gilbertlee-amd's avatar
gilbertlee-amd committed
263

gilbertlee-amd's avatar
gilbertlee-amd committed
264
265
      size_t const loop3Stride = nTeams * nWaves * warpSize;
      for( size_t idx = numFloat4 * 4 + (teamIdx * teamStride2 + waveIdx * waveStride2) * warpSize + tIdx; idx < p.N; idx += loop3Stride)
gilbertlee-amd's avatar
gilbertlee-amd committed
266
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
267
268
269
270
271
272
        if (numSrcs)
        {
          val = p.src[0][idx];
          for (int s = 1; s < numSrcs; s++)
            val += p.src[s][idx];
        }
gilbertlee-amd's avatar
gilbertlee-amd committed
273

gilbertlee-amd's avatar
gilbertlee-amd committed
274
275
        for (int d = 0; d < numDsts; d++)
          p.dst[d][idx] = val;
gilbertlee-amd's avatar
gilbertlee-amd committed
276
277
      }
    }
Gilbert Lee's avatar
Gilbert Lee committed
278
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
279

gilbertlee-amd's avatar
gilbertlee-amd committed
280
281
  // Wait for all threads to finish
  __syncthreads();
gilbertlee-amd's avatar
gilbertlee-amd committed
282
283
  if (threadIdx.x == 0)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
284
    __threadfence_system();
285
    p.stopCycle  = GetTimestamp();
gilbertlee-amd's avatar
gilbertlee-amd committed
286
287
288
    p.startCycle = startCycle;
    p.xccId      = xccId;
    __trace_hwreg();
gilbertlee-amd's avatar
gilbertlee-amd committed
289
  }
Gilbert Lee's avatar
Gilbert Lee committed
290
}
gilbertlee-amd's avatar
gilbertlee-amd committed
291

gilbertlee-amd's avatar
gilbertlee-amd committed
292
typedef void (*GpuKernelFuncPtr)(SubExecParam*, int);
gilbertlee-amd's avatar
gilbertlee-amd committed
293

gilbertlee-amd's avatar
gilbertlee-amd committed
294
295
296
297
298
299
300
301
302
#define GPU_KERNEL_UNROLL_DECL(BLOCKSIZE) \
  {GpuReduceKernel<BLOCKSIZE, 1>,  \
   GpuReduceKernel<BLOCKSIZE, 2>,  \
   GpuReduceKernel<BLOCKSIZE, 3>,  \
   GpuReduceKernel<BLOCKSIZE, 4>,  \
   GpuReduceKernel<BLOCKSIZE, 5>,  \
   GpuReduceKernel<BLOCKSIZE, 6>,  \
   GpuReduceKernel<BLOCKSIZE, 7>,  \
   GpuReduceKernel<BLOCKSIZE, 8>}
gilbertlee-amd's avatar
gilbertlee-amd committed
303

gilbertlee-amd's avatar
gilbertlee-amd committed
304
GpuKernelFuncPtr GpuKernelTable[MAX_WAVEGROUPS][MAX_UNROLL] =
gilbertlee-amd's avatar
gilbertlee-amd committed
305
{
gilbertlee-amd's avatar
gilbertlee-amd committed
306
307
308
309
310
311
312
313
  GPU_KERNEL_UNROLL_DECL(64),
  GPU_KERNEL_UNROLL_DECL(128),
  GPU_KERNEL_UNROLL_DECL(192),
  GPU_KERNEL_UNROLL_DECL(256),
  GPU_KERNEL_UNROLL_DECL(320),
  GPU_KERNEL_UNROLL_DECL(384),
  GPU_KERNEL_UNROLL_DECL(448),
  GPU_KERNEL_UNROLL_DECL(512)
gilbertlee-amd's avatar
gilbertlee-amd committed
314
};