Проверка device_vector внутри ядра CUDA не работает

Вы в основном ищите синхронное поведение глобальных потоков. Это нет-нет в программировании на GPU. В идеале каждый блок потоков является независимым и может выполнять работу на основе собственных данных и обработки. Создание потоковых блоков, которые зависят от результатов других потоковых блоков, чтобы завершить их работу, создает возможность условия взаимоблокировки. Предположим, у меня есть графический процессор с 14 SM (блоки выполнения потоков), и предположим, что я создаю 100 блоков потоков. Теперь предположим, что блоки потоков 0-13 ожидают, чтобы блок 99 потоков освободил блокировку (например, записать нулевое значение в конкретное место). Теперь предположим, что эти первые 14 потоковых блоков начинают выполняться на 14 SM, возможно, зацикливаясь, вращаясь на значении блокировки. В графическом процессоре нет механизма, который гарантировал бы, что блок 99 потока будет выполняться первым или даже вообще выполняться, если блоки потоков 0-13 имеют связанные с SM блоки.

Давайте не будем вдаваться в вопросы о том, «как насчет остановок GMEM, которые вынуждают вытеснять блоки потоков 0–13», потому что ничто из этого не гарантирует, что блок потоков 99 получит приоритет для выполнения в любой момент. Единственное, что гарантирует выполнение потокового блока 99, — это удаление (то есть завершение) других потоковых блоков. Но если другие потоковые блоки вращаются, ожидая результатов от потокового блока 99, это может никогда не произойти.

Хороший, совместимый с прямым масштабированием код GPU зависит от независимой параллельной работы. Поэтому вам рекомендуется переделать свой алгоритм, чтобы сделать работу, которую вы пытаетесь выполнить, независимой, по крайней мере, на уровне межпоточных блоков.

Если вам необходимо выполнить глобальную синхронизацию потоков, запуск ядра — единственная действительно гарантированная точка для этого, и, таким образом, ваш первый подход — это рабочий подход.

Чтобы помочь с этим, может быть полезно изучить, как алгоритмы сокращения реализуются на графическом процессоре. Различные типы сокращений имеют зависимости во всех потоках, но, создав промежуточные результаты, мы можем разбить работу на независимые части. Затем можно объединить независимые фрагменты с использованием многоядерного подхода (или некоторых других более продвинутых подходов), чтобы ускорить то, что составляет последовательный алгоритм.

Ваше ядро ​​на самом деле мало что делает. Он устанавливает один массив, равный его индексу, то есть a [i] = i; и он устанавливает другой массив во все нули (хотя и последовательно) b [i] = 0 ;. Чтобы показать пример вашего первого «ускоренного» кода, вы можете сделать что-то вроде этого:

    // include the header files
#include <iostream>
#include <stdio.h>
#include <time.h>

#include "cuda.h"#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

using namespace std;

//kernel function
__global__
void dosomething(int *d_bPtr, int count, int* d_bStopPtr)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
while(tid<count)
{
d_bPtr[tid]=tid;
while(d_bStopPtr[tid]!=0)
// only if the arrary cell before it is 0, then change it to 0 too
if (tid==0) d_bStopPtr[tid] =0;
else if (d_bStopPtr[tid-1]==0 )
d_bStopPtr[tid]=0;
tid += blockDim.x;
}
}

int main()
{
int count=100000;
// define the vectors
thrust::host_vector <int> h_a(count);
thrust::device_vector <int> d_b(count,0);
int* d_bPtr=thrust::raw_pointer_cast(&d_b[0]);
thrust::device_vector <int> d_bStop(count,1);
int* d_bStopPtr=thrust::raw_pointer_cast(&d_bStop[0]);
// get the device property
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);

//    int threadsPerBlock = prop.maxThreadsDim[0];
int threadsPerBlock = 32;
//    int blocksPerGrid = min(prop.maxGridSize[0], (count + threadsPerBlock - 1) / threadsPerBlock);
int blocksPerGrid = 1;
//copy device to host
thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//run the kernel
//    while(d_bStop[count-1])
//    {
dosomething<<<blocksPerGrid, threadsPerBlock>>>(d_bPtr,count,d_bStopPtr);
//    }
//copy device back to host again
cudaDeviceSynchronize();
thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//wait to see the console output
int x;
cin>>x;
return 0;
}

На моей машине это ускоряет время выполнения с 10 секунд до почти мгновенного (намного меньше 1 секунды). Обратите внимание, что это не лучший пример программирования CUDA, потому что я запускаю только один блок из 32 потоков. Этого недостаточно для эффективного использования машины. Но работа, выполняемая вашим ядром, настолько тривиальна, что я не уверен, что будет хорошим примером. Я мог бы просто создать ядро, которое устанавливает один массив в его индекс a [i] = i; а другой массив к нулю b [i] = 0; все параллельно. Это было бы еще быстрее, и мы могли бы использовать всю машину таким образом.