OpenCL FIR

gistfile1.txt

#define num_taps 150

__constant float taps[num_taps] = {
4.5528412686211801e-04, 4.7143479283810269e-04, 4.9265664783479932e-04, 5.1935133206458756e-04, 5.5190883344588668e-04, 5.9070568604295396e-04, 6.3610320042731044e-04, 6.8844573336344591e-04, 7.4805900411431905e-04, 8.1524846452430972e-04, 8.9029772986915878e-04, 9.7346707726454198e-04, 1.0649920181980451e-03, 1.1650819514941403e-03, 1.2739189027393427e-03, 1.3916563558860344e-03, 1.5184181824196682e-03, 1.6542976731164991e-03, 1.7993566770389934e-03, 1.9536248520151279e-03, 2.1170990304274768e-03, 2.2897427036999980e-03, 2.4714856284164220e-03, 2.6622235565360503e-03, 2.8618180916923654e-03, 3.0700966730691317e-03, 3.2868526878496000e-03, 3.5118457127290654e-03, 3.7448018844713667e-03, 3.9854143989781774e-03, 4.2333441378280190e-03, 4.4882204207321872e-03, 4.7496418818490762e-03, 5.0171774673990199e-03, 5.2903675515306690e-03, 5.5687251669092235e-03, 5.8517373460285097e-03, 6.1388665687948886e-03, 6.4295523114932876e-03, 6.7232126918258892e-03, 7.0192462043144020e-03, 7.3170335399784384e-03, 7.6159394838475883e-03, 7.9153148835343929e-03, 8.2144986817911259e-03, 8.5128200056965395e-03, 8.8096003048703988e-03, 9.1041555308952089e-03, 9.3957983499366408e-03, 9.6838403803978373e-03, 9.9675944473185795e-03, 1.0246376845139083e-02, 1.0519509600389887e-02, 1.0786322725844880e-02, 1.1046156457683476e-02, 1.1298363467250896e-02, 1.1542311039081835e-02, 1.1777383206962653e-02, 1.2002982839949956e-02, 1.2218533670438616e-02, 1.2423482256579165e-02, 1.2617299871582523e-02, 1.2799484312717831e-02, 1.2969561623106257e-02, 1.3127087719738292e-02, 1.3271649921493347e-02, 1.3402868371316995e-02, 1.3520397347111102e-02, 1.3623926456314295e-02, 1.3713181709592454e-02, 1.3787926469519981e-02, 1.3847962270610194e-02, 1.3893129507545419e-02, 1.3923307988962729e-02, 1.3938417354666922e-02, 1.3938417354666922e-02, 1.3923307988962729e-02, 1.3893129507545419e-02, 1.3847962270610194e-02, 1.3787926469519981e-02, 1.3713181709592454e-02, 1.3623926456314295e-02, 1.3520397347111102e-02, 1.3402868371316995e-02, 1.3271649921493347e-02, 1.3127087719738292e-02, 1.2969561623106257e-02, 1.2799484312717831e-02, 1.2617299871582523e-02, 1.2423482256579165e-02, 1.2218533670438616e-02, 1.2002982839949956e-02, 1.1777383206962653e-02, 1.1542311039081835e-02, 1.1298363467250896e-02, 1.1046156457683476e-02, 1.0786322725844880e-02, 1.0519509600389887e-02, 1.0246376845139083e-02, 9.9675944473185795e-03, 9.6838403803978373e-03, 9.3957983499366408e-03, 9.1041555308952089e-03, 8.8096003048703988e-03, 8.5128200056965395e-03, 8.2144986817911259e-03, 7.9153148835343929e-03, 7.6159394838475883e-03, 7.3170335399784384e-03, 7.0192462043144020e-03, 6.7232126918258892e-03, 6.4295523114932876e-03, 6.1388665687948886e-03, 5.8517373460285097e-03, 5.5687251669092235e-03, 5.2903675515306690e-03, 5.0171774673990199e-03, 4.7496418818490762e-03, 4.4882204207321872e-03, 4.2333441378280190e-03, 3.9854143989781774e-03, 3.7448018844713667e-03, 3.5118457127290654e-03, 3.2868526878496000e-03, 3.0700966730691317e-03, 2.8618180916923654e-03, 2.6622235565360503e-03, 2.4714856284164220e-03, 2.2897427036999980e-03, 2.1170990304274768e-03, 1.9536248520151279e-03, 1.7993566770389934e-03, 1.6542976731164991e-03, 1.5184181824196682e-03, 1.3916563558860344e-03, 1.2739189027393427e-03, 1.1650819514941403e-03, 1.0649920181980451e-03, 9.7346707726454198e-04, 8.9029772986915878e-04, 8.1524846452430972e-04, 7.4805900411431905e-04, 6.8844573336344591e-04, 6.3610320042731044e-04, 5.9070568604295396e-04, 5.5190883344588668e-04, 5.1935133206458756e-04, 4.9265664783479932e-04, 4.7143479283810269e-04, 4.5528412686211801e-04};


@kernel {
	int i = get_global_id(0);

	output[i] = 0;

	// The output at time 0 is the sum of the
	// previous num_taps samples times the taps
	int j;
	for(j = 0; j < num_taps; j++) {
		float2 sample = 0;

		int sample_index = i - (num_taps - 1);

		if (sample_index < 0) {
			sample = lastBlock[2048+sample_index];
		} else {
			sample = input[sample_index];
		}

		output[i] += sample * (float2)taps[j];
	}

	// Copy the input to the last block,
	// this ensures continuity
	lastBlock[i] = input[i];
}