xgiovio/MultiPar
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Add files via upload

Browse Source
This commit is contained in:
Yutaka Sawada
2023-12-26 18:53:12 +09:00
committed by GitHub
parent 79d0b184b8
commit 1552fb8ec8
13 changed files with 920 additions and 430 deletions
Download Patch File Download Diff File Expand all files Collapse all files

194
source/par2j/source.cl
View File

@@ -1,10 +1,14 @@
void calc_table(__local uint *mtab, int id, int factor)
{
int i, sum = 0;
int i, sum, mask;
for (i = 0; i < 8; i++){
sum = (id & (1 << i)) ? (sum ^ factor) : sum;
factor = (factor & 0x8000) ? ((factor << 1) ^ 0x1100B) : (factor << 1);
mask = (id & 1) ? 0xFFFF : 0;
sum = mask & factor;
for (i = 1; i < 8; i++){
mask = (factor & 0x8000) ? 0x1100B : 0;
factor = (factor << 1) ^ mask;
mask = (id & (1 << i)) ? 0xFFFF : 0;
sum ^= mask & factor;
}
mtab[id] = sum;
@@ -14,6 +18,32 @@ void calc_table(__local uint *mtab, int id, int factor)
mtab[id + 256] = sum;
}
void calc_table2(__local uint *mtab, int id, int factor, int factor2)
{
int i, sum, sum2, mask;
mask = (id & 1) ? 0xFFFF : 0;
sum = mask & factor;
sum2 = mask & factor2;
for (i = 1; i < 8; i++){
mask = (factor & 0x8000) ? 0x1100B : 0;
factor = (factor << 1) ^ mask;
mask = (factor2 & 0x8000) ? 0x1100B : 0;
factor2 = (factor2 << 1) ^ mask;
mask = (id & (1 << i)) ? 0xFFFF : 0;
sum ^= mask & factor;
sum2 ^= mask & factor2;
}
mtab[id] = sum | (sum2 << 16);
sum = (sum << 4) ^ (((sum << 16) >> 31) & 0x88058) ^ (((sum << 17) >> 31) & 0x4402C) ^ (((sum << 18) >> 31) & 0x22016) ^ (((sum << 19) >> 31) & 0x1100B);
sum = (sum << 4) ^ (((sum << 16) >> 31) & 0x88058) ^ (((sum << 17) >> 31) & 0x4402C) ^ (((sum << 18) >> 31) & 0x22016) ^ (((sum << 19) >> 31) & 0x1100B);
sum2 = (sum2 << 4) ^ (((sum2 << 16) >> 31) & 0x88058) ^ (((sum2 << 17) >> 31) & 0x4402C) ^ (((sum2 << 18) >> 31) & 0x22016) ^ (((sum2 << 19) >> 31) & 0x1100B);
sum2 = (sum2 << 4) ^ (((sum2 << 16) >> 31) & 0x88058) ^ (((sum2 << 17) >> 31) & 0x4402C) ^ (((sum2 << 18) >> 31) & 0x22016) ^ (((sum2 << 19) >> 31) & 0x1100B);
mtab[id + 256] = sum | (sum2 << 16);
}
__kernel void method1(
__global uint *src,
__global uint *dst,
@@ -31,6 +61,7 @@ __kernel void method1(
dst[i] = 0;
for (blk = 0; blk < blk_num; blk++){
barrier(CLK_LOCAL_MEM_FENCE);
calc_table(mtab, table_id, factors[blk]);
barrier(CLK_LOCAL_MEM_FENCE);
@@ -42,7 +73,6 @@ __kernel void method1(
dst[i] ^= sum;
}
src += BLK_SIZE;
barrier(CLK_LOCAL_MEM_FENCE);
}
}
@@ -65,6 +95,7 @@ __kernel void method2(
}
for (blk = 0; blk < blk_num; blk++){
barrier(CLK_LOCAL_MEM_FENCE);
calc_table(mtab, table_id, factors[blk]);
barrier(CLK_LOCAL_MEM_FENCE);
@@ -82,11 +113,10 @@ __kernel void method2(
dst[pos + 4] ^= ((sum1 & 0xFF00FF00) >> 8) | (sum2 & 0xFF00FF00);
}
src += BLK_SIZE;
barrier(CLK_LOCAL_MEM_FENCE);
}
}
__kernel void method3(
__kernel void method4(
__global uint4 *src,
__global uint4 *dst,
__global ushort *factors,
@@ -106,6 +136,7 @@ __kernel void method3(
}
for (blk = 0; blk < blk_num; blk++){
barrier(CLK_LOCAL_MEM_FENCE);
calc_table(mtab, table_id, factors[blk]);
barrier(CLK_LOCAL_MEM_FENCE);
@@ -124,11 +155,140 @@ __kernel void method3(
dst[i + 1] ^= as_uint4((uchar16)(r0.y, r0.w, r1.y, r1.w, r2.y, r2.w, r3.y, r3.w, r4.y, r4.w, r5.y, r5.w, r6.y, r6.w, r7.y, r7.w));
}
src += BLK_SIZE / 4;
barrier(CLK_LOCAL_MEM_FENCE);
}
}
__kernel void method4(
__kernel void method9(
__global uint *src,
__global uint *dst,
__global ushort *factors,
int blk_num)
{
__local uint mtab[512];
int i, blk;
uint v, sum, sum2;
const int work_id = get_global_id(0);
const int work_size = get_global_size(0);
const int table_id = get_local_id(0);
for (i = work_id; i < BLK_SIZE; i += work_size){
dst[i] = 0;
dst[i + BLK_SIZE] = 0;
}
for (blk = 0; blk < blk_num; blk++){
barrier(CLK_LOCAL_MEM_FENCE);
calc_table2(mtab, table_id, factors[blk], factors[blk_num + blk]);
barrier(CLK_LOCAL_MEM_FENCE);
for (i = work_id; i < BLK_SIZE; i += work_size){
v = src[i];
sum = mtab[(uchar)v] ^ mtab[256 + (uchar)(v >> 8)];
sum2 = mtab[(uchar)(v >> 16)] ^ mtab[256 + (v >> 24)];
dst[i] ^= (sum & 0xFFFF) | (sum2 << 16);
dst[i + BLK_SIZE] ^= (sum >> 16) | (sum2 & 0xFFFF0000);
}
src += BLK_SIZE;
}
}
__kernel void method10(
__global uint *src,
__global uint *dst,
__global ushort *factors,
int blk_num)
{
__local uint mtab[512];
int i, blk, pos;
uint lo, hi, sum1, sum2, sum3, sum4;
const int work_id = get_global_id(0) * 2;
const int work_size = get_global_size(0) * 2;
const int table_id = get_local_id(0);
for (i = work_id; i < BLK_SIZE; i += work_size){
dst[i ] = 0;
dst[i + 1] = 0;
dst[i + BLK_SIZE ] = 0;
dst[i + BLK_SIZE + 1] = 0;
}
for (blk = 0; blk < blk_num; blk++){
barrier(CLK_LOCAL_MEM_FENCE);
calc_table2(mtab, table_id, factors[blk], factors[blk_num + blk]);
barrier(CLK_LOCAL_MEM_FENCE);
for (i = work_id; i < BLK_SIZE; i += work_size){
pos = (i & ~7) + ((i & 7) >> 1);
lo = src[pos ];
hi = src[pos + 4];
sum1 = mtab[(uchar)lo] ^ mtab[256 + (uchar)hi];
sum2 = mtab[(uchar)(lo >> 8)] ^ mtab[256 + (uchar)(hi >> 8)];
sum3 = mtab[(uchar)(lo >> 16)] ^ mtab[256 + (uchar)(hi >> 16)];
sum4 = mtab[lo >> 24] ^ mtab[256 + (hi >> 24)];
dst[pos ] ^= (sum1 & 0xFF) | ((sum2 & 0xFF) << 8) | ((sum3 & 0xFF) << 16) | (sum4 << 24);
dst[pos + 4] ^= ((sum1 >> 8) & 0xFF) | (sum2 & 0xFF00) | ((sum3 & 0xFF00) << 8) | ((sum4 & 0xFF00) << 16);
dst[pos + BLK_SIZE ] ^= ((sum1 >> 16) & 0xFF) | ((sum2 >> 8) & 0xFF00) | (sum3 & 0xFF0000) | ((sum4 & 0xFF0000) << 8);
dst[pos + BLK_SIZE + 4] ^= (sum1 >> 24) | ((sum2 >> 16) & 0xFF00) | ((sum3 >> 8) & 0xFF0000) | (sum4 & 0xFF000000);
}
src += BLK_SIZE;
}
}
__kernel void method12(
__global uint4 *src,
__global uint4 *dst,
__global ushort *factors,
int blk_num)
{
__local uint mtab[512];
int i, blk;
uchar4 r0, r1, r2, r3, r4, r5, r6, r7, r8, r9, rA, rB, rC, rD, rE, rF;
uchar16 lo, hi;
const int work_id = get_global_id(0) * 2;
const int work_size = get_global_size(0) * 2;
const int table_id = get_local_id(0);
for (i = work_id; i < BLK_SIZE / 4; i += work_size){
dst[i ] = 0;
dst[i + 1] = 0;
dst[i + BLK_SIZE / 4 ] = 0;
dst[i + BLK_SIZE / 4 + 1] = 0;
}
for (blk = 0; blk < blk_num; blk++){
barrier(CLK_LOCAL_MEM_FENCE);
calc_table2(mtab, table_id, factors[blk], factors[blk_num + blk]);
barrier(CLK_LOCAL_MEM_FENCE);
for (i = work_id; i < BLK_SIZE / 4; i += work_size){
lo = as_uchar16(src[i ]);
hi = as_uchar16(src[i + 1]);
r0 = as_uchar4(mtab[lo.s0] ^ mtab[256 + hi.s0]);
r1 = as_uchar4(mtab[lo.s1] ^ mtab[256 + hi.s1]);
r2 = as_uchar4(mtab[lo.s2] ^ mtab[256 + hi.s2]);
r3 = as_uchar4(mtab[lo.s3] ^ mtab[256 + hi.s3]);
r4 = as_uchar4(mtab[lo.s4] ^ mtab[256 + hi.s4]);
r5 = as_uchar4(mtab[lo.s5] ^ mtab[256 + hi.s5]);
r6 = as_uchar4(mtab[lo.s6] ^ mtab[256 + hi.s6]);
r7 = as_uchar4(mtab[lo.s7] ^ mtab[256 + hi.s7]);
r8 = as_uchar4(mtab[lo.s8] ^ mtab[256 + hi.s8]);
r9 = as_uchar4(mtab[lo.s9] ^ mtab[256 + hi.s9]);
rA = as_uchar4(mtab[lo.sa] ^ mtab[256 + hi.sa]);
rB = as_uchar4(mtab[lo.sb] ^ mtab[256 + hi.sb]);
rC = as_uchar4(mtab[lo.sc] ^ mtab[256 + hi.sc]);
rD = as_uchar4(mtab[lo.sd] ^ mtab[256 + hi.sd]);
rE = as_uchar4(mtab[lo.se] ^ mtab[256 + hi.se]);
rF = as_uchar4(mtab[lo.sf] ^ mtab[256 + hi.sf]);
dst[i ] ^= as_uint4((uchar16)(r0.x, r1.x, r2.x, r3.x, r4.x, r5.x, r6.x, r7.x, r8.x, r9.x, rA.x, rB.x, rC.x, rD.x, rE.x, rF.x));
dst[i + 1] ^= as_uint4((uchar16)(r0.y, r1.y, r2.y, r3.y, r4.y, r5.y, r6.y, r7.y, r8.y, r9.y, rA.y, rB.y, rC.y, rD.y, rE.y, rF.y));
dst[i + BLK_SIZE / 4 ] ^= as_uint4((uchar16)(r0.z, r1.z, r2.z, r3.z, r4.z, r5.z, r6.z, r7.z, r8.z, r9.z, rA.z, rB.z, rC.z, rD.z, rE.z, rF.z));
dst[i + BLK_SIZE / 4 + 1] ^= as_uint4((uchar16)(r0.w, r1.w, r2.w, r3.w, r4.w, r5.w, r6.w, r7.w, r8.w, r9.w, rA.w, rB.w, rC.w, rD.w, rE.w, rF.w));
}
src += BLK_SIZE / 4;
}
}
__kernel void method16(
__global uint *src,
__global uint *dst,
__global ushort *factors,
@@ -136,7 +296,7 @@ __kernel void method4(
{
__local int table[16];
__local uint cache[256];
int i, j, blk, pos, sht, mask;
int i, j, blk, pos, mask, tmp;
uint sum;
const int work_id = get_global_id(0);
const int work_size = get_global_size(0);
@@ -146,11 +306,12 @@ __kernel void method4(
for (blk = 0; blk < blk_num; blk++){
if (get_local_id(0) == 0){
pos = factors[blk] << 16;
table[0] = pos;
tmp = factors[blk];
table[0] = tmp;
for (j = 1; j < 16; j++){
pos = (pos << 1) ^ ((pos >> 31) & 0x100B0000);
table[j] = pos;
mask = (tmp & 0x8000) ? 0x1100B : 0;
tmp = (tmp << 1) ^ mask;
table[j] = tmp;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
@@ -161,10 +322,11 @@ __kernel void method4(
barrier(CLK_LOCAL_MEM_FENCE);
sum = 0;
sht = (i & 60) >> 2;
tmp = (i & 60) >> 2;
tmp = 0x8000 >> tmp;
pos &= ~60;
for (j = 15; j >= 0; j--){
mask = (table[j] << sht) >> 31;
mask = (table[j] & tmp) ? 0xFFFFFFFF : 0;
sum ^= mask & cache[pos];
pos += 4;
}