include/ck/utility/amd_buffer_addressing.hpp Source File

include/ck/utility/amd_buffer_addressing.hpp Source File#

Composable Kernel: include/ck/utility/amd_buffer_addressing.hpp Source File
amd_buffer_addressing.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 #include "data_type.hpp"
6 
7 namespace ck {
8 
9 template <typename T>
10 union BufferResource
11 {
12  __device__ constexpr BufferResource() : content{} {}
13 
14  // 128 bit SGPRs to supply buffer resource in buffer instructions
15  // https://rocm-documentation.readthedocs.io/en/latest/GCN_ISA_Manuals/testdocbook.html#vector-memory-buffer-instructions
16  int32x4_t content;
17  StaticallyIndexedArray<T*, 2> address;
18  StaticallyIndexedArray<int32_t, 4> range;
19  StaticallyIndexedArray<int32_t, 4> config;
20 };
21 
22 template <typename T>
23 __device__ int32x4_t make_wave_buffer_resource(T* p_wave, index_t element_space_size)
24 {
25  BufferResource<T> wave_buffer_resource;
26 
27  // wavewise base address (64 bit)
28  wave_buffer_resource.address(Number<0>{}) = const_cast<remove_cv_t<T>*>(p_wave);
29  // wavewise range (32 bit)
30  wave_buffer_resource.range(Number<2>{}) = element_space_size * sizeof(T);
31  // wavewise setting (32 bit)
32  wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD;
33 
34  return wave_buffer_resource.content;
35 }
36 
37 template <typename T>
38 __device__ int32x4_t make_wave_buffer_resource_with_default_range(T* p_wave)
39 {
40  BufferResource<T> wave_buffer_resource;
41 
42  // wavewise base address (64 bit)
43  wave_buffer_resource.address(Number<0>{}) = const_cast<remove_cv_t<T>*>(p_wave);
44  // wavewise range (32 bit)
45  wave_buffer_resource.range(Number<2>{}) = 0xffffffff; // max possible range
46  // wavewise setting (32 bit)
47  wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD;
48 
49  return wave_buffer_resource.content;
50 }
51 
52 // buffer load i8
53 __device__ int8_t
54 llvm_amdgcn_raw_buffer_load_i8(int32x4_t srsrc,
55  index_t voffset,
56  index_t soffset,
57  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i8");
58 
59 __device__ int8x2_t
60 llvm_amdgcn_raw_buffer_load_i8x2(int32x4_t srsrc,
61  index_t voffset,
62  index_t soffset,
63  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i8");
64 
65 __device__ int8x4_t
66 llvm_amdgcn_raw_buffer_load_i8x4(int32x4_t srsrc,
67  index_t voffset,
68  index_t soffset,
69  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i8");
70 
71 // buffer load i16
72 __device__ bhalf_t
73 llvm_amdgcn_raw_buffer_load_i16(int32x4_t srsrc,
74  index_t voffset,
75  index_t soffset,
76  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i16");
77 
78 __device__ bhalf2_t
79 llvm_amdgcn_raw_buffer_load_i16x2(int32x4_t srsrc,
80  index_t voffset,
81  index_t soffset,
82  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i16");
83 
84 __device__ bhalf4_t
85 llvm_amdgcn_raw_buffer_load_i16x4(int32x4_t srsrc,
86  index_t voffset,
87  index_t soffset,
88  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i16");
89 
90 // buffer load i32
91 __device__ int32_t
92 llvm_amdgcn_raw_buffer_load_i32(int32x4_t srsrc,
93  index_t voffset,
94  index_t soffset,
95  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i32");
96 
97 __device__ int32x2_t
98 llvm_amdgcn_raw_buffer_load_i32x2(int32x4_t srsrc,
99  index_t voffset,
100  index_t soffset,
101  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i32");
102 
103 __device__ int32x4_t
104 llvm_amdgcn_raw_buffer_load_i32x4(int32x4_t srsrc,
105  index_t voffset,
106  index_t soffset,
107  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i32");
108 
109 // buffer load fp16
110 __device__ half_t
111 llvm_amdgcn_raw_buffer_load_fp16(int32x4_t srsrc,
112  index_t voffset,
113  index_t soffset,
114  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.f16");
115 
116 __device__ half2_t
117 llvm_amdgcn_raw_buffer_load_fp16x2(int32x4_t srsrc,
118  index_t voffset,
119  index_t soffset,
120  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2f16");
121 
122 __device__ half4_t
123 llvm_amdgcn_raw_buffer_load_fp16x4(int32x4_t srsrc,
124  index_t voffset,
125  index_t soffset,
126  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4f16");
127 
128 // buffer load fp32
129 __device__ float
130 llvm_amdgcn_raw_buffer_load_fp32(int32x4_t srsrc,
131  index_t voffset,
132  index_t soffset,
133  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.f32");
134 
135 __device__ float2_t
136 llvm_amdgcn_raw_buffer_load_fp32x2(int32x4_t srsrc,
137  index_t voffset,
138  index_t soffset,
139  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2f32");
140 
141 __device__ float4_t
142 llvm_amdgcn_raw_buffer_load_fp32x4(int32x4_t srsrc,
143  index_t voffset,
144  index_t soffset,
145  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4f32");
146 
147 // buffer store i8
148 __device__ void
149 llvm_amdgcn_raw_buffer_store_i8(int8_t vdata,
150  int32x4_t rsrc,
151  index_t voffset,
152  index_t soffset,
153  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i8");
154 
155 __device__ void
156 llvm_amdgcn_raw_buffer_store_i8x2(int8x2_t vdata,
157  int32x4_t rsrc,
158  index_t voffset,
159  index_t soffset,
160  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i8");
161 
162 __device__ void
163 llvm_amdgcn_raw_buffer_store_i8x4(int8x4_t vdata,
164  int32x4_t rsrc,
165  index_t voffset,
166  index_t soffset,
167  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i8");
168 
169 // buffer store i16
170 __device__ void
171 llvm_amdgcn_raw_buffer_store_i16(bhalf_t vdata,
172  int32x4_t rsrc,
173  index_t voffset,
174  index_t soffset,
175  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i16");
176 
177 __device__ void
178 llvm_amdgcn_raw_buffer_store_i16x2(bhalf2_t vdata,
179  int32x4_t rsrc,
180  index_t voffset,
181  index_t soffset,
182  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i16");
183 
184 __device__ void
185 llvm_amdgcn_raw_buffer_store_i16x4(bhalf4_t vdata,
186  int32x4_t rsrc,
187  index_t voffset,
188  index_t soffset,
189  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i16");
190 
191 // buffer store i32
192 __device__ void
193 llvm_amdgcn_raw_buffer_store_i32(int32_t vdata,
194  int32x4_t rsrc,
195  index_t voffset,
196  index_t soffset,
197  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i32");
198 
199 __device__ void
200 llvm_amdgcn_raw_buffer_store_i32x2(int32x2_t vdata,
201  int32x4_t rsrc,
202  index_t voffset,
203  index_t soffset,
204  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i32");
205 
206 __device__ void
207 llvm_amdgcn_raw_buffer_store_i32x4(int32x4_t vdata,
208  int32x4_t rsrc,
209  index_t voffset,
210  index_t soffset,
211  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i32");
212 
213 // buffer store fp16
214 __device__ void
215 llvm_amdgcn_raw_buffer_store_fp16(half_t vdata,
216  int32x4_t rsrc,
217  index_t voffset,
218  index_t soffset,
219  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.f16");
220 
221 __device__ void
222 llvm_amdgcn_raw_buffer_store_fp16x2(half2_t vdata,
223  int32x4_t rsrc,
224  index_t voffset,
225  index_t soffset,
226  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2f16");
227 
228 __device__ void
229 llvm_amdgcn_raw_buffer_store_fp16x4(half4_t vdata,
230  int32x4_t rsrc,
231  index_t voffset,
232  index_t soffset,
233  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4f16");
234 
235 // buffer store fp32
236 __device__ void
237 llvm_amdgcn_raw_buffer_store_fp32(float vdata,
238  int32x4_t rsrc,
239  index_t voffset,
240  index_t soffset,
241  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.f32");
242 
243 __device__ void
244 llvm_amdgcn_raw_buffer_store_fp32x2(float2_t vdata,
245  int32x4_t rsrc,
246  index_t voffset,
247  index_t soffset,
248  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2f32");
249 
250 __device__ void
251 llvm_amdgcn_raw_buffer_store_fp32x4(float4_t vdata,
252  int32x4_t rsrc,
253  index_t voffset,
254  index_t soffset,
255  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4f32");
256 
257 // buffer atomic-add fp16
258 __device__ half2_t llvm_amdgcn_raw_buffer_atomic_add_fp16x2(
259  half2_t vdata,
260  int32x4_t rsrc,
261  index_t voffset,
262  index_t soffset,
263  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.v2f16");
264 
265 // buffer atomic-add i32
266 __device__ int32_t llvm_amdgcn_raw_buffer_atomic_add_i32(
267  int32_t vdata,
268  int32x4_t rsrc,
269  index_t voffset,
270  index_t soffset,
271  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.add.i32");
272 
273 // buffer atomic-add fp32
274 __device__ float llvm_amdgcn_raw_buffer_atomic_add_fp32(
275  float vdata,
276  int32x4_t rsrc,
277  index_t voffset,
278  index_t soffset,
279  index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.f32");
280 
281 // buffer atomic-add fp32
282 __device__ double
283 llvm_amdgcn_raw_buffer_atomic_max_fp64(double vdata,
284  int32x4_t rsrc, // dst_wave_buffer_resource
285  int voffset, // dst_thread_addr_offset
286  int soffset, // dst_wave_addr_offset
287  int glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fmax.f64");
288 
289 // memory coherency bit for buffer store/load instruction
290 // check ISA manual for each GFX target
291 // e.g. for
292 // https://www.amd.com/system/files/TechDocs/instinct-mi200-cdna2-instruction-set-architecture.pdf,
293 // page 67~68
294 enum struct AmdBufferCoherenceEnum
295 {
296  DefaultCoherence = 0, // default value
297  GLC = 1,
298  SLC = 2,
299  GLC_SLC = 3,
300  // gfx94: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
301  // SC[1:0] System Cache level: 0=wave, 1=group, 2=device, 3=system
302  // NT Non-Temporal: 0=expect temporal reuse; 1=do not expect temporal reuse
303  WAVE_NT0 = 0,
304  WAVE_NT1 = 2,
305  GROUP_NT0 = 1,
306  GROUP_NT1 = 3,
307  DEVICE_NT0 = 8,
308  DEVICE_NT1 = 10,
309  SYSTEM_NT0 = 9,
310  SYSTEM_NT1 = 11,
311 };
312 
313 template <index_t N, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
314 __device__ typename vector_type<int8_t, N>::type
315 amd_buffer_load_impl_raw(int32x4_t src_wave_buffer_resource,
316  index_t src_thread_addr_offset,
317  index_t src_wave_addr_offset)
318 {
319  static_assert(N == 1 || N == 2 || N == 4 || N == 8 || N == 16 || N == 32 || N == 64,
320  "wrong! not implemented");
321 
322  if constexpr(N == 1)
323  {
324  return llvm_amdgcn_raw_buffer_load_i8(src_wave_buffer_resource,
325  src_thread_addr_offset,
326  src_wave_addr_offset,
327  static_cast<index_t>(coherence));
328  }
329  else if constexpr(N == 2)
330  {
331 
332  int16_t tmp = llvm_amdgcn_raw_buffer_load_i16(src_wave_buffer_resource,
333  src_thread_addr_offset,
334  src_wave_addr_offset,
335  static_cast<index_t>(coherence));
336 
337  return bit_cast<int8x2_t>(tmp);
338  }
339  else if constexpr(N == 4)
340  {
341  int32_t tmp = llvm_amdgcn_raw_buffer_load_i32(src_wave_buffer_resource,
342  src_thread_addr_offset,
343  src_wave_addr_offset,
344  static_cast<index_t>(coherence));
345 
346  return bit_cast<int8x4_t>(tmp);
347  }
348  else if constexpr(N == 8)
349  {
350  int32x2_t tmp = llvm_amdgcn_raw_buffer_load_i32x2(src_wave_buffer_resource,
351  src_thread_addr_offset,
352  src_wave_addr_offset,
353  static_cast<index_t>(coherence));
354 
355  return bit_cast<int8x8_t>(tmp);
356  }
357  else if constexpr(N == 16)
358  {
359  int32x4_t tmp = llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource,
360  src_thread_addr_offset,
361  src_wave_addr_offset,
362  static_cast<index_t>(coherence));
363  return bit_cast<int8x16_t>(tmp);
364  }
365  else if constexpr(N == 32)
366  {
367  int32x4_t tmp0 = llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource,
368  src_thread_addr_offset,
369  src_wave_addr_offset,
370  static_cast<index_t>(coherence));
371  int32x4_t tmp1 =
372  llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource,
373  src_thread_addr_offset,
374  src_wave_addr_offset + 4 * sizeof(int32_t),
375  static_cast<index_t>(coherence));
376  vector_type<int32_t, 8> tmp;
377 
378  tmp.AsType<int32x4_t>()(Number<0>{}) = tmp0;
379  tmp.AsType<int32x4_t>()(Number<1>{}) = tmp1;
380 
381  return bit_cast<int8x32_t>(tmp);
382  }
383  else if constexpr(N == 64)
384  {
385  int32x4_t tmp0 = llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource,
386  src_thread_addr_offset,
387  src_wave_addr_offset,
388  static_cast<index_t>(coherence));
389  int32x4_t tmp1 =
390  llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource,
391  src_thread_addr_offset,
392  src_wave_addr_offset + 4 * sizeof(int32_t),
393  static_cast<index_t>(coherence));
394  int32x4_t tmp2 =
395  llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource,
396  src_thread_addr_offset,
397  src_wave_addr_offset + 8 * sizeof(int32_t),
398  static_cast<index_t>(coherence));
399  int32x4_t tmp3 =
400  llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource,
401  src_thread_addr_offset,
402  src_wave_addr_offset + 12 * sizeof(int32_t),
403  static_cast<index_t>(coherence));
404 
405  vector_type<int32_t, 16> tmp;
406 
407  tmp.AsType<int32x4_t>()(Number<0>{}) = tmp0;
408  tmp.AsType<int32x4_t>()(Number<1>{}) = tmp1;
409  tmp.AsType<int32x4_t>()(Number<2>{}) = tmp2;
410  tmp.AsType<int32x4_t>()(Number<3>{}) = tmp3;
411 
412  return bit_cast<int8x64_t>(tmp);
413  }
414 }
415 
416 template <typename T,
417  index_t N,
418  AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
419 __device__ typename vector_type<T, N>::type amd_buffer_load_impl(int32x4_t src_wave_buffer_resource,
420  index_t src_thread_addr_offset,
421  index_t src_wave_addr_offset)
422 {
423  static_assert(
424  (is_same<T, double>::value && (N == 1 || N == 2 || N == 4 || N == 8)) ||
425  (is_same<T, float>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
426  (is_same<T, half_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
427  (is_same<T, bhalf_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
428  (is_same<T, int32_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
429  (is_same<T, f8_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
430  (is_same<T, bf8_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
431  (is_same<T, int8_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
432  (is_same<T, uint8_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
433  (is_same<T, pk_i4_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)),
434  "wrong! not implemented");
435 
436  using r_t = typename vector_type<T, N>::type;
437  auto raw_data = amd_buffer_load_impl_raw<sizeof(T) * N, coherence>(
438  src_wave_buffer_resource, src_thread_addr_offset, src_wave_addr_offset);
439  return bit_cast<r_t>(raw_data);
440 }
441 
442 template <index_t N, AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
443 __device__ void
444 amd_buffer_store_impl_raw(const typename vector_type<int8_t, N>::type src_thread_data,
445  int32x4_t dst_wave_buffer_resource,
446  index_t dst_thread_addr_offset,
447  index_t dst_wave_addr_offset)
448 {
449  static_assert(N == 1 || N == 2 || N == 4 || N == 8 || N == 16 || N == 32 || N == 64,
450  "wrong! not implemented");
451 
452  if constexpr(N == 1)
453  {
454  llvm_amdgcn_raw_buffer_store_i8(src_thread_data,
455  dst_wave_buffer_resource,
456  dst_thread_addr_offset,
457  dst_wave_addr_offset,
458  static_cast<index_t>(coherence));
459  }
460  else if constexpr(N == 2)
461  {
462 
463  llvm_amdgcn_raw_buffer_store_i16(bit_cast<int16_t>(src_thread_data),
464  dst_wave_buffer_resource,
465  dst_thread_addr_offset,
466  dst_wave_addr_offset,
467  static_cast<index_t>(coherence));
468  }
469  else if constexpr(N == 4)
470  {
471  llvm_amdgcn_raw_buffer_store_i32(bit_cast<int32_t>(src_thread_data),
472  dst_wave_buffer_resource,
473  dst_thread_addr_offset,
474  dst_wave_addr_offset,
475  static_cast<index_t>(coherence));
476  }
477  else if constexpr(N == 8)
478  {
479  llvm_amdgcn_raw_buffer_store_i32x2(bit_cast<int32x2_t>(src_thread_data),
480  dst_wave_buffer_resource,
481  dst_thread_addr_offset,
482  dst_wave_addr_offset,
483  static_cast<index_t>(coherence));
484  }
485  else if constexpr(N == 16)
486  {
487  llvm_amdgcn_raw_buffer_store_i32x4(bit_cast<int32x4_t>(src_thread_data),
488  dst_wave_buffer_resource,
489  dst_thread_addr_offset,
490  dst_wave_addr_offset,
491  static_cast<index_t>(coherence));
492  }
493  else if constexpr(N == 32)
494  {
495  vector_type<int32_t, 8> tmp{bit_cast<int32x8_t>(src_thread_data)};
496 
497  llvm_amdgcn_raw_buffer_store_i32x4(tmp.template AsType<int32x4_t>()[Number<0>{}],
498  dst_wave_buffer_resource,
499  dst_thread_addr_offset,
500  dst_wave_addr_offset,
501  static_cast<index_t>(coherence));
502 
503  llvm_amdgcn_raw_buffer_store_i32x4(tmp.template AsType<int32x4_t>()[Number<1>{}],
504  dst_wave_buffer_resource,
505  dst_thread_addr_offset,
506  dst_wave_addr_offset + sizeof(int32_t) * 4,
507  static_cast<index_t>(coherence));
508  }
509  else if constexpr(N == 64)
510  {
511  vector_type<int32_t, 16> tmp{bit_cast<int32x16_t>(src_thread_data)};
512 
513  llvm_amdgcn_raw_buffer_store_i32x4(tmp.template AsType<int32x4_t>()[Number<0>{}],
514  dst_wave_buffer_resource,
515  dst_thread_addr_offset,
516  dst_wave_addr_offset,
517  static_cast<index_t>(coherence));
518 
519  llvm_amdgcn_raw_buffer_store_i32x4(tmp.template AsType<int32x4_t>()[Number<1>{}],
520  dst_wave_buffer_resource,
521  dst_thread_addr_offset,
522  dst_wave_addr_offset + sizeof(int32_t) * 4,
523  static_cast<index_t>(coherence));
524 
525  llvm_amdgcn_raw_buffer_store_i32x4(tmp.template AsType<int32x4_t>()[Number<2>{}],
526  dst_wave_buffer_resource,
527  dst_thread_addr_offset,
528  dst_wave_addr_offset + sizeof(int32_t) * 8,
529  static_cast<index_t>(coherence));
530 
531  llvm_amdgcn_raw_buffer_store_i32x4(tmp.template AsType<int32x4_t>()[Number<3>{}],
532  dst_wave_buffer_resource,
533  dst_thread_addr_offset,
534  dst_wave_addr_offset + sizeof(int32_t) * 12,
535  static_cast<index_t>(coherence));
536  }
537 }
538 
539 template <typename T,
540  index_t N,
541  AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
542 __device__ void amd_buffer_store_impl(const typename vector_type<T, N>::type src_thread_data,
543  int32x4_t dst_wave_buffer_resource,
544  index_t dst_thread_addr_offset,
545  index_t dst_wave_addr_offset)
546 {
547  static_assert(
548  (is_same<T, double>::value && (N == 1 || N == 2 || N == 4 || N == 8)) ||
549  (is_same<T, float>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
550  (is_same<T, half_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
551  (is_same<T, bhalf_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
552  (is_same<T, int32_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
553  (is_same<T, f8_fnuz_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
554  (is_same<T, bf8_fnuz_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
555  (is_same<T, fp8_storage_t>::value &&
556  (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) ||
557  (is_same<T, int8_t>::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)),
558  "wrong! not implemented");
559 
560  using r_t = typename vector_type<int8_t, sizeof(T) * N>::type;
561 
562  amd_buffer_store_impl_raw<sizeof(T) * N, coherence>(bit_cast<r_t>(src_thread_data),
563  dst_wave_buffer_resource,
564  dst_thread_addr_offset,
565  dst_wave_addr_offset);
566 }
567 
568 template <typename T, index_t N>
569 __device__ void amd_global_atomic_add_impl(const typename vector_type<T, N>::type src_thread_data,
570  T* addr)
571 {
572  static_assert((is_same<T, bhalf_t>::value && (N == 2 || N == 4 || N == 8)) ||
573  (is_same<T, half_t>::value && (N == 2 || N == 4 || N == 8)),
574  "wrong! not implemented");
575 
576  if constexpr(is_same<T, half_t>::value)
577  {
578  vector_type<half_t, N> tmp{src_thread_data};
579  static_for<0, N / 2, 1>{}([&](auto i) {
580  __builtin_amdgcn_global_atomic_fadd_v2f16(bit_cast<half2_t*>(addr) + i,
581  tmp.template AsType<half2_t>()[i]);
582  });
583  }
584 #if defined(__gfx942__) || defined(__gfx950__)
585  else if constexpr(is_same<T, bhalf_t>::value)
586  {
587  vector_type<bhalf_t, N> tmp{src_thread_data};
588  static_for<0, N / 2, 1>{}([&](auto i) {
589  __builtin_amdgcn_global_atomic_fadd_v2bf16(bit_cast<bhalf2_t*>(addr) + i,
590  tmp.template AsType<bhalf2_t>()[i]);
591  });
592  }
593 #endif
594 }
595 
596 template <typename T, index_t N>
597 __device__ void amd_buffer_atomic_add_impl(const typename vector_type<T, N>::type src_thread_data,
598  int32x4_t dst_wave_buffer_resource,
599  index_t dst_thread_addr_offset,
600  index_t dst_wave_addr_offset)
601 {
602  static_assert((is_same<T, float>::value && (N == 1 || N == 2 || N == 4)) ||
603  (is_same<T, half_t>::value && (N == 2 || N == 4 || N == 8)) ||
604  (is_same<T, int32_t>::value && (N == 1 || N == 2 || N == 4)),
605  "wrong! not implemented");
606 
607  if constexpr(is_same<T, float>::value)
608  {
609  if constexpr(N == 1)
610  {
611  llvm_amdgcn_raw_buffer_atomic_add_fp32(src_thread_data,
612  dst_wave_buffer_resource,
613  dst_thread_addr_offset,
614  dst_wave_addr_offset,
615  0);
616  }
617  else if constexpr(N == 2)
618  {
619  vector_type<float, 2> tmp{src_thread_data};
620 
621  llvm_amdgcn_raw_buffer_atomic_add_fp32(tmp.AsType<float>()[Number<0>{}],
622  dst_wave_buffer_resource,
623  dst_thread_addr_offset,
624  dst_wave_addr_offset,
625  0);
626 
627  llvm_amdgcn_raw_buffer_atomic_add_fp32(tmp.AsType<float>()[Number<1>{}],
628  dst_wave_buffer_resource,
629  dst_thread_addr_offset,
630  dst_wave_addr_offset + sizeof(float),
631  0);
632  }
633  else if constexpr(N == 4)
634  {
635  vector_type<float, 4> tmp{src_thread_data};
636 
637  llvm_amdgcn_raw_buffer_atomic_add_fp32(tmp.AsType<float>()[Number<0>{}],
638  dst_wave_buffer_resource,
639  dst_thread_addr_offset,
640  dst_wave_addr_offset,
641  0);
642 
643  llvm_amdgcn_raw_buffer_atomic_add_fp32(tmp.AsType<float>()[Number<1>{}],
644  dst_wave_buffer_resource,
645  dst_thread_addr_offset,
646  dst_wave_addr_offset + sizeof(float),
647  0);
648 
649  llvm_amdgcn_raw_buffer_atomic_add_fp32(tmp.AsType<float>()[Number<2>{}],
650  dst_wave_buffer_resource,
651  dst_thread_addr_offset,
652  dst_wave_addr_offset + 2 * sizeof(float),
653  0);
654 
655  llvm_amdgcn_raw_buffer_atomic_add_fp32(tmp.AsType<float>()[Number<3>{}],
656  dst_wave_buffer_resource,
657  dst_thread_addr_offset,
658  dst_wave_addr_offset + 3 * sizeof(float),
659  0);
660  }
661  }
662  else if constexpr(is_same<T, half_t>::value)
663  {
664  if constexpr(N == 2)
665  {
666  llvm_amdgcn_raw_buffer_atomic_add_fp16x2(src_thread_data,
667  dst_wave_buffer_resource,
668  dst_thread_addr_offset,
669  dst_wave_addr_offset,
670  0);
671  }
672  else if constexpr(N == 4)
673  {
674  vector_type<half_t, 4> tmp{src_thread_data};
675 
676  static_for<0, 2, 1>{}([&](auto i) {
677  llvm_amdgcn_raw_buffer_atomic_add_fp16x2(tmp.AsType<half2_t>()[i],
678  dst_wave_buffer_resource,
679  dst_thread_addr_offset,
680  dst_wave_addr_offset + i * sizeof(half2_t),
681  0);
682  });
683  }
684  else if constexpr(N == 8)
685  {
686  vector_type<half_t, 8> tmp{src_thread_data};
687 
688  static_for<0, 4, 1>{}([&](auto i) {
689  llvm_amdgcn_raw_buffer_atomic_add_fp16x2(tmp.AsType<half2_t>()[i],
690  dst_wave_buffer_resource,
691  dst_thread_addr_offset,
692  dst_wave_addr_offset + i * sizeof(half2_t),
693  0);
694  });
695  }
696  }
697  else if constexpr(is_same<T, int32_t>::value)
698  {
699  if constexpr(N == 1)
700  {
701  llvm_amdgcn_raw_buffer_atomic_add_i32(src_thread_data,
702  dst_wave_buffer_resource,
703  dst_thread_addr_offset,
704  dst_wave_addr_offset,
705  0);
706  }
707  else if constexpr(N == 2)
708  {
709  vector_type<int32_t, 2> tmp{src_thread_data};
710 
711  llvm_amdgcn_raw_buffer_atomic_add_i32(tmp.AsType<int32_t>()[Number<0>{}],
712  dst_wave_buffer_resource,
713  dst_thread_addr_offset,
714  dst_wave_addr_offset,
715  0);
716 
717  llvm_amdgcn_raw_buffer_atomic_add_i32(tmp.AsType<int32_t>()[Number<1>{}],
718  dst_wave_buffer_resource,
719  dst_thread_addr_offset,
720  dst_wave_addr_offset + sizeof(int32_t),
721  0);
722  }
723  else if constexpr(N == 4)
724  {
725  vector_type<int32_t, 4> tmp{src_thread_data};
726 
727  llvm_amdgcn_raw_buffer_atomic_add_i32(tmp.AsType<int32_t>()[Number<0>{}],
728  dst_wave_buffer_resource,
729  dst_thread_addr_offset,
730  dst_wave_addr_offset,
731  0);
732 
733  llvm_amdgcn_raw_buffer_atomic_add_i32(tmp.AsType<int32_t>()[Number<1>{}],
734  dst_wave_buffer_resource,
735  dst_thread_addr_offset,
736  dst_wave_addr_offset + sizeof(int32_t),
737  0);
738 
739  llvm_amdgcn_raw_buffer_atomic_add_i32(tmp.AsType<int32_t>()[Number<2>{}],
740  dst_wave_buffer_resource,
741  dst_thread_addr_offset,
742  dst_wave_addr_offset + 2 * sizeof(int32_t),
743  0);
744 
745  llvm_amdgcn_raw_buffer_atomic_add_i32(tmp.AsType<int32_t>()[Number<3>{}],
746  dst_wave_buffer_resource,
747  dst_thread_addr_offset,
748  dst_wave_addr_offset + 3 * sizeof(int32_t),
749  0);
750  }
751  }
752 }
753 
754 template <typename T, index_t N>
755 __device__ void amd_buffer_atomic_max_impl(const typename vector_type<T, N>::type src_thread_data,
756  int32x4_t dst_wave_buffer_resource,
757  index_t dst_thread_addr_offset,
758  index_t dst_wave_addr_offset)
759 {
760  static_assert((is_same<T, double>::value && (N == 1 || N == 2 || N == 4)),
761  "wrong! not implemented");
762  if constexpr(is_same<T, double>::value)
763  {
764  if constexpr(N == 1)
765  {
766  llvm_amdgcn_raw_buffer_atomic_max_fp64(src_thread_data,
767  dst_wave_buffer_resource,
768  dst_thread_addr_offset,
769  dst_wave_addr_offset,
770  0);
771  }
772  else if constexpr(N == 2)
773  {
774  vector_type<double, 2> tmp{src_thread_data};
775 
776  llvm_amdgcn_raw_buffer_atomic_max_fp64(tmp.AsType<double>()[Number<0>{}],
777  dst_wave_buffer_resource,
778  dst_thread_addr_offset,
779  dst_wave_addr_offset,
780  0);
781 
782  llvm_amdgcn_raw_buffer_atomic_max_fp64(tmp.AsType<double>()[Number<1>{}],
783  dst_wave_buffer_resource,
784  dst_thread_addr_offset,
785  dst_wave_addr_offset + sizeof(double),
786  0);
787  }
788  else if constexpr(N == 4)
789  {
790  vector_type<double, 4> tmp{src_thread_data};
791 
792  llvm_amdgcn_raw_buffer_atomic_max_fp64(tmp.AsType<double>()[Number<0>{}],
793  dst_wave_buffer_resource,
794  dst_thread_addr_offset,
795  dst_wave_addr_offset,
796  0);
797 
798  llvm_amdgcn_raw_buffer_atomic_max_fp64(tmp.AsType<double>()[Number<1>{}],
799  dst_wave_buffer_resource,
800  dst_thread_addr_offset,
801  dst_wave_addr_offset + sizeof(double),
802  0);
803 
804  llvm_amdgcn_raw_buffer_atomic_max_fp64(tmp.AsType<double>()[Number<2>{}],
805  dst_wave_buffer_resource,
806  dst_thread_addr_offset,
807  dst_wave_addr_offset + 2 * sizeof(double),
808  0);
809 
810  llvm_amdgcn_raw_buffer_atomic_max_fp64(tmp.AsType<double>()[Number<3>{}],
811  dst_wave_buffer_resource,
812  dst_thread_addr_offset,
813  dst_wave_addr_offset + 3 * sizeof(double),
814  0);
815  }
816  }
817 }
818 
819 // buffer_load requires:
820 // 1) p_src_wave must point to global memory space
821 // 2) p_src_wave must be a wavewise pointer.
822 // It is user's responsibility to make sure that is true.
823 template <typename T,
824  index_t N,
825  AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
826 __device__ typename vector_type_maker<T, N>::type::type
827 amd_buffer_load_invalid_element_return_zero(const T* p_src_wave,
828  index_t src_thread_element_offset,
829  bool src_thread_element_valid,
830  index_t src_element_space_size)
831 {
832  const int32x4_t src_wave_buffer_resource =
833  make_wave_buffer_resource(p_src_wave, src_element_space_size);
834 
835  index_t src_thread_addr_offset = src_thread_element_offset * sizeof(T);
836 
837  using vector_t = typename vector_type_maker<T, N>::type::type;
838  using scalar_t = typename scalar_type<vector_t>::type;
839 
840  constexpr index_t vector_size = scalar_type<vector_t>::vector_size;
841 
842 #if CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK
843  uint32_t src_addr_shift = src_thread_element_valid ? 0 : 0x80000000;
844  return amd_buffer_load_impl<scalar_t, vector_size, coherence>(
845  src_wave_buffer_resource, src_addr_shift + src_thread_addr_offset, 0);
846 
847 #else
848 
849  vector_t tmp{amd_buffer_load_impl<scalar_t, vector_size, coherence>(
850  src_wave_buffer_resource, src_thread_addr_offset, 0)};
851  return src_thread_element_valid ? tmp : vector_t(0);
852 #endif
853 }
854 
855 // buffer_load requires:
856 // 1) p_src_wave must point to global memory space
857 // 2) p_src_wave must be a wavewise pointer.
858 // It is user's responsibility to make sure that is true.
859 template <typename T,
860  index_t N,
861  AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
862 __device__ typename vector_type_maker<T, N>::type::type
863 amd_buffer_load_invalid_element_return_customized_value(const T* p_src_wave,
864  index_t src_thread_element_offset,
865  bool src_thread_element_valid,
866  index_t src_element_space_size,
867  T customized_value)
868 {
869  const int32x4_t src_wave_buffer_resource =
870  make_wave_buffer_resource(p_src_wave, src_element_space_size);
871 
872  index_t src_thread_addr_offset = src_thread_element_offset * sizeof(T);
873 
874  using vector_t = typename vector_type_maker<T, N>::type::type;
875  using scalar_t = typename scalar_type<vector_t>::type;
876 
877  constexpr index_t vector_size = scalar_type<vector_t>::vector_size;
878 
879  vector_t tmp{amd_buffer_load_impl<scalar_t, vector_size, coherence>(
880  src_wave_buffer_resource, src_thread_addr_offset, 0)};
881 
882  return src_thread_element_valid ? tmp : vector_t(customized_value);
883 }
884 
885 // buffer_store requires:
886 // 1) p_dst_wave must point to global memory
887 // 2) p_dst_wave must be a wavewise pointer.
888 // It is user's responsibility to make sure that is true.
889 template <typename T,
890  index_t N,
891  AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
892 __device__ void amd_buffer_store(const typename vector_type_maker<T, N>::type::type src_thread_data,
893  T* p_dst_wave,
894  const index_t dst_thread_element_offset,
895  const bool dst_thread_element_valid,
896  const index_t dst_element_space_size)
897 {
898  const int32x4_t dst_wave_buffer_resource =
899  make_wave_buffer_resource(p_dst_wave, dst_element_space_size);
900 
901  index_t dst_thread_addr_offset = dst_thread_element_offset * sizeof(T);
902 
903  using vector_t = typename vector_type_maker<T, N>::type::type;
904  using scalar_t = typename scalar_type<vector_t>::type;
905  constexpr index_t vector_size = scalar_type<vector_t>::vector_size;
906 
907 #if CK_EXPERIMENTAL_USE_BUFFER_STORE_OOB_CHECK_OFFSET_TRICK
908  uint32_t dst_addr_shift = dst_thread_element_valid ? 0 : 0x80000000;
909  amd_buffer_store_impl<scalar_t, vector_size, coherence>(
910  src_thread_data, dst_wave_buffer_resource, dst_addr_shift + dst_thread_addr_offset, 0);
911 #else
912  if(dst_thread_element_valid)
913  {
914  amd_buffer_store_impl<scalar_t, vector_size, coherence>(
915  src_thread_data, dst_wave_buffer_resource, dst_thread_addr_offset, 0);
916  }
917 #endif
918 }
919 
920 // buffer_atomic_add requires:
921 // 1) p_dst_wave must point to global memory
922 // 2) p_dst_wave must be a wavewise pointer.
923 // It is user's responsibility to make sure that is true.
924 template <typename T, index_t N>
925 __device__ void
926 amd_buffer_atomic_add(const typename vector_type_maker<T, N>::type::type src_thread_data,
927  T* p_dst_wave,
928  const index_t dst_thread_element_offset,
929  const bool dst_thread_element_valid,
930  const index_t dst_element_space_size)
931 {
932  const int32x4_t dst_wave_buffer_resource =
933  make_wave_buffer_resource(p_dst_wave, dst_element_space_size);
934 
935  index_t dst_thread_addr_offset = dst_thread_element_offset * sizeof(T);
936 
937  using vector_t = typename vector_type_maker<T, N>::type::type;
938  using scalar_t = typename scalar_type<vector_t>::type;
939  constexpr index_t vector_size = scalar_type<vector_t>::vector_size;
940 
941  if constexpr(is_same<T, bhalf_t>::value)
942  {
943  if(dst_thread_element_valid)
944  {
945  amd_global_atomic_add_impl<scalar_t, vector_size>(
946  src_thread_data, p_dst_wave + dst_thread_element_offset);
947  }
948  }
949  else
950  {
951 #if CK_EXPERIMENTAL_USE_BUFFER_ATOMIC_ADD_OOB_CHECK_OFFSET_TRICK
952  uint32_t dst_addr_shift = dst_thread_element_valid ? 0 : 0x80000000;
953 
954  amd_buffer_atomic_add_impl<scalar_t, vector_size>(
955  src_thread_data, dst_wave_buffer_resource, dst_addr_shift + dst_thread_addr_offset, 0);
956 #else
957  if(dst_thread_element_valid)
958  {
959  amd_buffer_atomic_add_impl<scalar_t, vector_size>(
960  src_thread_data, dst_wave_buffer_resource, dst_thread_addr_offset, 0);
961  }
962 #endif
963  }
964 }
965 
966 // buffer_atomic_max requires:
967 // 1) p_dst_wave must point to global memory
968 // 2) p_dst_wave must be a wavewise pointer.
969 // It is user's responsibility to make sure that is true.
970 template <typename T, index_t N>
971 __device__ void
972 amd_buffer_atomic_max(const typename vector_type_maker<T, N>::type::type src_thread_data,
973  T* p_dst_wave,
974  const index_t dst_thread_element_offset,
975  const bool dst_thread_element_valid,
976  const index_t dst_element_space_size)
977 {
978  const int32x4_t dst_wave_buffer_resource =
979  make_wave_buffer_resource(p_dst_wave, dst_element_space_size);
980 
981  index_t dst_thread_addr_offset = dst_thread_element_offset * sizeof(T);
982 
983  using vector_t = typename vector_type_maker<T, N>::type::type;
984  using scalar_t = typename scalar_type<vector_t>::type;
985  constexpr index_t vector_size = scalar_type<vector_t>::vector_size;
986 
987 #if CK_EXPERIMENTAL_USE_BUFFER_ATOMIC_MAX_OOB_CHECK_OFFSET_TRICK
988  uint32_t dst_addr_shift = dst_thread_element_valid ? 0 : 0x80000000;
989 
990  amd_buffer_atomic_max_impl<scalar_t, vector_size>(
991  src_thread_data, dst_wave_buffer_resource, dst_addr_shift + dst_thread_addr_offset, 0);
992 #else
993  if(dst_thread_element_valid)
994  {
995  amd_buffer_atomic_max_impl<scalar_t, vector_size>(
996  src_thread_data, dst_wave_buffer_resource, dst_thread_addr_offset, 0);
997  }
998 #endif
999 }
1000 
1001 // Direct loads from global to LDS.
1002 __device__ void
1003 llvm_amdgcn_raw_buffer_load_lds(int32x4_t rsrc,
1004  __attribute__((address_space(3))) uint32_t* lds_ptr,
1005  index_t size,
1006  index_t voffset,
1007  index_t soffset,
1008  index_t offset,
1009  index_t aux) __asm("llvm.amdgcn.raw.buffer.load.lds");
1010 
1011 template <typename T, index_t NumElemsPerThread>
1012 __device__ void amd_direct_load_global_to_lds(const T* global_base_ptr,
1013  const index_t global_offset,
1014  T* lds_base_ptr,
1015  const index_t lds_offset,
1016  const bool is_valid,
1017  const index_t src_element_space_size)
1018 {
1019  // Direct loads require that each thread reads and writes exactly a single DWORD.
1020  constexpr auto dword_bytes = 4;
1021  constexpr auto bytes_per_thread = sizeof(T) * NumElemsPerThread;
1022  static_assert(bytes_per_thread == dword_bytes);
1023 
1024 #ifndef CK_CODE_GEN_RTC
1025  const uint32_t* global_ptr =
1026  reinterpret_cast<uint32_t*>(reinterpret_cast<uintptr_t>(global_base_ptr));
1027 #else
1028  const uint32_t* global_ptr =
1029  reinterpret_cast<uint32_t*>(reinterpret_cast<size_t>(global_base_ptr));
1030 #endif
1031  const int32x4_t src_resource = make_wave_buffer_resource(global_ptr, src_element_space_size);
1032  const index_t global_offset_bytes = is_valid ? global_offset * sizeof(T) : 0x80000000;
1033 
1034 #if CK_USE_AMD_LDS_DIRECT_LOAD_INLINE_ASM
1035  T* lds_ptr = lds_base_ptr + lds_offset;
1036 #ifndef CK_CODE_GEN_RTC
1037  auto const lds_ptr_sgpr =
1038  __builtin_amdgcn_readfirstlane((reinterpret_cast<uintptr_t>(lds_ptr)));
1039 #else
1040  auto const lds_ptr_sgpr = __builtin_amdgcn_readfirstlane((reinterpret_cast<size_t>(lds_ptr)));
1041 #endif
1042  asm volatile("s_mov_b32 m0, %0; \n\t"
1043  "buffer_load_dword %1, %2, 0 offen lds;\n\t" ::"s"(lds_ptr_sgpr),
1044  "v"(global_offset_bytes),
1045  "s"(src_resource)
1046  : "memory");
1047 #else
1048  // LDS pointer must be attributed with the LDS address space.
1049  __attribute__((address_space(3))) uint32_t* lds_ptr =
1050 #ifndef CK_CODE_GEN_RTC
1051  reinterpret_cast<__attribute__((address_space(3))) uint32_t*>(
1052  reinterpret_cast<uintptr_t>(lds_base_ptr + lds_offset));
1053 #else
1054  reinterpret_cast<__attribute__((address_space(3))) uint32_t*>(
1055  reinterpret_cast<size_t>(lds_base_ptr + lds_offset));
1056 #endif
1057 
1058  llvm_amdgcn_raw_buffer_load_lds(
1059  src_resource, lds_ptr, sizeof(uint32_t), global_offset_bytes, 0, 0, 0);
1060 #endif
1061 }
1062 
1063 } // namespace ck
CK_BUFFER_RESOURCE_3RD_DWORD
#define CK_BUFFER_RESOURCE_3RD_DWORD
Definition: ck.hpp:82
data_type.hpp
ck_tile::int8_t
int8_t int8_t
Definition: int8.hpp:20
ck
Definition: ck.hpp:264
ck::llvm_amdgcn_raw_buffer_store_fp32
__device__ void llvm_amdgcn_raw_buffer_store_fp32(float vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.f32")
ck::make_wave_buffer_resource_with_default_range
__device__ int32x4_t make_wave_buffer_resource_with_default_range(T *p_wave)
Definition: amd_buffer_addressing.hpp:38
ck::StaticallyIndexedArray
typename detail::StaticallyIndexedArrayImpl< T, N >::type StaticallyIndexedArray
Definition: statically_indexed_array.hpp:45
ck::amd_buffer_store
__device__ void amd_buffer_store(const typename vector_type_maker< T, N >::type::type src_thread_data, T *p_dst_wave, const index_t dst_thread_element_offset, const bool dst_thread_element_valid, const index_t dst_element_space_size)
Definition: amd_buffer_addressing.hpp:892
ck::amd_direct_load_global_to_lds
__device__ void amd_direct_load_global_to_lds(const T *global_base_ptr, const index_t global_offset, T *lds_base_ptr, const index_t lds_offset, const bool is_valid, const index_t src_element_space_size)
Definition: amd_buffer_addressing.hpp:1012
ck::amd_buffer_atomic_max
__device__ void amd_buffer_atomic_max(const typename vector_type_maker< T, N >::type::type src_thread_data, T *p_dst_wave, const index_t dst_thread_element_offset, const bool dst_thread_element_valid, const index_t dst_element_space_size)
Definition: amd_buffer_addressing.hpp:972
ck::llvm_amdgcn_raw_buffer_load_i32x4
__device__ int32x4_t llvm_amdgcn_raw_buffer_load_i32x4(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i32")
ck::llvm_amdgcn_raw_buffer_store_fp32x4
__device__ void llvm_amdgcn_raw_buffer_store_fp32x4(float4_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4f32")
ck::llvm_amdgcn_raw_buffer_store_i8x2
__device__ void llvm_amdgcn_raw_buffer_store_i8x2(int8x2_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i8")
ck::amd_buffer_store_impl
__device__ void amd_buffer_store_impl(const typename vector_type< T, N >::type src_thread_data, int32x4_t dst_wave_buffer_resource, index_t dst_thread_addr_offset, index_t dst_wave_addr_offset)
Definition: amd_buffer_addressing.hpp:542
ck::llvm_amdgcn_raw_buffer_load_fp32
__device__ float llvm_amdgcn_raw_buffer_load_fp32(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.f32")
ck::llvm_amdgcn_raw_buffer_load_i32x2
__device__ int32x2_t llvm_amdgcn_raw_buffer_load_i32x2(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i32")
ck::AmdBufferCoherenceEnum
AmdBufferCoherenceEnum
Definition: amd_buffer_addressing.hpp:295
ck::bhalf4_t
typename vector_type< bhalf_t, 4 >::type bhalf4_t
Definition: data_type.hpp:2498
ck::make_wave_buffer_resource
__device__ int32x4_t make_wave_buffer_resource(T *p_wave, index_t element_space_size)
Definition: amd_buffer_addressing.hpp:23
ck::int32x2_t
typename vector_type< int32_t, 2 >::type int32x2_t
Definition: data_type.hpp:2505
ck::llvm_amdgcn_raw_buffer_store_fp16
__device__ void llvm_amdgcn_raw_buffer_store_fp16(half_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.f16")
ck::int8x2_t
typename vector_type< int8_t, 2 >::type int8x2_t
Definition: data_type.hpp:2513
ck::llvm_amdgcn_raw_buffer_load_fp16
__device__ half_t llvm_amdgcn_raw_buffer_load_fp16(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.f16")
ck::llvm_amdgcn_raw_buffer_load_i16
__device__ bhalf_t llvm_amdgcn_raw_buffer_load_i16(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i16")
ck::llvm_amdgcn_raw_buffer_load_lds
__device__ void llvm_amdgcn_raw_buffer_load_lds(int32x4_t rsrc, uint32_t *lds_ptr, index_t size, index_t voffset, index_t soffset, index_t offset, index_t aux) __asm("llvm.amdgcn.raw.buffer.load.lds")
ck::llvm_amdgcn_raw_buffer_store_i8x4
__device__ void llvm_amdgcn_raw_buffer_store_i8x4(int8x4_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i8")
ck::llvm_amdgcn_raw_buffer_load_i8x4
__device__ int8x4_t llvm_amdgcn_raw_buffer_load_i8x4(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i8")
ck::amd_buffer_atomic_add_impl
__device__ void amd_buffer_atomic_add_impl(const typename vector_type< T, N >::type src_thread_data, int32x4_t dst_wave_buffer_resource, index_t dst_thread_addr_offset, index_t dst_wave_addr_offset)
Definition: amd_buffer_addressing.hpp:597
ck::llvm_amdgcn_raw_buffer_load_fp32x4
__device__ float4_t llvm_amdgcn_raw_buffer_load_fp32x4(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4f32")
ck::llvm_amdgcn_raw_buffer_load_i16x2
__device__ bhalf2_t llvm_amdgcn_raw_buffer_load_i16x2(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i16")
ck::float2_t
typename vector_type< float, 2 >::type float2_t
Definition: data_type.hpp:2481
ck::llvm_amdgcn_raw_buffer_store_i32x4
__device__ void llvm_amdgcn_raw_buffer_store_i32x4(int32x4_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i32")
ck::llvm_amdgcn_raw_buffer_atomic_add_fp32
__device__ float llvm_amdgcn_raw_buffer_atomic_add_fp32(float vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.f32")
ck::half4_t
typename vector_type< half_t, 4 >::type half4_t
Definition: data_type.hpp:2490
ck::amd_global_atomic_add_impl
__device__ void amd_global_atomic_add_impl(const typename vector_type< T, N >::type src_thread_data, T *addr)
Definition: amd_buffer_addressing.hpp:569
ck::half_t
_Float16 half_t
Definition: data_type.hpp:25
ck::llvm_amdgcn_raw_buffer_store_i32x2
__device__ void llvm_amdgcn_raw_buffer_store_i32x2(int32x2_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i32")
ck::llvm_amdgcn_raw_buffer_atomic_add_fp16x2
__device__ half2_t llvm_amdgcn_raw_buffer_atomic_add_fp16x2(half2_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.v2f16")
ck::bhalf_t
ushort bhalf_t
Definition: data_type.hpp:24
ck::amd_buffer_load_impl
__device__ vector_type< T, N >::type amd_buffer_load_impl(int32x4_t src_wave_buffer_resource, index_t src_thread_addr_offset, index_t src_wave_addr_offset)
Definition: amd_buffer_addressing.hpp:419
ck::bhalf2_t
typename vector_type< bhalf_t, 2 >::type bhalf2_t
Definition: data_type.hpp:2497
ck::llvm_amdgcn_raw_buffer_store_i16
__device__ void llvm_amdgcn_raw_buffer_store_i16(bhalf_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i16")
ck::amd_buffer_load_invalid_element_return_customized_value
__device__ vector_type_maker< T, N >::type::type amd_buffer_load_invalid_element_return_customized_value(const T *p_src_wave, index_t src_thread_element_offset, bool src_thread_element_valid, index_t src_element_space_size, T customized_value)
Definition: amd_buffer_addressing.hpp:863
ck::llvm_amdgcn_raw_buffer_store_i16x2
__device__ void llvm_amdgcn_raw_buffer_store_i16x2(bhalf2_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i16")
ck::float4_t
typename vector_type< float, 4 >::type float4_t
Definition: data_type.hpp:2482
ck::llvm_amdgcn_raw_buffer_load_fp16x2
__device__ half2_t llvm_amdgcn_raw_buffer_load_fp16x2(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2f16")
ck::amd_buffer_atomic_add
__device__ void amd_buffer_atomic_add(const typename vector_type_maker< T, N >::type::type src_thread_data, T *p_dst_wave, const index_t dst_thread_element_offset, const bool dst_thread_element_valid, const index_t dst_element_space_size)
Definition: amd_buffer_addressing.hpp:926
ck::llvm_amdgcn_raw_buffer_store_fp32x2
__device__ void llvm_amdgcn_raw_buffer_store_fp32x2(float2_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2f32")
ck::llvm_amdgcn_raw_buffer_atomic_max_fp64
__device__ double llvm_amdgcn_raw_buffer_atomic_max_fp64(double vdata, int32x4_t rsrc, int voffset, int soffset, int glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fmax.f64")
ck::llvm_amdgcn_raw_buffer_load_i16x4
__device__ bhalf4_t llvm_amdgcn_raw_buffer_load_i16x4(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i16")
ck::amd_buffer_load_invalid_element_return_zero
__device__ vector_type_maker< T, N >::type::type amd_buffer_load_invalid_element_return_zero(const T *p_src_wave, index_t src_thread_element_offset, bool src_thread_element_valid, index_t src_element_space_size)
Definition: amd_buffer_addressing.hpp:827
ck::llvm_amdgcn_raw_buffer_load_i8
__device__ int8_t llvm_amdgcn_raw_buffer_load_i8(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i8")
ck::llvm_amdgcn_raw_buffer_load_fp32x2
__device__ float2_t llvm_amdgcn_raw_buffer_load_fp32x2(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2f32")
ck::half2_t
typename vector_type< half_t, 2 >::type half2_t
Definition: data_type.hpp:2489
ck::amd_buffer_atomic_max_impl
__device__ void amd_buffer_atomic_max_impl(const typename vector_type< T, N >::type src_thread_data, int32x4_t dst_wave_buffer_resource, index_t dst_thread_addr_offset, index_t dst_wave_addr_offset)
Definition: amd_buffer_addressing.hpp:755
ck::llvm_amdgcn_raw_buffer_atomic_add_i32
__device__ int32_t llvm_amdgcn_raw_buffer_atomic_add_i32(int32_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.add.i32")
ck::int32x4_t
typename vector_type< int32_t, 4 >::type int32x4_t
Definition: data_type.hpp:2506
ck::index_t
int32_t index_t
Definition: ck.hpp:289
ck::llvm_amdgcn_raw_buffer_store_fp16x2
__device__ void llvm_amdgcn_raw_buffer_store_fp16x2(half2_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2f16")
ck::int8x4_t
typename vector_type< int8_t, 4 >::type int8x4_t
Definition: data_type.hpp:2514
ck::llvm_amdgcn_raw_buffer_store_i8
__device__ void llvm_amdgcn_raw_buffer_store_i8(int8_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i8")
ck::llvm_amdgcn_raw_buffer_load_i32
__device__ int32_t llvm_amdgcn_raw_buffer_load_i32(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i32")
ck::amd_buffer_store_impl_raw
__device__ void amd_buffer_store_impl_raw(const typename vector_type< int8_t, N >::type src_thread_data, int32x4_t dst_wave_buffer_resource, index_t dst_thread_addr_offset, index_t dst_wave_addr_offset)
Definition: amd_buffer_addressing.hpp:444
ck::llvm_amdgcn_raw_buffer_store_i16x4
__device__ void llvm_amdgcn_raw_buffer_store_i16x4(bhalf4_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i16")
ck::amd_buffer_load_impl_raw
__device__ vector_type< int8_t, N >::type amd_buffer_load_impl_raw(int32x4_t src_wave_buffer_resource, index_t src_thread_addr_offset, index_t src_wave_addr_offset)
Definition: amd_buffer_addressing.hpp:315
ck::llvm_amdgcn_raw_buffer_load_i8x2
__device__ int8x2_t llvm_amdgcn_raw_buffer_load_i8x2(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i8")
ck::remove_cv_t
typename remove_cv< T >::type remove_cv_t
Definition: type.hpp:298
ck::llvm_amdgcn_raw_buffer_store_i32
__device__ void llvm_amdgcn_raw_buffer_store_i32(int32_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i32")
ck::llvm_amdgcn_raw_buffer_load_fp16x4
__device__ half4_t llvm_amdgcn_raw_buffer_load_fp16x4(int32x4_t srsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4f16")
ck::llvm_amdgcn_raw_buffer_store_fp16x4
__device__ void llvm_amdgcn_raw_buffer_store_fp16x4(half4_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4f16")
ck::integral_constant
Definition: integral_constant.hpp:10
ck::is_same
Definition: type.hpp:177
ck::scalar_type
Definition: data_type.hpp:394
ck::static_for
Definition: functional2.hpp:31
ck::vector_type_maker
Definition: data_type.hpp:367
ck::vector_type
Definition: data_type.hpp:347
ck::BufferResource
Definition: amd_buffer_addressing.hpp:11
ck::BufferResource::config
StaticallyIndexedArray< int32_t, 4 > config
Definition: amd_buffer_addressing.hpp:19
ck::BufferResource::content
int32x4_t content
Definition: amd_buffer_addressing.hpp:16
ck::BufferResource::range
StaticallyIndexedArray< int32_t, 4 > range
Definition: amd_buffer_addressing.hpp:18
ck::BufferResource::BufferResource
constexpr __device__ BufferResource()
Definition: amd_buffer_addressing.hpp:12
ck::BufferResource::address
StaticallyIndexedArray< T *, 2 > address
Definition: amd_buffer_addressing.hpp:17