Danylo Lytovchenko
403 linhas
15 KiB
C
403 linhas
15 KiB
C
/* Copyright (c) 2010 - 2021 Advanced Micro Devices, Inc.
|
|
|
|
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. */
|
|
|
|
static const char* local_atomics_histogram =
|
|
"#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
|
|
"#define MIN(a,b) ((a) < (b)) ? (a) : (b) \n"
|
|
"#define MAX(a,b) ((a) > (b)) ? (a) : (b) \n"
|
|
"__kernel __attribute__((reqd_work_group_size(256,1,1)))\n"
|
|
"void local_atomics_histogram(__global uint4 *Image,\n"
|
|
"__global uint *Histogram,\n"
|
|
"uint n4VectorsPerThread)\n"
|
|
"{\n"
|
|
" __local __attribute__((aligned(16))) uint subhists[NBANKS * NBINS];\n"
|
|
"\n"
|
|
" uint tid = get_global_id(0);\n"
|
|
" uint ltid = get_local_id(0);\n"
|
|
" uint Stride = get_global_size(0);\n"
|
|
"\n"
|
|
" uint i, idx;\n"
|
|
" uint4 temp, temp2;\n"
|
|
" const uint shft = (uint) BITS_PER_PIX;\n"
|
|
" const uint msk = (uint) (NBINS-1);\n"
|
|
" uint offset = (uint) ltid % (uint) (NBANKS);\n"
|
|
"\n"
|
|
" uint lmem_items = NBANKS * NBINS;\n"
|
|
" uint lmem_items_per_thread;\n"
|
|
" uint lmem_max_threads;\n"
|
|
"\n"
|
|
" // parallel LDS clear\n"
|
|
" // first, calculate threads per item, at least 1:\n"
|
|
" lmem_max_threads = MIN( 1, get_local_size(0) / lmem_items );\n"
|
|
" // but no more than we have items:\n"
|
|
" lmem_max_threads = MAX( 1, lmem_max_threads / lmem_items );\n"
|
|
" // calculate threads total:\n"
|
|
" lmem_max_threads = lmem_items / lmem_max_threads;\n"
|
|
" // but no more than LDS banks:\n"
|
|
" lmem_max_threads = MIN( get_local_size(0), lmem_max_threads );\n"
|
|
"\n"
|
|
" lmem_items_per_thread = lmem_items / lmem_max_threads;\n"
|
|
"\n"
|
|
" // now, clear LDS\n"
|
|
" __local uint4 *p = (__local uint4 *) subhists;\n"
|
|
"\n"
|
|
" if( ltid < lmem_max_threads )\n"
|
|
" {\n"
|
|
" for(i=0, idx=ltid; i<lmem_items_per_thread/4; i++, "
|
|
"idx+=lmem_max_threads)\n"
|
|
" {\n"
|
|
" p[idx] = 0;\n"
|
|
" }\n"
|
|
" }\n"
|
|
"\n"
|
|
" barrier( CLK_LOCAL_MEM_FENCE );\n"
|
|
"\n"
|
|
" // read & scatter phase\n"
|
|
"\n"
|
|
" for( i=0, idx=tid; i<n4VectorsPerThread; i++, idx += Stride )\n"
|
|
" {\n"
|
|
" temp = Image[idx];\n"
|
|
" temp2 = (temp & msk) * (uint4) NBANKS + offset;\n"
|
|
"\n"
|
|
" (void) atom_inc( subhists + temp2.x );\n"
|
|
" (void) atom_inc( subhists + temp2.y );\n"
|
|
" (void) atom_inc( subhists + temp2.z );\n"
|
|
" (void) atom_inc( subhists + temp2.w );\n"
|
|
"\n"
|
|
" temp = temp >> shft;\n"
|
|
" temp2 = (temp & msk) * (uint4) NBANKS + offset;\n"
|
|
"\n"
|
|
" (void) atom_inc( subhists + temp2.x );\n"
|
|
" (void) atom_inc( subhists + temp2.y );\n"
|
|
" (void) atom_inc( subhists + temp2.z );\n"
|
|
" (void) atom_inc( subhists + temp2.w );\n"
|
|
"\n"
|
|
" temp = temp >> shft;\n"
|
|
" temp2 = (temp & msk) * (uint4) NBANKS + offset;\n"
|
|
"\n"
|
|
" (void) atom_inc( subhists + temp2.x );\n"
|
|
" (void) atom_inc( subhists + temp2.y );\n"
|
|
" (void) atom_inc( subhists + temp2.z );\n"
|
|
" (void) atom_inc( subhists + temp2.w );\n"
|
|
"\n"
|
|
" temp = temp >> shft;\n"
|
|
" temp2 = (temp & msk) * (uint4) NBANKS + offset;\n"
|
|
"\n"
|
|
" (void) atom_inc( subhists + temp2.x );\n"
|
|
" (void) atom_inc( subhists + temp2.y );\n"
|
|
" (void) atom_inc( subhists + temp2.z );\n"
|
|
" (void) atom_inc( subhists + temp2.w );\n"
|
|
" }\n"
|
|
"\n"
|
|
" barrier( CLK_LOCAL_MEM_FENCE );\n"
|
|
"\n"
|
|
" // reduce __local banks to single histogram per work-group\n"
|
|
"\n"
|
|
" if( ltid < NBINS )\n"
|
|
" {\n"
|
|
" uint bin = 0;\n"
|
|
" for( i=0; i<NBANKS; i++ )\n"
|
|
" {\n"
|
|
" bin += subhists[ (ltid * NBANKS) + i ];\n"
|
|
" }\n"
|
|
" Histogram[ (get_group_id(0) * NBINS) + ltid ] = bin;\n"
|
|
" }\n"
|
|
"}\n";
|
|
|
|
static const char* local_atomics_reduce =
|
|
" __kernel void local_atomics_reduce( __global uint *Histogram, uint "
|
|
"nSubHists )\n"
|
|
"{\n"
|
|
" uint tid = get_global_id(0);\n"
|
|
" uint bin = 0;\n"
|
|
" // Reduce work-group histograms into single histogram,\n"
|
|
" // one thread for each bin.\n"
|
|
" for( int i=0; i < nSubHists; i++ )\n"
|
|
" bin += Histogram[ (i * NBINS) + tid ];\n"
|
|
" Histogram[ tid ] = bin;\n"
|
|
"}\n";
|
|
|
|
static const char* global_atomics_histogram =
|
|
"#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\n"
|
|
"__kernel __attribute__((reqd_work_group_size(256,1,1)))\n"
|
|
"void global_atomics_histogram(uint ItemsPerThread,\n"
|
|
"__global uint *Input,\n"
|
|
"__global uint *Histogram)\n"
|
|
"{\n"
|
|
" uint tid = get_global_id(0);\n"
|
|
" const uint shft = (uint) BITS_PER_PIX;\n"
|
|
" const uint msk = (uint) (NBINS-1);\n"
|
|
" uint Stride = get_global_size(0);\n"
|
|
" for( int i = 0; i < ItemsPerThread; i++)\n"
|
|
" {\n"
|
|
" uint temp = Input[tid];\n"
|
|
" atom_inc( &(Histogram[ (temp & msk) ]) );\n"
|
|
" temp = temp >> shft;\n"
|
|
" atom_inc( &(Histogram[ (temp & msk) ]) );\n"
|
|
" temp = temp >> shft;\n"
|
|
" atom_inc( &(Histogram[ (temp & msk) ]) );\n"
|
|
" temp = temp >> shft;\n"
|
|
" atom_inc( &(Histogram[ (temp & msk) ]) );\n"
|
|
" tid += Stride;"
|
|
" }\n"
|
|
"}\n";
|
|
|
|
static const char* global_vec4_atomics_histogram =
|
|
"#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\n"
|
|
"__kernel __attribute__((reqd_work_group_size(256,1,1)))\n"
|
|
"void global_atomics_histogram(uint ItemsPerThread,\n"
|
|
"__global uint4 *Input,\n"
|
|
"__global uint *Histogram)\n"
|
|
"{\n"
|
|
" uint tid = get_global_id(0);\n"
|
|
" const uint shft = (uint) BITS_PER_PIX;\n"
|
|
" const uint msk = (uint) (NBINS-1);\n"
|
|
" uint Stride = get_global_size(0);\n"
|
|
" for( int i = 0; i < ItemsPerThread; i++)\n"
|
|
" {\n"
|
|
" uint4 temp = Input[tid];\n"
|
|
" atom_inc( &(Histogram[ (temp.x & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.y & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.z & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.w & msk) ]) );\n"
|
|
" temp = temp >> shft;\n"
|
|
" atom_inc( &(Histogram[ (temp.x & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.y & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.z & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.w & msk) ]) );\n"
|
|
" temp = temp >> shft;\n"
|
|
" atom_inc( &(Histogram[ (temp.x & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.y & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.z & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.w & msk) ]) );\n"
|
|
" temp = temp >> shft;\n"
|
|
" atom_inc( &(Histogram[ (temp.x & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.y & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.z & msk) ]) );\n"
|
|
" atom_inc( &(Histogram[ (temp.w & msk) ]) );\n"
|
|
" tid += Stride;"
|
|
" }\n"
|
|
"}\n";
|
|
|
|
static const char* global_atomics_sum_reduction_all_to_zero =
|
|
"#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\n"
|
|
" __kernel void global_atomics_sum_reduction_all_to_zero(uint "
|
|
"ItemsPerThread, __global uint *Input, __global int *Output )\n"
|
|
"{\n"
|
|
" uint sum = 0;\n"
|
|
" const uint msk = (uint)3;\n"
|
|
" const uint shft = (uint)8;\n"
|
|
" \n"
|
|
" uint tid = get_global_id(0);\n"
|
|
" uint Stride = get_global_size(0);\n"
|
|
" for( int i = 0; i < ItemsPerThread; i++)\n"
|
|
" {\n"
|
|
" uint data = Input[tid];\n"
|
|
" sum += data & msk;\n"
|
|
" data = data >> shft;"
|
|
" sum += data & msk;\n"
|
|
" data = data >> shft;"
|
|
" sum += data & msk;\n"
|
|
" data = data >> shft;"
|
|
" sum += data & msk;\n"
|
|
" tid += Stride;\n"
|
|
" }\n"
|
|
" atom_add( &(Output[0]), sum);\n"
|
|
"}\n";
|
|
|
|
static const char* global_atomics_sum_reduction_workgroup =
|
|
"#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\n"
|
|
" __kernel void global_atomics_sum_reduction_workgroup(uint "
|
|
"ItemsPerThread, __global uint *Input, __global int *Output )\n"
|
|
"{\n"
|
|
" uint sum = 0;\n"
|
|
" const uint msk = (uint)3;\n"
|
|
" const uint shft = (uint)8;\n"
|
|
" \n"
|
|
" uint tid = get_global_id(0);\n"
|
|
" uint Stride = get_global_size(0);\n"
|
|
" for( int i = 0; i < ItemsPerThread; i++)\n"
|
|
" {\n"
|
|
" uint data = Input[tid];\n"
|
|
" sum += data & msk;\n"
|
|
" data = data >> shft;"
|
|
" sum += data & msk;\n"
|
|
" data = data >> shft;"
|
|
" sum += data & msk;\n"
|
|
" data = data >> shft;"
|
|
" sum += data & msk;\n"
|
|
" tid += Stride;\n"
|
|
" }\n"
|
|
" atom_add( &(Output[get_group_id(0)]), sum);\n"
|
|
"}\n";
|
|
|
|
static const char* local_reduction =
|
|
"__kernel void local_reduction(__global uint* input, __global uint* "
|
|
"output, __local uint* sdata)\n"
|
|
"{\n"
|
|
" // load shared mem\n"
|
|
" const uint msk = (uint)3;\n"
|
|
" const uint shft = (uint)8;\n"
|
|
" unsigned int tid = get_local_id(0);\n"
|
|
"\n"
|
|
" unsigned int localSize = get_local_size(0);\n"
|
|
" unsigned int stride = get_global_id(0) * 2;\n"
|
|
" unsigned int data1 = input[stride];\n"
|
|
" unsigned int data2 = input[stride + 1];\n"
|
|
" unsigned int sum = 0;\n"
|
|
" for( int i = 0; i < 4; i++)\n"
|
|
" {\n"
|
|
" sum += (data1 & msk) + (data2 & msk);\n"
|
|
" data1 = data1 >> shft;\n"
|
|
" data2 = data2 >> shft;\n"
|
|
" }\n"
|
|
" sdata[tid] = sum;"
|
|
"\n"
|
|
" barrier(CLK_LOCAL_MEM_FENCE);\n"
|
|
" // do reduction in shared mem\n"
|
|
" for(unsigned int s = localSize >> 1; s > 0; s >>= 1)\n"
|
|
" {\n"
|
|
" if(tid < s) \n"
|
|
" {\n"
|
|
" sdata[tid] += sdata[tid + s];\n"
|
|
" }\n"
|
|
" barrier(CLK_LOCAL_MEM_FENCE);\n"
|
|
" }\n"
|
|
"\n"
|
|
" // write result for this block to global mem\n"
|
|
" if(tid == 0) output[get_group_id(0)] = sdata[0];\n"
|
|
"}\n";
|
|
|
|
static const char* local_vec4_reduction =
|
|
"__kernel void local_reduction(__global uint4* input, __global uint4* "
|
|
"output, __local uint4* sdata)\n"
|
|
"{\n"
|
|
" // load shared mem\n"
|
|
" const uint msk = (uint)3;\n"
|
|
" const uint shft = (uint)8;\n"
|
|
" unsigned int tid = get_local_id(0);\n"
|
|
"\n"
|
|
" unsigned int localSize = get_local_size(0);\n"
|
|
" unsigned int stride = get_global_id(0) * 2;\n"
|
|
" uint4 data1 = input[stride];\n"
|
|
" uint4 data2 = input[stride + 1];\n"
|
|
" uint4 sum = 0;\n"
|
|
" for( int i = 0; i < 4; i++)\n"
|
|
" {\n"
|
|
" sum += (data1 & msk) + (data2 & msk);\n"
|
|
" data1 = data1 >> shft;\n"
|
|
" data2 = data2 >> shft;\n"
|
|
" }\n"
|
|
" sdata[tid] = sum;"
|
|
"\n"
|
|
" barrier(CLK_LOCAL_MEM_FENCE);\n"
|
|
" // do reduction in shared mem\n"
|
|
" for(unsigned int s = localSize >> 1; s > 0; s >>= 1)\n"
|
|
" {\n"
|
|
" if(tid < s) \n"
|
|
" {\n"
|
|
" sdata[tid] += sdata[tid + s];\n"
|
|
" }\n"
|
|
" barrier(CLK_LOCAL_MEM_FENCE);\n"
|
|
" }\n"
|
|
"\n"
|
|
" // write result for this block to global mem\n"
|
|
" if(tid == 0) output[get_group_id(0)] = sdata[0];\n"
|
|
"}\n";
|
|
|
|
static const char* local_atomics_reduction =
|
|
"#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
|
|
"__kernel void local_reduction(__global uint* input, __global uint* "
|
|
"output, __local uint* sdata)\n"
|
|
"{\n"
|
|
" // load shared mem\n"
|
|
" const uint msk = (uint)3;\n"
|
|
" const uint shft = (uint)8;\n"
|
|
" unsigned int tid = get_local_id(0);\n"
|
|
"\n"
|
|
" unsigned int localSize = get_local_size(0);\n"
|
|
" unsigned int stride = get_global_id(0) * 2;\n"
|
|
" unsigned int data1 = input[stride];\n"
|
|
" unsigned int data2 = input[stride + 1];\n"
|
|
" unsigned int sum = 0;\n"
|
|
" for( int i = 0; i < 4; i++)\n"
|
|
" {\n"
|
|
" sum += (data1 & msk) + (data2 & msk);\n"
|
|
" data1 = data1 >> shft;\n"
|
|
" data2 = data2 >> shft;\n"
|
|
" }\n"
|
|
" sdata[tid] = sum;"
|
|
"\n"
|
|
" barrier(CLK_LOCAL_MEM_FENCE);\n"
|
|
" // do reduction in shared mem\n"
|
|
" for(unsigned int s = localSize >> 1; s > 0; s >>= 1)\n"
|
|
" {\n"
|
|
" if(tid < s) \n"
|
|
" {\n"
|
|
" atom_add( &(sdata[tid]), sdata[tid + s]);\n"
|
|
" }\n"
|
|
" barrier(CLK_LOCAL_MEM_FENCE);\n"
|
|
" }\n"
|
|
"\n"
|
|
" // write result for this block to global mem\n"
|
|
" if(tid == 0) output[get_group_id(0)] = sdata[0];\n"
|
|
"}\n";
|
|
|
|
static const char* local_vec4_atomics_reduction =
|
|
"#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
|
|
"__kernel void local_reduction(__global uint4* input, __global uint4* "
|
|
"output, __local uint4* sdata)\n"
|
|
"{\n"
|
|
" // load shared mem\n"
|
|
" const uint msk = (uint)3;\n"
|
|
" const uint shft = (uint)8;\n"
|
|
" unsigned int tid = get_local_id(0);\n"
|
|
"\n"
|
|
" unsigned int localSize = get_local_size(0);\n"
|
|
" unsigned int stride = get_global_id(0) * 2;\n"
|
|
" uint4 data1 = input[stride];\n"
|
|
" uint4 data2 = input[stride + 1];\n"
|
|
" uint4 sum = 0;\n"
|
|
" for( int i = 0; i < 4; i++)\n"
|
|
" {\n"
|
|
" sum += (data1 & msk) + (data2 & msk);\n"
|
|
" data1 = data1 >> shft;\n"
|
|
" data2 = data2 >> shft;\n"
|
|
" }\n"
|
|
" sdata[tid] = sum;"
|
|
"\n"
|
|
" barrier(CLK_LOCAL_MEM_FENCE);\n"
|
|
" // do reduction in shared mem\n"
|
|
" for(unsigned int s = localSize >> 1; s > 0; s >>= 1)\n"
|
|
" {\n"
|
|
" if(tid < s) \n"
|
|
" {\n"
|
|
" atom_add( &(sdata[tid]).x, sdata[tid + s].x);\n"
|
|
" atom_add( &(sdata[tid]).y, sdata[tid + s].y);\n"
|
|
" atom_add( &(sdata[tid]).z, sdata[tid + s].z);\n"
|
|
" atom_add( &(sdata[tid]).w, sdata[tid + s].w);\n"
|
|
" }\n"
|
|
" barrier(CLK_LOCAL_MEM_FENCE);\n"
|
|
" }\n"
|
|
"\n"
|
|
" // write result for this block to global mem\n"
|
|
" if(tid == 0) output[get_group_id(0)] = sdata[0];\n"
|
|
"}\n";
|