Skip to content

O
opencv
Commit ce5bae1c authored by Alexander Alekhin's avatar Alexander Alekhin Committed by OpenCV Buildbot
Browse files
Options

Merge pull request #2871 from vbystricky:oclopt_integral

parents 1db9cc3f 606df046
Hide whitespace changes
Inline Side-by-side
Showing with 94 additions and 218 deletions
modules/imgproc/src/opencl/integral_sum.cl
View file @ ce5bae1c
...@@ -61,115 +61,64 @@ ...@@ -61,115 +61,64 @@
#define GET_CONFLICT_OFFSET(lid) ((lid) >> LOG_NUM_BANKS) #define GET_CONFLICT_OFFSET(lid) ((lid) >> LOG_NUM_BANKS)
#if sdepth == 4 #if sdepth == 4
#define sumT int
#define vecSumT int4
#define convertToSum4 convert_int4
#elif sdepth == 5
#define sumT float
#define vecSumT float4
#define convertToSum4 convert_float4
#endif
kernel void integral_sum_cols(__global uchar4 *src, __global int *sum, kernel void integral_sum_cols(__global const uchar4 *src, __global uchar *sum_ptr,
int src_offset, int pre_invalid, int rows, int cols, int src_step, int dst_step) int src_offset, int rows, int cols, int src_step, int dst_step)
{ {
__global sumT *sum = (__global sumT *)sum_ptr;
int lid = get_local_id(0); int lid = get_local_id(0);
int gid = get_group_id(0); int gid = get_group_id(0);
int4 src_t[2], sum_t[2]; vecSumT src_t[2], sum_t[2];
__local int4 lm_sum[2][LSIZE + LOG_LSIZE]; __local vecSumT lm_sum[2][LSIZE + LOG_LSIZE];
__local int* sum_p; __local sumT* sum_p;
src_step = src_step >> 2; src_step = src_step >> 2;
gid = gid << 1; gid = gid << 1;
for(int i = 0; i < rows; i =i + LSIZE_1) int lid_prim = ((lid & 127) << 1) + 1;
for (int i = 0; i < rows; i += LSIZE_1)
{ {
src_t[0] = (i + lid < rows ? convert_int4(src[src_offset + (lid+i) * src_step + gid]) : 0); if (i + lid < rows)
src_t[1] = (i + lid < rows ? convert_int4(src[src_offset + (lid+i) * src_step + gid + 1]) : 0);
sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]);
sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]);
barrier(CLK_LOCAL_MEM_FENCE);
int bf_loc = lid + GET_CONFLICT_OFFSET(lid);
lm_sum[0][bf_loc] = src_t[0];
lm_sum[1][bf_loc] = src_t[1];
int offset = 1;
for(int d = LSIZE >> 1 ; d > 0; d>>=1)
{ {
barrier(CLK_LOCAL_MEM_FENCE); int src_index = mad24((lid+i), src_step, gid + src_offset);
int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; src_t[0] = convertToSum4(src[src_index]);
ai += GET_CONFLICT_OFFSET(ai); src_t[1] = convertToSum4(src[src_index + 1]);
bi += GET_CONFLICT_OFFSET(bi);
if((lid & 127) < d)
{
lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
}
offset <<= 1;
} }
barrier(CLK_LOCAL_MEM_FENCE); else
if(lid < 2)
{ {
lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; src_t[0] = (vecSumT)0;
src_t[1] = (vecSumT)0;
} }
for(int d = 1; d < LSIZE; d <<= 1)
{
barrier(CLK_LOCAL_MEM_FENCE);
offset >>= 1;
int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset;
ai += GET_CONFLICT_OFFSET(ai);
bi += GET_CONFLICT_OFFSET(bi);
if((lid & 127) < d) if (i == 0)
{ {
lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; sum_t[0] = (vecSumT)0;
lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai]; sum_t[1] = (vecSumT)0;
}
} }
barrier(CLK_LOCAL_MEM_FENCE); else
if(lid > 0 && (i+lid) <= rows)
{ {
int loc_s0 = gid * dst_step + i + lid - 1 - pre_invalid * dst_step / 4, loc_s1 = loc_s0 + dst_step ; sum_t[0] = lm_sum[0][LSIZE_2 + LOG_LSIZE];
lm_sum[0][bf_loc] += sum_t[0]; sum_t[1] = lm_sum[1][LSIZE_2 + LOG_LSIZE];
lm_sum[1][bf_loc] += sum_t[1];
sum_p = (__local int*)(&(lm_sum[0][bf_loc]));
for(int k = 0; k < 4; k++)
{
if(gid * 4 + k >= cols + pre_invalid || gid * 4 + k < pre_invalid) continue;
sum[loc_s0 + k * dst_step / 4] = sum_p[k];
}
sum_p = (__local int*)(&(lm_sum[1][bf_loc]));
for(int k = 0; k < 4; k++)
{
if(gid * 4 + k + 4 >= cols + pre_invalid) break;
sum[loc_s1 + k * dst_step / 4] = sum_p[k];
}
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
}
}
kernel void integral_sum_rows(__global int4 *srcsum, __global int *sum,
int rows, int cols, int src_step, int sum_step, int sum_offset)
{
int lid = get_local_id(0);
int gid = get_group_id(0);
int4 src_t[2], sum_t[2];
__local int4 lm_sum[2][LSIZE + LOG_LSIZE];
__local int *sum_p;
src_step = src_step >> 4;
for(int i = 0; i < rows; i =i + LSIZE_1)
{
src_t[0] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2] : 0;
src_t[1] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2 + 1] : 0;
sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]);
sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]);
barrier(CLK_LOCAL_MEM_FENCE);
int bf_loc = lid + GET_CONFLICT_OFFSET(lid); int bf_loc = lid + GET_CONFLICT_OFFSET(lid);
lm_sum[0][bf_loc] = src_t[0];
lm_sum[0][bf_loc] = src_t[0];
lm_sum[1][bf_loc] = src_t[1]; lm_sum[1][bf_loc] = src_t[1];
int offset = 1; int offset = 1;
for(int d = LSIZE >> 1 ; d > 0; d>>=1) for (int d = LSIZE >> 1 ; d > 0; d>>=1)
{ {
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; int ai = offset * lid_prim - 1,bi = ai + offset;
ai += GET_CONFLICT_OFFSET(ai); ai += GET_CONFLICT_OFFSET(ai);
bi += GET_CONFLICT_OFFSET(bi); bi += GET_CONFLICT_OFFSET(bi);
...@@ -180,15 +129,15 @@ kernel void integral_sum_rows(__global int4 *srcsum, __global int *sum, ...@@ -180,15 +129,15 @@ kernel void integral_sum_rows(__global int4 *srcsum, __global int *sum,
offset <<= 1; offset <<= 1;
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
if(lid < 2) if (lid < 2)
{ {
lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
} }
for(int d = 1; d < LSIZE; d <<= 1) for (int d = 1; d < LSIZE; d <<= 1)
{ {
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
offset >>= 1; offset >>= 1;
int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; int ai = offset * lid_prim - 1,bi = ai + offset;
ai += GET_CONFLICT_OFFSET(ai); ai += GET_CONFLICT_OFFSET(ai);
bi += GET_CONFLICT_OFFSET(bi); bi += GET_CONFLICT_OFFSET(bi);
...@@ -199,152 +148,78 @@ kernel void integral_sum_rows(__global int4 *srcsum, __global int *sum, ...@@ -199,152 +148,78 @@ kernel void integral_sum_rows(__global int4 *srcsum, __global int *sum,
} }
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
if(gid == 0 && (i + lid) <= rows) if (lid > 0 && (i+lid) <= rows)
{
sum[sum_offset + i + lid] = 0;
}
if(i + lid == 0)
{
int loc0 = gid * 2 * sum_step;
for(int k = 1; k <= 8; k++)
{
if(gid * 8 + k > cols) break;
sum[sum_offset + loc0 + k * sum_step / 4] = 0;
}
}
if(lid > 0 && (i+lid) <= rows)
{ {
int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ; int loc_s0 = mad24(gid, dst_step, i + lid - 1), loc_s1 = loc_s0 + dst_step;
lm_sum[0][bf_loc] += sum_t[0]; lm_sum[0][bf_loc] += sum_t[0];
lm_sum[1][bf_loc] += sum_t[1]; lm_sum[1][bf_loc] += sum_t[1];
sum_p = (__local int*)(&(lm_sum[0][bf_loc])); sum_p = (__local sumT*)(&(lm_sum[0][bf_loc]));
for(int k = 0; k < 4; k++) for (int k = 0; k < 4; k++)
{ {
if(gid * 8 + k >= cols) break; if (gid * 4 + k >= cols)
sum[loc_s0 + k * sum_step / 4] = sum_p[k]; break;
sum[loc_s0 + k * dst_step / 4] = sum_p[k];
} }
sum_p = (__local int*)(&(lm_sum[1][bf_loc])); sum_p = (__local sumT*)(&(lm_sum[1][bf_loc]));
for(int k = 0; k < 4; k++) for (int k = 0; k < 4; k++)
{ {
if(gid * 8 + 4 + k >= cols) break; if (gid * 4 + k + 4 >= cols)
sum[loc_s1 + k * sum_step / 4] = sum_p[k]; break;
sum[loc_s1 + k * dst_step / 4] = sum_p[k];
} }
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
} }
#elif sdepth == 5
kernel void integral_sum_cols(__global uchar4 *src, __global float *sum, kernel void integral_sum_rows(__global const uchar *srcsum_ptr, __global uchar *sum_ptr,
int src_offset, int pre_invalid, int rows, int cols, int src_step, int dst_step) int rows, int cols, int src_step, int sum_step, int sum_offset)
{ {
__global const vecSumT *srcsum = (__global const vecSumT *)srcsum_ptr;
__global sumT *sum = (__global sumT *)sum_ptr;
int lid = get_local_id(0); int lid = get_local_id(0);
int gid = get_group_id(0); int gid = get_group_id(0);
float4 src_t[2], sum_t[2]; vecSumT src_t[2], sum_t[2];
__local float4 lm_sum[2][LSIZE + LOG_LSIZE]; __local vecSumT lm_sum[2][LSIZE + LOG_LSIZE];
__local float* sum_p; __local sumT *sum_p;
src_step = src_step >> 2; src_step = src_step >> 4;
gid = gid << 1; int lid_prim = ((lid & 127) << 1) + 1;
for(int i = 0; i < rows; i =i + LSIZE_1) for (int i = 0; i < rows; i += LSIZE_1)
{ {
src_t[0] = (i + lid < rows ? convert_float4(src[src_offset + (lid+i) * src_step + gid]) : (float4)0); if (i + lid < rows)
src_t[1] = (i + lid < rows ? convert_float4(src[src_offset + (lid+i) * src_step + gid + 1]) : (float4)0);
sum_t[0] = (i == 0 ? (float4)0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]);
sum_t[1] = (i == 0 ? (float4)0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]);
barrier(CLK_LOCAL_MEM_FENCE);
int bf_loc = lid + GET_CONFLICT_OFFSET(lid);
lm_sum[0][bf_loc] = src_t[0];
lm_sum[1][bf_loc] = src_t[1];
int offset = 1;
for(int d = LSIZE >> 1 ; d > 0; d>>=1)
{ {
barrier(CLK_LOCAL_MEM_FENCE); int sum_idx = mad24(lid + i, src_step, gid * 2);
int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; src_t[0] = srcsum[sum_idx];
ai += GET_CONFLICT_OFFSET(ai); src_t[1] = srcsum[sum_idx + 1];
bi += GET_CONFLICT_OFFSET(bi);
if((lid & 127) < d)
{
lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
}
offset <<= 1;
} }
barrier(CLK_LOCAL_MEM_FENCE); else
if(lid < 2)
{ {
lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; src_t[0] = 0;
src_t[1] = 0;
} }
for(int d = 1; d < LSIZE; d <<= 1) if (i == 0)
{ {
barrier(CLK_LOCAL_MEM_FENCE); sum_t[0] = 0;
offset >>= 1; sum_t[1] = 0;
int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset;
ai += GET_CONFLICT_OFFSET(ai);
bi += GET_CONFLICT_OFFSET(bi);
if((lid & 127) < d)
{
lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai];
}
} }
barrier(CLK_LOCAL_MEM_FENCE); else
if(lid > 0 && (i+lid) <= rows)
{ {
int loc_s0 = gid * dst_step + i + lid - 1 - pre_invalid * dst_step / 4, loc_s1 = loc_s0 + dst_step ; sum_t[0] = lm_sum[0][LSIZE_2 + LOG_LSIZE];
lm_sum[0][bf_loc] += sum_t[0]; sum_t[1] = lm_sum[1][LSIZE_2 + LOG_LSIZE];
lm_sum[1][bf_loc] += sum_t[1];
sum_p = (__local float*)(&(lm_sum[0][bf_loc]));
for(int k = 0; k < 4; k++)
{
if(gid * 4 + k >= cols + pre_invalid || gid * 4 + k < pre_invalid) continue;
sum[loc_s0 + k * dst_step / 4] = sum_p[k];
}
sum_p = (__local float*)(&(lm_sum[1][bf_loc]));
for(int k = 0; k < 4; k++)
{
if(gid * 4 + k + 4 >= cols + pre_invalid) break;
sum[loc_s1 + k * dst_step / 4] = sum_p[k];
}
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
}
}
kernel void integral_sum_rows(__global float4 *srcsum, __global float *sum,
int rows, int cols, int src_step, int sum_step, int sum_offset)
{
int lid = get_local_id(0);
int gid = get_group_id(0);
float4 src_t[2], sum_t[2];
__local float4 lm_sum[2][LSIZE + LOG_LSIZE];
__local float *sum_p;
src_step = src_step >> 4;
for(int i = 0; i < rows; i =i + LSIZE_1)
{
src_t[0] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2] : (float4)0;
src_t[1] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2 + 1] : (float4)0;
sum_t[0] = (i == 0 ? (float4)0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]);
sum_t[1] = (i == 0 ? (float4)0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]);
barrier(CLK_LOCAL_MEM_FENCE);
int bf_loc = lid + GET_CONFLICT_OFFSET(lid); int bf_loc = lid + GET_CONFLICT_OFFSET(lid);
lm_sum[0][bf_loc] = src_t[0];
lm_sum[0][bf_loc] = src_t[0];
lm_sum[1][bf_loc] = src_t[1]; lm_sum[1][bf_loc] = src_t[1];
int offset = 1; int offset = 1;
for(int d = LSIZE >> 1 ; d > 0; d>>=1) for (int d = LSIZE >> 1 ; d > 0; d>>=1)
{ {
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; int ai = offset * lid_prim - 1, bi = ai + offset;
ai += GET_CONFLICT_OFFSET(ai); ai += GET_CONFLICT_OFFSET(ai);
bi += GET_CONFLICT_OFFSET(bi); bi += GET_CONFLICT_OFFSET(bi);
...@@ -355,59 +230,60 @@ kernel void integral_sum_rows(__global float4 *srcsum, __global float *sum, ...@@ -355,59 +230,60 @@ kernel void integral_sum_rows(__global float4 *srcsum, __global float *sum,
offset <<= 1; offset <<= 1;
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
if(lid < 2) if (lid < 2)
{ {
lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
} }
for(int d = 1; d < LSIZE; d <<= 1) for (int d = 1; d < LSIZE; d <<= 1)
{ {
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
offset >>= 1; offset >>= 1;
int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; int ai = offset * lid_prim - 1,bi = ai + offset;
ai += GET_CONFLICT_OFFSET(ai); ai += GET_CONFLICT_OFFSET(ai);
bi += GET_CONFLICT_OFFSET(bi); bi += GET_CONFLICT_OFFSET(bi);
if((lid & 127) < d) if ((lid & 127) < d)
{ {
lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai];
lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai]; lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai];
} }
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
if(gid == 0 && (i + lid) <= rows) if (gid == 0 && (i + lid) <= rows)
{ {
sum[sum_offset + i + lid] = 0; sum[sum_offset + i + lid] = 0;
} }
if(i + lid == 0) if (i + lid == 0)
{ {
int loc0 = gid * 2 * sum_step; int loc0 = gid * 2 * sum_step;
for(int k = 1; k <= 8; k++) for(int k = 1; k <= 8; k++)
{ {
if(gid * 8 + k > cols) break; if (gid * 8 + k > cols)
break;
sum[sum_offset + loc0 + k * sum_step / 4] = 0; sum[sum_offset + loc0 + k * sum_step / 4] = 0;
} }
} }
if(lid > 0 && (i+lid) <= rows) if (lid > 0 && (i+lid) <= rows)
{ {
int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ; int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ;
lm_sum[0][bf_loc] += sum_t[0]; lm_sum[0][bf_loc] += sum_t[0];
lm_sum[1][bf_loc] += sum_t[1]; lm_sum[1][bf_loc] += sum_t[1];
sum_p = (__local float*)(&(lm_sum[0][bf_loc])); sum_p = (__local sumT*)(&(lm_sum[0][bf_loc]));
for(int k = 0; k < 4; k++) for(int k = 0; k < 4; k++)
{ {
if(gid * 8 + k >= cols) break; if (gid * 8 + k >= cols)
break;
sum[loc_s0 + k * sum_step / 4] = sum_p[k]; sum[loc_s0 + k * sum_step / 4] = sum_p[k];
} }
sum_p = (__local float*)(&(lm_sum[1][bf_loc])); sum_p = (__local sumT*)(&(lm_sum[1][bf_loc]));
for(int k = 0; k < 4; k++) for(int k = 0; k < 4; k++)
{ {
if(gid * 8 + 4 + k >= cols) break; if (gid * 8 + 4 + k >= cols)
break;
sum[loc_s1 + k * sum_step / 4] = sum_p[k]; sum[loc_s1 + k * sum_step / 4] = sum_p[k];
} }
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
} }
#endif
modules/imgproc/src/sumpixels.cpp
View file @ ce5bae1c
...@@ -254,19 +254,19 @@ static bool ocl_integral( InputArray _src, OutputArray _sum, int sdepth ) ...@@ -254,19 +254,19 @@ static bool ocl_integral( InputArray _src, OutputArray _sum, int sdepth )
UMat src = _src.getUMat(), t_sum(t_size, sdepth), sum = _sum.getUMat(); UMat src = _src.getUMat(), t_sum(t_size, sdepth), sum = _sum.getUMat();
t_sum = t_sum(Range::all(), Range(0, size.height)); t_sum = t_sum(Range::all(), Range(0, size.height));
int offset = (int)src.offset / vlen, pre_invalid = (int)src.offset % vlen; int offset = (int)src.offset / vlen;
int vcols = (pre_invalid + src.cols + vlen - 1) / vlen; int vcols = (src.cols + vlen - 1) / vlen;
int sum_offset = (int)sum.offset / vlen; int sum_offset = (int)sum.offset / vlen;
k1.args(ocl::KernelArg::PtrReadOnly(src), ocl::KernelArg::PtrWriteOnly(t_sum), k1.args(ocl::KernelArg::PtrReadOnly(src), ocl::KernelArg::PtrWriteOnly(t_sum),
offset, pre_invalid, src.rows, src.cols, (int)src.step, (int)t_sum.step); offset, src.rows, src.cols, (int)src.step, (int)t_sum.step);
size_t gt = ((vcols + 1) / 2) * 256, lt = 256; size_t gt = ((vcols + 1) / 2) * 256, lt = 256;
if (!k1.run(1, &gt, &lt, false)) if (!k1.run(1, &gt, &lt, false))
return false; return false;
ocl::Kernel k2("integral_sum_rows", ocl::imgproc::integral_sum_oclsrc, ocl::Kernel k2("integral_sum_rows", ocl::imgproc::integral_sum_oclsrc,
format("-D sdepth=%d", sdepth)); format("-D sdepth=%d", sdepth));
k2.args(ocl::KernelArg::PtrReadWrite(t_sum), ocl::KernelArg::PtrWriteOnly(sum), k2.args(ocl::KernelArg::PtrReadOnly(t_sum), ocl::KernelArg::PtrWriteOnly(sum),
t_sum.rows, t_sum.cols, (int)t_sum.step, (int)sum.step, sum_offset); t_sum.rows, t_sum.cols, (int)t_sum.step, (int)sum.step, sum_offset);
size_t gt2 = t_sum.cols * 32, lt2 = 256; size_t gt2 = t_sum.cols * 32, lt2 = 256;
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment