Как уменьшить дублирование кода между ядрами OpenCL?
У меня есть несколько похожих ядер для генерации случайных данных и хранения их в глобальной памяти. Я всегда использую один и тот же алгоритм для рандомизации, но из-за проблем с переменной областью (мне нужно отслеживать данные) мне не удается избежать серьезных дублирований кода.
Есть ли способы избежать этого? Генерация случайных данных в OpenCL кажется довольно стандартной задачей, но это противоречит любым хорошим стандартам кодирования, чтобы иметь такой уровень дублирования кода. Например, вот два моих ядра:
////////////////////////////////////////////////////////////////////////////////
// OpenCL Kernel for Mersenne Twister RNG -- applied to AWGN channel
////////////////////////////////////////////////////////////////////////////////
__kernel void MersenneTwisterAWGN(__global double* d_Rand,
__global int* seeds,
__global long* inputcw,
int nPerRng, float sigma)
{
int globalID = get_global_id(0);
double c = 2.0/(sigma*sigma);
int iState, iState1, iStateM, iOut;
unsigned int mti, mti1, mtiM, x;
unsigned int mt[MT_NN];
//Initialize current state
mt[0] = seeds[globalID];
for (iState = 1; iState < MT_NN; iState++)
mt[iState] = (1812433253U*(mt[iState-1]^(mt[iState-1]>>30))+iState) & MT_WMASK;
iState = 0;
mti1 = mt[0];
for (iOut = 0; iOut < nPerRng; iOut=iOut+2) {
iState1 = iState + 1;
iStateM = iState + MT_MM;
if(iState1 >= MT_NN) iState1 -= MT_NN;
if(iStateM >= MT_NN) iStateM -= MT_NN;
mti = mti1;
mti1 = mt[iState1];
mtiM = mt[iStateM];
// MT recurrence
x = (mti & MT_UMASK) | (mti1 & MT_LMASK);
x = mtiM ^ (x >> 1) ^ ((x & 1) ? matrix_a : 0);
mt[iState] = x;
iState = iState1;
//Tempering transformation
x ^= (x >> MT_SHIFT0);
x ^= (x << MT_SHIFTB) & mask_b;
x ^= (x << MT_SHIFTC) & mask_c;
x ^= (x >> MT_SHIFT1);
double u1 = ((double)x + 1.0f) / 4294967296.0f;
iState1 = iState + 1;
iStateM = iState + MT_MM;
if(iState1 >= MT_NN) iState1 -= MT_NN;
if(iStateM >= MT_NN) iStateM -= MT_NN;
mti = mti1;
mti1 = mt[iState1];
mtiM = mt[iStateM];
// MT recurrence
x = (mti & MT_UMASK) | (mti1 & MT_LMASK);
x = mtiM ^ (x >> 1) ^ ((x & 1) ? matrix_a : 0);
mt[iState] = x;
iState = iState1;
//Tempering transformation
x ^= (x >> MT_SHIFT0);
x ^= (x << MT_SHIFTB) & mask_b;
x ^= (x << MT_SHIFTC) & mask_c;
x ^= (x >> MT_SHIFT1);
double u2 = ((double)x + 1.0f) / 4294967296.0f;
double r = sqrt(-2.0f * log(u1));
double phi = 2 * PI * u2;
u1 = r * cos(phi);
u1 = inputcw[iOut]+sigma*u1;
u1=1/(1+exp(-c*u1));
d_Rand[globalID * nPerRng + iOut]=log((1-u1)/u1);
if (iOut!=nPerRng-1) {
u2 = r * sin(phi);
u2 = inputcw[iOut+1]+sigma*u2;
u2=1/(1+exp(-c*u2));
u2=log((1-u2)/u2);
d_Rand[globalID * nPerRng + iOut+1]=u2;
}
}
}
а также
////////////////////////////////////////////////////////////////////////////////
// OpenCL Kernel for Mersenne Twister RNG -- applied to BSC channel
////////////////////////////////////////////////////////////////////////////////
__kernel void MersenneTwisterBSC(__global double* d_Rand,
__global int* seeds,
__global long* inputcw,
int nPerRng, float flipProb)
{
int globalID = get_global_id(0);
int iState, iState1, iStateM, iOut;
unsigned int mti, mti1, mtiM, x;
unsigned int mt[MT_NN];
//Initialize current state
mt[0] = seeds[globalID];
for (iState = 1; iState < MT_NN; iState++)
mt[iState] = (1812433253U*(mt[iState-1]^(mt[iState-1]>>30))+iState) & MT_WMASK;
iState = 0;
mti1 = mt[0];
for (iOut = 0; iOut < nPerRng; iOut=iOut+1) {
iState1 = iState + 1;
iStateM = iState + MT_MM;
if(iState1 >= MT_NN) iState1 -= MT_NN;
if(iStateM >= MT_NN) iStateM -= MT_NN;
mti = mti1;
mti1 = mt[iState1];
mtiM = mt[iStateM];
// MT recurrence
x = (mti & MT_UMASK) | (mti1 & MT_LMASK);
x = mtiM ^ (x >> 1) ^ ((x & 1) ? matrix_a : 0);
mt[iState] = x;
iState = iState1;
//Tempering transformation
x ^= (x >> MT_SHIFT0);
x ^= (x << MT_SHIFTB) & mask_b;
x ^= (x << MT_SHIFTC) & mask_c;
x ^= (x >> MT_SHIFT1);
double c = log((1-flipProb)/flipProb);
double u = ((double)x + 1.0f) / 4294967296.0f;
u = (2*isless(u,flipProb)-1)*inputcw[iOut]*c;
d_Rand[globalID * nPerRng + iOut]=u;
}
}
Есть ли способы, уловки или методы, чтобы избежать этого? Кажется, подпрограммы не могут правильно использовать переменные (особенно mt), так что мне не удалось сократить его так, как это позволили бы другие языки.
Или я должен просто принять это как необходимое зло в OpenCL и продолжать управлять 10 различными ядрами таким образом?
1 ответ
На сайте Хроноса написано
Программы OpenCL также могут содержать вспомогательные функции и постоянные данные, которые могут использоваться функциями __kernel.
Пример для генерации случайного числа между 0.0f и 1.0f на поток:
Основная функция для итерации семени:
uint wang_hash(uint seed)
{
seed = (seed ^ 61) ^ (seed >> 16);
seed *= 9;
seed = seed ^ (seed >> 4);
seed *= 0x27d4eb2d;
seed = seed ^ (seed >> 15);
return seed;
}
Инициализация и итерация каждого семени потоков:
// id=thread id, rnd=seed array
void wang_rnd_init(__global unsigned int * rnd,int id)
{
uint maxint=0;
maxint--; // could be a 0xFFFFFFFF
uint rndint=wang_hash(id);
rnd[id]=rndint;
}
// id=thread id, rnd=seed array
float wang_rnd(__global unsigned int * rnd,int id)
{
uint maxint=0;
maxint--; // could be a 0xFFFFFFFF
uint rndint=wang_hash(rnd[id]);
rnd[id]=rndint;
return ((float)rndint)/(float)maxint;
}
Использование в ядре генератора случайных оттенков серого цвета:
__kernel void rnd_1(__global unsigned int * rnd, __global int *rgba)
{
int id=get_global_id(0);
float rgba_register=wang_rnd(rnd,id);
rgba[id] = ((int)(rgba_register * 255) << 24) | ((int)(rgba_register * 255) << 16) | ((int)(rgba_register * 255) << 8) | ((int)(rgba_register * 255));
}
и wang_rnd() можно использовать в других ядрах, не определяя его дважды, если они находятся в одном и том же скомпилированном контексте, так же как все соответствующие ядра и функции помещаются в один и тот же файл для компиляции.
Вспомогательные функции не ограничиваются регистрами и глобальной памятью. Они также могут принимать параметры локальной и постоянной памяти. Поскольку они в основном работают с памятью на стороне устройства, они также могут принимать и возвращать структуры.