Kernels.hpp 13.6 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
26
#define PackedFloat_t   float4
#define WARP_SIZE       64
27
#define MAX_BLOCKSIZE   512
gilbertlee-amd's avatar
gilbertlee-amd committed
28
29
30
#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
31

gilbertlee-amd's avatar
gilbertlee-amd committed
32
33
34
35
// 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
36
{
37
  // Inputs
gilbertlee-amd's avatar
gilbertlee-amd committed
38
39
40
41
42
  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
43
44

  // Outputs
gilbertlee-amd's avatar
gilbertlee-amd committed
45
46
  long long startCycle;                         // Start timestamp for in-kernel timing (GPU-GFX executor)
  long long stopCycle;                          // Stop  timestamp for in-kernel timing (GPU-GFX executor)
47
  uint32_t  hwId;                               // Hardware ID
gilbertlee-amd's avatar
gilbertlee-amd committed
48
};
Gilbert Lee's avatar
Gilbert Lee committed
49

50
51
52
53
54
55
56
57
58
// 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
59
60
61
62
63
64
65
66
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)
67
      memset(p.dst[i], MEMSET_CHAR, p.N * sizeof(float));
gilbertlee-amd's avatar
gilbertlee-amd committed
68
69
70
71
72
73
  }
  else if (numSrcs == 1)
  {
    float const* __restrict__ src = p.src[0];
    for (int i = 0; i < numDsts; ++i)
    {
74
      memcpy(p.dst[i], src, p.N * sizeof(float));
gilbertlee-amd's avatar
gilbertlee-amd committed
75
76
77
78
79
80
81
82
83
84
85
86
87
    }
  }
  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;
    }
  }
}

88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
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);
}

// 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);
  }
}

gilbertlee-amd's avatar
gilbertlee-amd committed
110
111
112
113
114
115
116
// 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); }

// GPU copy kernel 0: 3 loops: unroll float 4, float4s, floats
template <int LOOP1_UNROLL>
117
__global__ void __launch_bounds__(MAX_BLOCKSIZE)
gilbertlee-amd's avatar
gilbertlee-amd committed
118
119
GpuReduceKernel(SubExecParam* params)
{
120
  int64_t startCycle = wall_clock64();
Gilbert Lee's avatar
Gilbert Lee committed
121
122

  // Operate on wavefront granularity
gilbertlee-amd's avatar
gilbertlee-amd committed
123
124
125
126
127
  SubExecParam& p    = params[blockIdx.x];
  int const numSrcs  = p.numSrcs;
  int const numDsts  = p.numDsts;
  int const waveId   = threadIdx.x / WARP_SIZE; // Wavefront number
  int const threadId = threadIdx.x % WARP_SIZE; // Thread index within wavefront
Gilbert Lee's avatar
Gilbert Lee committed
128
129
130

  // 1st loop - each wavefront operates on LOOP1_UNROLL x FLOATS_PER_PACK per thread per iteration
  // Determine the number of packed floats processed by the first loop
gilbertlee-amd's avatar
gilbertlee-amd committed
131
132
133
  size_t       Nrem        = p.N;
  size_t const loop1Npack  = (Nrem / (FLOATS_PER_PACK * LOOP1_UNROLL * WARP_SIZE)) * (LOOP1_UNROLL * WARP_SIZE);
  size_t const loop1Nelem  = loop1Npack * FLOATS_PER_PACK;
134
  size_t const loop1Inc    = blockDim.x * LOOP1_UNROLL;
gilbertlee-amd's avatar
gilbertlee-amd committed
135
  size_t       loop1Offset = waveId * LOOP1_UNROLL * WARP_SIZE + threadId;
Gilbert Lee's avatar
Gilbert Lee committed
136
137
138

  while (loop1Offset < loop1Npack)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
139
    PackedFloat_t vals[LOOP1_UNROLL] = {};
Gilbert Lee's avatar
Gilbert Lee committed
140

gilbertlee-amd's avatar
gilbertlee-amd committed
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
    if (numSrcs == 0)
    {
      #pragma unroll
      for (int u = 0; u < LOOP1_UNROLL; ++u) vals[u] = MemsetVal<float4>();
    }
    else
    {
      for (int i = 0; i < numSrcs; ++i)
      {
        PackedFloat_t const* __restrict__ packedSrc = (PackedFloat_t const*)(p.src[i]) + loop1Offset;
        #pragma unroll
        for (int u = 0; u < LOOP1_UNROLL; ++u)
          vals[u] += *(packedSrc + u * WARP_SIZE);
      }
    }
Gilbert Lee's avatar
Gilbert Lee committed
156

gilbertlee-amd's avatar
gilbertlee-amd committed
157
158
159
160
161
162
    for (int i = 0; i < numDsts; ++i)
    {
      PackedFloat_t* __restrict__ packedDst = (PackedFloat_t*)(p.dst[i]) + loop1Offset;
      #pragma unroll
      for (int u = 0; u < LOOP1_UNROLL; ++u) *(packedDst + u * WARP_SIZE) = vals[u];
    }
Gilbert Lee's avatar
Gilbert Lee committed
163
164
165
    loop1Offset += loop1Inc;
  }
  Nrem -= loop1Nelem;
