<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/54767>54767</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
AMDGPU: incorrect assembly generated when using block of inline assembly using same result operand
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
Epliz
</td>
</tr>
</table>
<pre>
Hi,
I am using clang 14 from ROCM 5.0.2:
```
HIP version: 5.0.13601-ded05588
AMD clang version 14.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-5.0.2 22065 030a405a181176f1a7749819092f4ef8ea5f0758)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm-5.0.2/llvm/bin
```
I am trying to implement some kind of matrix multiplication kernel, currently trying to get it to work for gfx1031.
I had to resort (against my will) to using inline assembly to get some kind of decent performance (the compiler misses quite a lot of potential of optimizations such as using scalar loads for matrix A, or fails to emit vectorized s_load/global_load operations).
But even using inline assembly seems to not lead to proper generated code (when compiling with `-O2`).
Here is my current code:
```
#ifndef __BLOCKED_REDUCE_MATMUL_V3_HPP__
#define __BLOCKED_REDUCE_MATMUL_V3_HPP__
#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>
__device__ inline int32_t vgpr_to_sgpr(int32_t vgpr_val) {
int32_t sgpr_val;
asm ("v_readfirstlane_b32 %0, %1" : "=s"(sgpr_val) : "v"(vgpr_val));
return sgpr_val;
}
__device__ inline float fvgpr_to_fsgpr(float vgpr_val) {
float sgpr_val;
asm ("v_readfirstlane_b32 %0, %1" : "=s"(sgpr_val) : "v"(vgpr_val));
return sgpr_val;
}
#define TILE_BLOCKED_REDV3_ROWS 16
#define TILE_BLOCKED_REDV3_COLS (warpSize * 2)
#define TILE_BLOCKED_REDV3_DEPTH 2
#define BLOCK_MUL_REDV3XX_(ii, av_off, avii) do { \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r1[ii]) : "s"(avii[0]), "v"(b_0.x)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r2[ii]) : "s"(avii[0]), "v"(b_0.y)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r1[ii]) : "s"(avii[1]), "v"(b_1.x)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r2[ii]) : "s"(avii[1]), "v"(b_1.y)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r1[ii]) : "s"(avii[2]), "v"(b_2.x)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r2[ii]) : "s"(avii[2]), "v"(b_2.y)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r1[ii]) : "s"(avii[3]), "v"(b_3.x)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r2[ii]) : "s"(avii[3]), "v"(b_3.y)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r1[ii]) : "s"(avii[4]), "v"(b_4.x)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r2[ii]) : "s"(avii[4]), "v"(b_4.y)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r1[ii]) : "s"(avii[5]), "v"(b_5.x)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r2[ii]) : "s"(avii[5]), "v"(b_5.y)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r1[ii]) : "s"(avii[6]), "v"(b_6.x)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r2[ii]) : "s"(avii[6]), "v"(b_6.y)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r1[ii]) : "s"(avii[7]), "v"(b_7.x)); \
asm volatile ("v_fmac_f32_e32 %0, %1, %2;" : "=v"(r2[ii]) : "s"(avii[7]), "v"(b_7.y)); \
} while(0)
#define BLOCK_MUL_REDV3XX(ii, av_off) BLOCK_MUL_REDV3XX_(ii, av_off, av##ii)
#define GLOBAL_ADD4V3(ii) do { \
c[((ii) + 0) * TILE_BLOCKED_REDV3_COLS + j] += r1[((ii) + 0)]; \
c[((ii) + 0) * TILE_BLOCKED_REDV3_COLS + j + 1] += r2[((ii) + 0)]; \
c[((ii) + 1) * TILE_BLOCKED_REDV3_COLS + j] += r1[((ii) + 1)]; \
c[((ii) + 1) * TILE_BLOCKED_REDV3_COLS + j + 1] += r2[((ii) + 1)]; \
c[((ii) + 2) * TILE_BLOCKED_REDV3_COLS + j] += r1[((ii) + 2)]; \
c[((ii) + 2) * TILE_BLOCKED_REDV3_COLS + j + 1] += r2[((ii) + 2)]; \
c[((ii) + 3) * TILE_BLOCKED_REDV3_COLS + j] += r1[((ii) + 3)]; \
c[((ii) + 3) * TILE_BLOCKED_REDV3_COLS + j + 1] += r2[((ii) + 3)]; \
} while (0)
typedef float mfloat8 __attribute__((ext_vector_type(8)));
__global__ void matmul_blocked_reducev3(int N, float* a, float* b, float* c) {
int32_t i = TILE_BLOCKED_REDV3_ROWS * blockIdx.x;
int32_t wid = vgpr_to_sgpr(threadIdx.x / warpSize);
int32_t lid = threadIdx.x % warpSize;
int32_t j = blockIdx.y * TILE_BLOCKED_REDV3_COLS + 2 * lid;
float r1[TILE_BLOCKED_REDV3_ROWS] = {0.0f};
float r2[TILE_BLOCKED_REDV3_ROWS] = {0.0f};
if (i + (TILE_BLOCKED_REDV3_ROWS - 1) < N) {
if (j + warpSize < N) {
int32_t Nk = N / TILE_BLOCKED_REDV3_DEPTH;
int32_t k_start = vgpr_to_sgpr(wid * Nk);
int32_t k_end = vgpr_to_sgpr(k_start + Nk);
int32_t av_off_0 = vgpr_to_sgpr(i* N + k_start);
b = &b[k_start * N];
mfloat8 t8;
#pragma nounroll
for (int32_t k = k_start; k + 7 < k_end; k += 8) {
float2 b_0;//= *((float2*) &b[j]);
asm volatile ("global_load_dwordx2 %0, %1, off;" : "=v"(b_0) : "v"((float2*) &b[j]));
b += N;
float2 b_1;// = *((float2*) &b[j]);
asm volatile ("global_load_dwordx2 %0, %1, off;" : "=v"(b_1) : "v"((float2*) &b[j]));
b += N;
float2 b_2;// = *((float2*) &b[j]);
asm volatile ("global_load_dwordx2 %0, %1, off;" : "=v"(b_2) : "v"((float2*) &b[j]));
b += N;
float2 b_3;// = *((float2*) &b[j]);
asm volatile ("global_load_dwordx2 %0, %1, off;" : "=v"(b_3) : "v"((float2*) &b[j]));
b += N;
float2 b_4;// = *((float2*) &b[j]);
asm volatile ("global_load_dwordx2 %0, %1, off;" : "=v"(b_4) : "v"((float2*) &b[j]));
b += N;
float2 b_5;// = *((float2*) &b[j]);
asm volatile ("global_load_dwordx2 %0, %1, off;" : "=v"(b_5) : "v"((float2*) &b[j]));
b += N;
float2 b_6;// = *((float2*) &b[j]);
asm volatile ("global_load_dwordx2 %0, %1, off;" : "=v"(b_6) : "v"((float2*) &b[j]));
b += N;
float2 b_7;// = *((float2*) &b[j]);
asm volatile ("global_load_dwordx2 %0, %1, off;" : "=v"(b_7) : "v"((float2*) &b[j]));
b += N;
int32_t av_off = av_off_0;
mfloat8 av0;
mfloat8 av1;
mfloat8 av2;
mfloat8 av3;
mfloat8 av4;
mfloat8 av5;
mfloat8 av6;
mfloat8 av7;
mfloat8 av8;
mfloat8 av9;
mfloat8 av10;
mfloat8 av11;
mfloat8 av12;
mfloat8 av13;
mfloat8 av14;
mfloat8 av15;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av0) : "s"(((mfloat8*) &a[av_off])));
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av1) : "s"(((mfloat8*) &a[av_off])));
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av2) : "s"(((mfloat8*) &a[av_off])));
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av3) : "s"(((mfloat8*) &a[av_off])));
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av4) : "s"(((mfloat8*) &a[av_off])));
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av5) : "s"(((mfloat8*) &a[av_off])));
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av6) : "s"(((mfloat8*) &a[av_off])));
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av7) : "s"(((mfloat8*) &a[av_off])));
asm volatile("s_waitcnt vmcnt(0) lgkmcnt(7);");
BLOCK_MUL_REDV3XX(0, av_off);
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av8) : "s"(((mfloat8*) &a[av_off])));
asm volatile("s_waitcnt lgkmcnt(7);");
BLOCK_MUL_REDV3XX(1, av_off);
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av9) : "s"(((mfloat8*) &a[av_off])));
asm volatile("s_waitcnt lgkmcnt(7);");
BLOCK_MUL_REDV3XX(2, av_off);
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av10) : "s"(((mfloat8*) &a[av_off])));
asm volatile("s_waitcnt lgkmcnt(7);");
BLOCK_MUL_REDV3XX(3, av_off);
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av11) : "s"(((mfloat8*) &a[av_off])));
asm volatile("s_waitcnt lgkmcnt(7);");
BLOCK_MUL_REDV3XX(4, av_off);
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av12) : "s"(((mfloat8*) &a[av_off])));
asm volatile("s_waitcnt lgkmcnt(7);");
BLOCK_MUL_REDV3XX(5, av_off);
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av13) : "s"(((mfloat8*) &a[av_off])));
asm volatile("s_waitcnt lgkmcnt(7);");
BLOCK_MUL_REDV3XX(6, av_off);
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av14) : "s"(((mfloat8*) &a[av_off])));
asm volatile("s_waitcnt lgkmcnt(7);");
BLOCK_MUL_REDV3XX(7, av_off);
av_off += N;
asm volatile("s_load_dwordx8 %0, %1;" : "=s"(av15) : "s"(((mfloat8*) &a[av_off])));
asm volatile("s_waitcnt lgkmcnt(7);");
BLOCK_MUL_REDV3XX(8, av_off);
asm volatile("s_waitcnt lgkmcnt(6);");
BLOCK_MUL_REDV3XX(9, av_off);
asm volatile("s_waitcnt lgkmcnt(5);");
BLOCK_MUL_REDV3XX(10, av_off);
asm volatile("s_waitcnt lgkmcnt(4);");
BLOCK_MUL_REDV3XX(11, av_off);
asm volatile("s_waitcnt lgkmcnt(3);");
BLOCK_MUL_REDV3XX(12, av_off);
asm volatile("s_waitcnt lgkmcnt(2);");
BLOCK_MUL_REDV3XX(13, av_off);
asm volatile("s_waitcnt lgkmcnt(1);");
BLOCK_MUL_REDV3XX(14, av_off);
asm volatile("s_waitcnt lgkmcnt(0);");
BLOCK_MUL_REDV3XX(15, av_off);
av_off_0 += 8;
}
}
}
__syncthreads();
// accumulate using global mem
//# pragma unroll_complete
//for (int32_t ii = 0; ii < TILE_BLOCKED_RED_ROWS; ii++) {
// scratchpad[ii * TILE_BLOCKED_RED_COLS + j] += r[ii];
//}
for (int32_t w = 0; w < TILE_BLOCKED_REDV3_DEPTH; w++) {
if (wid == w) {
GLOBAL_ADD4V3(0);
GLOBAL_ADD4V3(4);
GLOBAL_ADD4V3(8);
GLOBAL_ADD4V3(12);
}
__syncthreads();
}
}
void blocked_reduce_matmul_v3(size_t warpSize, size_t N, float* a_device, float* b_device, float* c_device) {
const size_t x_threads_per_block = warpSize;
size_t num_blocks_y = ((N + 2 * x_threads_per_block - 1) / (2 * x_threads_per_block));
size_t num_blocks_x = (N/TILE_BLOCKED_REDV3_ROWS);
const size_t y_threads_per_block = TILE_BLOCKED_REDV3_DEPTH;
dim3 block_dim(x_threads_per_block * y_threads_per_block);
dim3 grid_dim(num_blocks_x, num_blocks_y);
matmul_blocked_reducev3<<<grid_dim, block_dim, 0, 0>>>(N, a_device, b_device, c_device);
}
#endif /* __BLOCKED_REDUCE_MATMUL_V3_HPP__ */
```
The dest operands of the fmac operations from the `BLOCK_MUL_REDV3XX_` macros are not using the correct register.
In the generated assembly, I see the following:
``
;;#ASMSTART
s_load_dwordx8 s[8:15], s[80:81];
;;#ASMEND
.loc 10 194 59 ; ./blocked_reduce_matmul_v3.hpp:194:59
s_add_u32 s80, s82, s70
s_addc_u32 s81, s83, s71
.loc 10 187 9 ; ./blocked_reduce_matmul_v3.hpp:187:9
;;#ASMSTART
s_waitcnt vmcnt(0) lgkmcnt(7);
;;#ASMEND
.Ltmp869:
.loc 10 188 9 ; ./blocked_reduce_matmul_v3.hpp:188:9
;;#ASMSTART
v_fmac_f32_e32 v0, s0, v7;
;;#ASMEND
;;#ASMSTART
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s0, v8;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s1, v39;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s1, v40;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s2, v41;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s2, v42;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s3, v43;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s3, v44;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s4, v45;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s4, v46;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s5, v47;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s5, v48;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s6, v49;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v0, s6, v50;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v18, s7, v51;
;;#ASMEND
;;#ASMSTART
v_fmac_f32_e32 v9, s7, v52;
```
Here, `v0` as destination is not correct as far as I can tell. It should use `v18` and `v19` like at the two last instructions.
Any idea how to get it to work?
Thanks in advance.
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzVXFlz4roS_jXkRRXKK9gPeSCQnEndLFMzmXPPm0rYMmjihesFkvn1tyXZxgY5kBk4kBSFsZbur79uyx1Z8jTx366-sJ4x7mmTnja6QyRCRcbiGfJCAt-6hYI0idC3p_EDsvta3-iZI9m2N9DKjzj9cvcVLWmasSSGJqKtbg40_dKnvmbbjiObjR4mpeiyMaiAphrqGc48zxcZl2_cwmfG8nkx7XtJBCffiE-T-GlB43ESLYqcQlkYLqPLRZr8pF6O0sS7FPiQYWgDG2mmRizNJrqj68NBoJPh0HId3dVcI7Bo4FBiB9rQdnqGK4E9k3RGcw791RnggXVZxC9xsoovQxYXr5ezuCjbzVNKfBQlPg1560WSsdeSvTjLSRhSf8JSXgUYk0UO3wAukuhK2HCYsljJY8MRefrGPZEniEWLkEY0zlGWRBS9sNhHSYAikqfsFUVFmLNFyDySc0JfaBoDNmOMvCJNoVP41hAFRiKW81-rJH1BQZKiWfCqa6ber1TPwT6oT2mWpDl3DJkRBrah6A2tWAiiXV4v44TFQBBFJMtoNOWapIoWTp96HPuCpqAuIrFHudR8ThF4d8FCmqKIgYAM_a9gOQhDYZLzjoskh46MhPwEuGQR-yWMzFBWeHPQWqLIPBKSFLoRPxM2ldSMOA1wGhAWZhwbjcD4JURMkrJf1EcZ5n14uIXJlITiDDTRVKoBU0tarosc0SWNO8zOKI2EghiQh1QyCMEJkoCPmMsDbR5EDTd9NQdB0nYubAWhjiACLp8MHge1zvLKoilFLOPslw4VcrouxJ5hsiD2aYAwvr5_Gv_nZoK_3Ux-jG_ww-j54cc9_tvEX75-xbhuD425MXu2r7XEXlhwe8zxnC2AQvjGaQH-imh_3jNv9mqJyYI1W4tvjH26ZB7FuCKaxblpYHDdbJHiPMEZHIHIVvGSiNDsDa-lFFT3yqpqs64iWcQ90TOMJeZXdMDSLIeBieKpaUCNrfHQgaMOTZC8nGHwm2T8YDhZU6GsXMqaBhT-WWtMaV6k8TaW3nDyvu0BBGWOgsr0oLRdFndYLis_k90b0VXG5PPd_U0zLCEWvz399zvSB3u0HD_dfxeXG0kX3-Fyh98jZNRD_rt9Jzdfn79AYyUq0RTza0M0_ucfzGOR30cRWeIkCOQvXuIiP-GOQT17XDEivbBMQhhkQlq7I4iIhwMIWLrlCXE0OHUtp5TUp3rPvgZl9qThltJhAoV9rclKKaly2RRr_dfaX0dFaPw-wrd_B-FuDnU1Qv18OOxEeDYcGmqExvlw2InwbDg01QjN8-GwE-HZcGipEVrnw2EnwrPh0FYjtM-Hw06EZ8PhQI1wcD4cdiI8Gw6HaoTD8-GwE6GaQ8iM0WoO4KCVtk5XdyWhWzmou2-iCkL5v2msU9df90_Xo3s8mkysv81ShDqz9cBcQWnZpmdcI03-GL2To1-jn8AR_wHEIuF0lRTO46Y3f1-j-Nabeo0_1KsfxFL94Br3s_Qjeo2DWGocXON-ln5Er3kQS82Da9zPUrXeeoRBrSEGKsQxf1tQPoMkZxEicXAQxiTPUzYtcoqx1ENfcyxn0zDvAyVOOaI1pgCqeY1yig3D2Mt8PkMXFSGehon3Qn2cUr_w6NKUkzrokQ9MQi_ngbTOpq0zr2POhyHOSdf8gZDDVd_5r3CfaENdC1kBUC5mY9opF5PAoitIukXVBEN74qMSEpZC2r3sda_tLj9Fhxrg285oMEQL0LRlivShiMkOMkQIgTogUetrAZ-XMTcmkkRo_U53-VebF_B4YwIx_OhyzmU5qpljHgct71YyZPSvJ3aUTdd8Pr4IiI_CW12zPRuwq74vOMsJnwjfjgMRHsD740vb8-3uNFYGUS0XDGkLUImRt2usqSQxgUEIKoVuw5lKHxmDKXhyrRr6ydFhW3N12efOhihICxYpmUUExUkRp0kYNmv57HtjZlYSX8GCUehF4BwKlwlu6kLe0FE4sQxCA02xJpIw_nhIWjOSw5CsF6duZePPMufajEVF8teY-8f-Kkn91-0MkOdKHQkgh7U1HboDlQLYtOLgsfP62WBDr9lAZ0WHfho6jDOlwzgNHeaZ0mGehg7rTOmwTkOHfaZ02KehY3CmdAxOQ8fwTOkY_qt0tHMvwUSVhjUyJrdKlMiyo1hXFxvqYlNdbKmLbXXxQF08VBc76mK3w5wuMzvs1DsM1Tss1TtM1Vu2NsNLRlfWDCynHVibEVXPzzVTt6yOKPiUetdRRSCqykmzOrRUYV8Gy2aIHQ6y_vkgG58Psvn5IFufD7L9-SAPPh_k4aEgA6Tum3wFcEVY7sU5WkbwXc4vonD2Up4OS4GGoSJD9VRDaz3UOAF_zsFd_g5tf8KUfmqm3M_ClHFqpvTDZx5Hoso8OVWHz3iORJV1cqoOn2kdiSr75FQdPsM7ElWDk1N1-MzySFQNT07V4TPaI1Hl7KBqL7WDD6t1D6HW_nhmtCuJ3Euv9XG9u1KyvfSaH9e7K8HZS6_xcb27soW99Oof17vr1ruXXu3jenfcx1TjD3-KWz3ubLev90RsnDR-Ypy9xZ5cS5AJS5pKy0lb4nlFVIC1tNyvJGdeUUSjdsueYaLyYa58lIv5tqSQ5rTdbuOpLpOrK_hEoPw93nqqLtcGiHpuLP-0nuyWSFHmpST35gvii2V0SLXSoWOdTb3ubtP-Bl0bwFdr3Csl7MZiALRSAq-WIJRrQ7i81VYTtLVkTtuOjs0m1u4mzu4m-lbsbuy04UU7oqjuoNqjI5bvtNft4HI1j1i_k7FflFNdL4oZo7JoY1VPudOpvbhHVejVhRs0ewnfl1hKf8WlPXhBU7mwSHhbsc4GVX3iIpItM_xWPurg92q5mEIuqlGJrVao8OcjMEp2tVPc0bf1vlZ6gZ3brjU2rTkYpfVvHdbvtdrFZ5EpXYrhJ2BRcglGKrRsWihkzVLml6KapnKvNilX29W1NMwcy89a-LgJeow0-WXelB_B6bgVaM34aoTVGsVGxBsmjX244OXIMtq5SVI-K7ste6v29j7PKfIpOE5sM439jG9u5Zth-Qrkxt5TufOaV4AAxULegQZEeWmSIZJSse9UDvVyX22aim3RdMaynKbV3t5Y1K43pFb7Vzkbd3wTqwQC94FkBbI295jWc4DAFr9LmqPvD9-fR9-e65qNnDmDIRrucSNdrsQfywLw0MjR2-uOWjJvHmsnuH1wMBx0DemuhWwXqf_4eN3n-6o7xqX-fLHgQFwLvm23AZj4Pi5MA2VynjGTU0PZUGu38cpGumxkyka6CqczRF0w98YJCb05cvdifM_51neJ7t_n0cIZuGuPty1yDmCRs6dFG2vxl9It4rv9ALMrZLpFf7zHe2CcPwOjFi3ia2nudtlvy7a0Y8g2pGz9iLKNY8g2pWzziLKtY8i2pGz7iLIHx5BtS9l_eCG_K_so1-VAyj7KdSll2we-LnVH3qOk8ANfmG5TduPCVGU8_A0WYuZsoC15HX9vB8-AWCzfWMIykb5USQvUBiTlhzvkEUhZaBj20R1kuvOkCH1Ic0RKxO3jsvjqan7m8rOQvVBEcpHG5KsEhQQSLf7ykrTwRFbVerXGKH5DzKcEzZPV9stReuZts_HznMQvGQhDxF_yl5j0L_wr03dNl1zkLA_p1ehh8tfXH3zuj8VrY8o3hKyTLvH2D5mrycQa8r_N94mUrzUhEeXvYSnCOlu8KNLw6p035ZSvl2m-HAdOWZYVFP7Pu7Wt4WB4Mb-imgVZPPGGrje1DDvwjIE5HQa6SXxKjal7EZIpDbMrsaPEiOkKCRF8YsSeXLArQzMMzdJsyKdc3e3ToUk017W0ILDNYRD0LI1GhIV9jqOfpLOL9EpAmhazDCpDyEqzdSWYzWYxpUIdyCdFPk_Sq5tFyH5dCMVXAvj_AUS23Xg">