__kernel void kernel_scan(__global float* input, __global float* output) {
int global_idx = get_global_id(0);
int local_idx = get_local_id(0);
int block_size = get_local_size(0);
int group_id = get_group_id(0);
output[global_idx] = input[global_idx];
mem_fence(CLK_GLOBAL_MEM_FENCE);
for(int i = 1; i < block_size; i <<= 1) {
if(global_idx >= i)
output[global_idx] += output[global_idx - i];
mem_fence(CLK_GLOBAL_MEM_FENCE);
}
}
inline void compute_chest(task *task) {
symbol_data *symbolData = task->symbolData;
int rx = task->rx;
int layer = task->layer;
int res_power[4] = {0, 0, 0, 0};
mf(&symbolData->data->in_data[symbolData->slot][3][rx][symbolData->startSc], &symbolData->data->in_rs[symbolData->slot][symbolData->startSc][layer], symbolData->nmbSc, symbolData->layer_data[layer][rx], &symbolData->pow[rx]);
ifft(symbolData->layer_data[layer][rx], symbolData->nmbSc, symbolData->data->fftw[symbolData->slot]);
chest(symbolData->layer_data[layer][rx], symbolData->pow[rx], symbolData->nmbSc, symbolData->layer_data[layer][rx], &res_power[rx]);
/* Put power values in the R matrix */
symbolData->R[layer][rx] = cmake(res_power[rx],0);
fft(symbolData->layer_data[layer][rx], symbolData->nmbSc, symbolData->data->fftw[symbolData->slot]);
/* Mark the task as computed */
mem_fence();
task->computed = true;
}
void *thread1(void *num)
{
int i = *((int *)num);
for (int it = 0; it < AMOUNT; it++)
{
#ifdef LOCK_ENABLED
flag1 = 1;
turn = 0;
mem_fence();
while (flag0 == 1 && turn == 0)
wait();
#endif
// critical section
resource++;
// end of critical section
#ifdef LOCK_ENABLED
flag1 = 0;
#endif
}
printf("Contagem thread1: %d\n", resource);
}
inline void compute_symbol(task *task) {
symbol_data *symbolData = task->symbolData;
int ofdm = task->ofdm;
int ofdm_count = task->ofdm_count;
int layer = task->layer;
int nmbSc = symbolData->nmbSc;
int slot = symbolData->slot;
complex* in[4];
int index_out;
in[0] = &symbolData->data->in_data[symbolData->slot][ofdm][0][symbolData->startSc];
in[1] = &symbolData->data->in_data[symbolData->slot][ofdm][1][symbolData->startSc];
in[2] = &symbolData->data->in_data[symbolData->slot][ofdm][2][symbolData->startSc];
in[3] = &symbolData->data->in_data[symbolData->slot][ofdm][3][symbolData->startSc];
/* Put all demodulated symbols in one long vector */
index_out = nmbSc*ofdm_count + slot*(OFDM_IN_SLOT-1)*nmbSc + layer*2*(OFDM_IN_SLOT-1)*nmbSc;
ant_comb(in, symbolData->combWeight[layer], nmbSc, &symbolData->symbols[index_out]);
/* Now transform data back to time plane */
ifft(&symbolData->symbols[index_out], nmbSc, symbolData->data->fftw[symbolData->slot]);
/* Mark the task as computed */
mem_fence();
task->computed = true;
}