gilbertlee-amd's avatar
gilbertlee-amd committed
166

Gilbert Lee's avatar
Gilbert Lee committed
167
168
169
  if (Nrem > 0)
  {
    // 2nd loop - Each thread operates on FLOATS_PER_PACK per iteration
gilbertlee-amd's avatar
gilbertlee-amd committed
170
171
172
    // NOTE: Using int32_t due to smaller size requirements
    int32_t const loop2Npack  = Nrem / FLOATS_PER_PACK;
    int32_t const loop2Nelem  = loop2Npack * FLOATS_PER_PACK;
173
    int32_t const loop2Inc    = blockDim.x;
gilbertlee-amd's avatar
gilbertlee-amd committed
174
    int32_t       loop2Offset = threadIdx.x;
Gilbert Lee's avatar
Gilbert Lee committed
175
176
177

    while (loop2Offset < loop2Npack)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
      PackedFloat_t val;
      if (numSrcs == 0)
      {
        val = MemsetVal<float4>();
      }
      else
      {
        val = {};
        for (int i = 0; i < numSrcs; ++i)
        {
          PackedFloat_t const* __restrict__ packedSrc = (PackedFloat_t const*)(p.src[i] + loop1Nelem) + loop2Offset;
          val += *packedSrc;
        }
      }

      for (int i = 0; i < numDsts; ++i)
      {
        PackedFloat_t* __restrict__ packedDst = (PackedFloat_t*)(p.dst[i] + loop1Nelem) + loop2Offset;
        *packedDst = val;
      }
Gilbert Lee's avatar
Gilbert Lee committed
198
199
200
201
202
203
204
205
      loop2Offset += loop2Inc;
    }
    Nrem -= loop2Nelem;

    // Deal with leftovers less than FLOATS_PER_PACK)
    if (threadIdx.x < Nrem)
    {
      int offset = loop1Nelem + loop2Nelem + threadIdx.x;
gilbertlee-amd's avatar
gilbertlee-amd committed
206
207
208
209
210
211
212
213
      float val = 0;
      if (numSrcs == 0)
      {
        val = MEMSET_VAL;
      }
      else
      {
        for (int i = 0; i < numSrcs; ++i)
214
          val += p.src[i][offset];
gilbertlee-amd's avatar
gilbertlee-amd committed
215
216
217
      }

      for (int i = 0; i < numDsts; ++i)
218
        p.dst[i][offset] = val;
Gilbert Lee's avatar
Gilbert Lee committed
219
220
221
    }
  }

gilbertlee-amd's avatar
gilbertlee-amd committed
222
  __syncthreads();
Gilbert Lee's avatar
Gilbert Lee committed
223
  if (threadIdx.x == 0)
gilbertlee-amd's avatar
gilbertlee-amd committed
224
  {
225
    p.stopCycle  = wall_clock64();
226
    p.startCycle = startCycle;
227
    __trace_hwreg();
gilbertlee-amd's avatar
gilbertlee-amd committed
228
  }
Gilbert Lee's avatar
Gilbert Lee committed
229
230
}

