Freeverb kernel attempt #2

[Ryan Stewart]

Archana,

Here's the code I came up with for a monolithic Freeverb kernel.  Needs tweaking but should save some time unrolling the previous code.  As you can see it is really long and I'm pessimistic whether it will provide any significant benefits over the previous attempt, but it's worth a shot I suppose.  This doesn't include the setup code, haven't had time to set that up.

// This was converted to work with long values instead of float values. I don't know if it'll still work.

// May want to experiment with the C version to get parameter values right.

const char *KernelSource = "\n"\

"__kernel void freeverb (__global long *input, __global long *output, int numsamples, \n"\

" __global long *lbcf_fs, __global long *lbcf_size, __global long *lbcf_buffer0, \n"\

" __global long *lbcf_buffer1, __global long *lbcf_buffer2, __global long *lbcf_buffer3,\n"\

" __global long *lbcf_buffer4, __global long *lbcf_buffer5, __global long *lbcf_buffer6,\n"\

" __global long *lbcf_buffer7, __global long *lbcf_idx, __global long *apf_idx, \n"\

" __global long *apf_buffer0, __global long *apf_buffer1, __global long *apf_buffer2, \n"\

" __global long *apf_size) { \n"\

" int i = 0, j = 0; \n"\

" long out, in, outReg; \n"\ 

" long gain = 1, wetLevel=1, dryLevel=4, damp=1, lbcf_feedback=84, apf_feedback=84;\n"\ // Parameters

" while (numsamples--) {\n"\

" out = 0;\n"\

" in = input[i]*gain;"

" // low pass feedback comb filters \n"\

" j = lbcf_idx[0];\n"\

" outReg = lbcf_buffer0[j];\n"\

" lbcf_fs[0] = out*damp + lbcf_fs[0]*(1-damp);\n"\

" lbcf_buffer0[j] = in +lbcf_fs[0]*lbcf_feedback;\n"\

" lbcf_idx[0] = (j + 1) % lbcf_size[0];\n"\

" out += outReg;\n"\

" j = lbcf_idx[1];\n"\

" outReg = lbcf_buffer1[j];\n"\

" lbcf_fs[1] = out*damp + lbcf_fs[1]*(1-damp);\n"\

" lbcf_buffer1[j] = in +lbcf_fs[1]*lbcf_feedback;\n"\

" lbcf_idx[1] = (j + 1) % lbcf_size[1];\n"\

" j = lbcf_idx[2];\n"\

" outReg = lbcf_buffer2[j];\n"\

" lbcf_fs[2] = out*damp + lbcf_fs[2]*(1-damp);\n"\

" lbcf_buffer2[j] = in +lbcf_fs[2]*lbcf_feedback;\n"\

" lbcf_idx[2] = (j + 1) % lbcf_size[2];\n"\

" out += outReg;\n"\

" j = lbcf_idx[3];\n"\

" outReg = lbcf_buffer3[j];\n"\

" lbcf_fs[3] = out*damp + lbcf_fs[3]*(1-damp);\n"\

" lbcf_buffer3[j] = in +lbcf_fs[3]*lbcf_feedback;\n"\

" lbcf_idx[3] = (j + 1) % lbcf_size[3];\n"\

" out += outReg;\n"\

" j = lbcf_idx[4];\n"\

" outReg = lbcf_buffer4[j];\n"\

" lbcf_fs[4] = out*damp + lbcf_fs[4]*(1-damp);\n"\

" lbcf_buffer4[j] = in +lbcf_fs[4]*lbcf_feedback;\n"\

" lbcf_idx[4] = (j + 1) % lbcf_size[4];\n"\

" out += outReg;\n"\

" j = lbcf_idx[5];\n"\

" outReg = lbcf_buffer5[j];\n"\

" lbcf_fs[5] = out*damp + lbcf_fs[5]*(1-damp);\n"\

" lbcf_buffer5[j] = in +lbcf_fs[5]*lbcf_feedback;\n"\

" lbcf_idx[5] = (j + 1) % lbcf_size[5];\n"\

" out += outReg;\n"\

" j = lbcf_idx[6];\n"\

" outReg = lbcf_buffer6[j];\n"\

" lbcf_fs[6] = out*damp + lbcf_fs[6]*(1-damp);\n"\

" lbcf_buffer6[j] = in +lbcf_fs[6]*lbcf_feedback;\n"\

" lbcf_idx[6] = (j + 1) % lbcf_size[6];\n"\

" out += outReg;\n"\

" j = lbcf_idx[7];\n"\

" outReg = lbcf_buffer7[j];\n"\

" lbcf_fs[7] = out*damp + lbcf_fs[7]*(1-damp);\n"\

" lbcf_buffer7[j] = in +lbcf_fs[7]*lbcf_feedback;\n"\

" lbcf_idx[7] = (j + 1) % lbcf_size[7];\n"\

" out += outReg;\n"\

" // all pass filters \n"\

" in = out;\n"\

" j = apf_idx[0];\n"\

" outReg = apf_buffer0[j];\n"\

" apf_buffer0[j] = in + outReg*apf_feedback;\n"\

" outReg = outReg - apf_buffer0[j]*apf_feedback;\n"\

" apf_idx[0] = (j + 1) % apf_size[0];\n"\

" in = outReg;\n"\

" j = apf_idx[1];\n"\

" outReg = apf_buffer1[j];\n"\

" apf_buffer1[j] = in + outReg*apf_feedback;\n"\

" outReg = outReg - apf_buffer1[j]*apf_feedback;\n"\

" apf_idx[1] = (j + 1) % apf_size[1];\n"\

" in = outReg;\n"\

" j = apf_idx[2];\n"\

" outReg = apf_buffer2[j];\n"\

" apf_buffer2[j] = in + outReg*apf_feedback;\n"\

" outReg = outReg - apf_buffer2[j]*apf_feedback;\n"\

" apf_idx[2] = (j + 1) % apf_size[2];\n"\

" out = outReg;\n"\

" \n"\

" out = out*wetLevel + in*dryLevel;\n"\

" out++;\n"\

" }\n"\

"}\n"\

[archana raja]

no thats k..i will integrate it wth the code tomm and let u knw