Introdução ao SYCL

menotti
75.4K views

Open Source Your Knowledge, Become a Contributor

Technology knowledge has to be shared and made accessible for free. Join the movement.

Create Content

Memória e sincronização SYCL: loop de redução

auto n_wgroups = (len + part_size - 1) / part_size;

Dentro do ciclo de redução, encontramos primeiro o número de grupos de trabalho para essa etapa de redução. É o tamanho que resta a ser reduzido dividido pelo número de elementos que cada grupo de trabalho reduz.

sycl::accessor<int32_t, 1, sycl::access::mode::read_write, sycl::access::target::local> local_mem(sycl::range<1>(wgroup_size)

Em seguida, no grupo de comandos, alocamos uma parte da memória local criando um acessador com access::target::local e um intervalo igual ao tamanho do grupo de trabalho. Verificamos o tamanho da memória anteriormente, então sabemos que ele está disponível. Como mencionado acima, essa região da memória parece diferente para cada grupo de trabalho e seu uso é para armazenamento temporário.

Você pode se perguntar: por que nos preocupamos em usar a memória local quando podemos realizar toda a operação em nível global? A resposta é que é muito mais rápido. A memória local (geralmente) é fisicamente mais próxima do chip do que a global e não sofre de problemas como o compartilhamento falso, uma vez que é exclusiva para cada unidade de computação. Portanto, é uma boa idéia sempre executar todas as operações temporárias na memória local para obter o melhor desempenho.

auto global_mem = buf.get_access<sycl::access::mode::read_write>(cgh);

Também obtemos um acessador para os dados disponíveis na memória global. Dessa vez, o get_access é explicitamente qualificado com access::target::global_buffer, enquanto anteriormente assumia esse valor por padrão.

cgh.parallel_for<class reduction_kernel>(
      sycl::nd_range<1>(n_wgroups * wgroup_size, wgroup_size),
      [=] (sycl::nd_item<1> item)

Por fim, lançamos um kernel paralelo. Usamos a variante nd_range, que nos permite especificar o tamanho global e local. O construtor nd_range recebe dois objetos de intervalo da mesma dimensionalidade que ele. O primeiro descreve o número de itens de trabalho por dimensão (lembre-se de que pode haver até três dimensões). O segundo argumento de intervalo para nd_range<n> descreve o número de itens de trabalho em um grupo de trabalho. Para encontrar o número de grupos de trabalho por dimensão, divida o primeiro argumento pelo segundo. Nesse caso, o resultado é n_wgroups, que é quantos grupos de trabalho serão instanciados. Nesta variante, o kernel lambda usa um argumento nd_item. Ele representa o item de trabalho atual e os métodos de recursos para obter informações detalhadas, como informações locais, globais e de grupo de trabalho.

Como cada etapa do ciclo de redução produz um número por grupo de trabalho, configuramos o len como n_wgroups a cada iteração, o que continuará reduzindo nos resultados.

Redução paralela
Open Source Your Knowledge: become a Contributor and help others learn. Create New Content