gilbertlee-amd's avatar
gilbertlee-amd committed
231
232
233
234
template <typename FLOAT_TYPE, int UNROLL_FACTOR>
__device__ size_t GpuReduceFuncImpl2(SubExecParam const &p, size_t const offset, size_t const N)
{
  int    constexpr numFloatsPerPack = sizeof(FLOAT_TYPE) / sizeof(float); // Number of floats handled at a time per thread
235
  size_t constexpr loopPackInc      = blockDim.x * UNROLL_FACTOR;
gilbertlee-amd's avatar
gilbertlee-amd committed
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
  size_t constexpr numPacksPerWave  = WARP_SIZE * UNROLL_FACTOR;
  int    const     waveId           = threadIdx.x / WARP_SIZE;            // Wavefront number
  int    const     threadId         = threadIdx.x % WARP_SIZE;            // Thread index within wavefront
  int    const     numSrcs          = p.numSrcs;
  int    const     numDsts          = p.numDsts;
  size_t const     numPacksDone     = (numFloatsPerPack == 1 && UNROLL_FACTOR == 1) ? N : (N / (FLOATS_PER_PACK * numPacksPerWave)) * numPacksPerWave;
  size_t const     numFloatsLeft    = N - numPacksDone * numFloatsPerPack;
  size_t           loopPackOffset   = waveId * numPacksPerWave + threadId;

  while (loopPackOffset < numPacksDone)
  {
    FLOAT_TYPE vals[UNROLL_FACTOR];

    if (numSrcs == 0)
    {
      #pragma unroll UNROLL_FACTOR
      for (int u = 0; u < UNROLL_FACTOR; ++u) vals[u] = MemsetVal<FLOAT_TYPE>();
    }
    else
    {
      FLOAT_TYPE const* __restrict__ src0Ptr = ((FLOAT_TYPE const*)(p.src[0] + offset)) + loopPackOffset;
      #pragma unroll UNROLL_FACTOR
      for (int u = 0; u < UNROLL_FACTOR; ++u)
        vals[u] = *(src0Ptr + u * WARP_SIZE);

      for (int i = 1; i < numSrcs; ++i)
      {
        FLOAT_TYPE const* __restrict__ srcPtr = ((FLOAT_TYPE const*)(p.src[i] + offset)) + loopPackOffset;

        #pragma unroll UNROLL_FACTOR
        for (int u = 0; u < UNROLL_FACTOR; ++u)
          vals[u] += *(srcPtr + u * WARP_SIZE);
      }
    }

    for (int i = 0; i < numDsts; ++i)
    {
      FLOAT_TYPE* __restrict__ dstPtr = (FLOAT_TYPE*)(p.dst[i + offset]) + loopPackOffset;
      #pragma unroll UNROLL_FACTOR
      for (int u = 0; u < UNROLL_FACTOR; ++u)
        *(dstPtr + u * WARP_SIZE) = vals[u];
    }
    loopPackOffset += loopPackInc;
  }

  return numFloatsLeft;
}

template <typename FLOAT_TYPE, int UNROLL_FACTOR>
__device__ size_t GpuReduceFuncImpl(SubExecParam const &p, size_t const offset, size_t const N)
Gilbert Lee's avatar
Gilbert Lee committed
286
{
gilbertlee-amd's avatar
gilbertlee-amd committed
287
288
  // Each thread in the block works on UNROLL_FACTOR FLOAT_TYPEs during each iteration of the loop
  int    constexpr numFloatsPerRead      = sizeof(FLOAT_TYPE) / sizeof(float);
289
290
  size_t const     numFloatsPerInnerLoop = blockDim.x * numFloatsPerRead;
  size_t const     numFloatsPerOuterLoop = numFloatsPerInnerLoop * UNROLL_FACTOR;
gilbertlee-amd's avatar
gilbertlee-amd committed
291
292
293
294
  size_t const     numFloatsLeft         = (numFloatsPerRead == 1 && UNROLL_FACTOR == 1) ? 0 : N % numFloatsPerOuterLoop;
  size_t const     numFloatsDone         = N - numFloatsLeft;
  int    const     numSrcs               = p.numSrcs;
  int    const     numDsts               = p.numDsts;
Gilbert Lee's avatar
Gilbert Lee committed
295

gilbertlee-amd's avatar
gilbertlee-amd committed
296
  for (size_t idx = threadIdx.x * numFloatsPerRead; idx < numFloatsDone; idx += numFloatsPerOuterLoop)
Gilbert Lee's avatar
Gilbert Lee committed
297
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
    FLOAT_TYPE tmp[UNROLL_FACTOR];

    if (numSrcs == 0)
    {
        #pragma unroll UNROLL_FACTOR
        for (int u = 0; u < UNROLL_FACTOR; ++u)
          tmp[u] = MemsetVal<FLOAT_TYPE>();
    }
    else
    {
      #pragma unroll UNROLL_FACTOR
      for (int u = 0; u < UNROLL_FACTOR; ++u)
        tmp[u] = *((FLOAT_TYPE*)(&p.src[0][offset + idx + u * numFloatsPerInnerLoop]));

      for (int i = 1; i < numSrcs; ++i)
      {
        #pragma unroll UNROLL_FACTOR
        for (int u = 0; u < UNROLL_FACTOR; ++u)
          tmp[u] += *((FLOAT_TYPE*)(&p.src[i][offset + idx + u * numFloatsPerInnerLoop]));
      }
    }

    for (int i = 0; i < numDsts; ++i)
    {
      for (int u = 0; u < UNROLL_FACTOR; ++u)
      {
        *((FLOAT_TYPE*)(&p.dst[i][offset + idx + u * numFloatsPerInnerLoop])) = tmp[u];
      }
    }
Gilbert Lee's avatar
Gilbert Lee committed
327
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
328
  return numFloatsLeft;
Gilbert Lee's avatar
Gilbert Lee committed
329
330
}

