Как уменьшить дублирование кода между ядрами OpenCL?

Question

У меня есть несколько подобных ядер для генерации случайных данных и сохранения их в глобальной памяти. Я всегда использую один и тот же алгоритм для рандомизации, но из-за проблем с переменной областью (мне нужно отслеживать данные) я не могу избежать серьезных дублирования кода.

Есть ли способы избежать этого? Генерация случайных данных в 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 разных ядер?

Answer 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() могут использоваться в других ядрах, не определяя их дважды, если они находятся в одном и том же скомпилированном контексте, так же как и сложение всех соответствующих ядер и функций в одном файле.

Вспомогательные функции не ограничиваются регистрами и глобальной памятью. Они также могут принимать локальные и постоянные параметры памяти. Поскольку они работают с памятью на стороне устройства в основном, они могут принимать и возвращать структуры.