gilbertlee-amd's avatar
gilbertlee-amd committed
331
332
template <typename FLOAT_TYPE>
__device__ size_t GpuReduceFunc(SubExecParam const &p, size_t const offset, size_t const N, int const unroll)
Gilbert Lee's avatar
Gilbert Lee committed
333
{
gilbertlee-amd's avatar
gilbertlee-amd committed
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
  switch (unroll)
  {
  case  1: return GpuReduceFuncImpl<FLOAT_TYPE,  1>(p, offset, N);
  case  2: return GpuReduceFuncImpl<FLOAT_TYPE,  2>(p, offset, N);
  case  3: return GpuReduceFuncImpl<FLOAT_TYPE,  3>(p, offset, N);
  case  4: return GpuReduceFuncImpl<FLOAT_TYPE,  4>(p, offset, N);
  case  5: return GpuReduceFuncImpl<FLOAT_TYPE,  5>(p, offset, N);
  case  6: return GpuReduceFuncImpl<FLOAT_TYPE,  6>(p, offset, N);
  case  7: return GpuReduceFuncImpl<FLOAT_TYPE,  7>(p, offset, N);
  case  8: return GpuReduceFuncImpl<FLOAT_TYPE,  8>(p, offset, N);
  case  9: return GpuReduceFuncImpl<FLOAT_TYPE,  9>(p, offset, N);
  case 10: return GpuReduceFuncImpl<FLOAT_TYPE, 10>(p, offset, N);
  case 11: return GpuReduceFuncImpl<FLOAT_TYPE, 11>(p, offset, N);
  case 12: return GpuReduceFuncImpl<FLOAT_TYPE, 12>(p, offset, N);
  case 13: return GpuReduceFuncImpl<FLOAT_TYPE, 13>(p, offset, N);
  case 14: return GpuReduceFuncImpl<FLOAT_TYPE, 14>(p, offset, N);
  case 15: return GpuReduceFuncImpl<FLOAT_TYPE, 15>(p, offset, N);
  case 16: return GpuReduceFuncImpl<FLOAT_TYPE, 16>(p, offset, N);
  default: return GpuReduceFuncImpl<FLOAT_TYPE,  1>(p, offset, N);
  }
Gilbert Lee's avatar
Gilbert Lee committed
354
355
}

gilbertlee-amd's avatar
gilbertlee-amd committed
356
// GPU copy kernel
357
__global__ void __launch_bounds__(MAX_BLOCKSIZE)
gilbertlee-amd's avatar
gilbertlee-amd committed
358
GpuReduceKernel2(SubExecParam* params)
Gilbert Lee's avatar
Gilbert Lee committed
359
{
360
  int64_t startCycle = wall_clock64();
gilbertlee-amd's avatar
gilbertlee-amd committed
361
362
363
364
365
366
367
368
369
370
371
372
373
  SubExecParam& p = params[blockIdx.x];

  size_t numFloatsLeft = GpuReduceFunc<float4>(p, 0, p.N, 8);
  if (numFloatsLeft)
    numFloatsLeft = GpuReduceFunc<float4>(p, p.N - numFloatsLeft, numFloatsLeft, 1);

  if (numFloatsLeft)
  GpuReduceFunc<float>(p, p.N - numFloatsLeft, numFloatsLeft, 1);

  __threadfence_system();
  if (threadIdx.x == 0)
  {
    p.startCycle = startCycle;
374
    p.stopCycle  = wall_clock64();
gilbertlee-amd's avatar
gilbertlee-amd committed
375
  }
Gilbert Lee's avatar
Gilbert Lee committed
376
}
gilbertlee-amd's avatar
gilbertlee-amd committed
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423

#define NUM_GPU_KERNELS 18
typedef void (*GpuKernelFuncPtr)(SubExecParam*);

GpuKernelFuncPtr GpuKernelTable[NUM_GPU_KERNELS] =
{
  GpuReduceKernel<8>,
  GpuReduceKernel<1>,
  GpuReduceKernel<2>,
  GpuReduceKernel<3>,
  GpuReduceKernel<4>,
  GpuReduceKernel<5>,
  GpuReduceKernel<6>,
  GpuReduceKernel<7>,
  GpuReduceKernel<8>,
  GpuReduceKernel<9>,
  GpuReduceKernel<10>,
  GpuReduceKernel<11>,
  GpuReduceKernel<12>,
  GpuReduceKernel<13>,
  GpuReduceKernel<14>,
  GpuReduceKernel<15>,
  GpuReduceKernel<16>,
  GpuReduceKernel2
};

std::string GpuKernelNames[NUM_GPU_KERNELS] =
{
  "Default - 8xUnroll",
  "Unroll x1",
  "Unroll x2",
  "Unroll x3",
  "Unroll x4",
  "Unroll x5",
  "Unroll x6",
  "Unroll x7",
  "Unroll x8",
  "Unroll x9",
  "Unroll x10",
  "Unroll x11",
  "Unroll x12",
  "Unroll x13",
  "Unroll x14",
  "Unroll x15",
  "Unroll x16",
  "8xUnrollB",
};