Для вашего примера кода итерации цикла кажутся независимыми, что облегчает задачу.
Мы можем начать с подсчета общего числа итераций: 8 * (100 + 100) = 1600.
Поэтому мы бы хотели ускорить сетку из как минимум 1600 потоков, если мы хотим полностью распределить работу.
Ваша переменная m
не кажется зависимой от цикла, поэтому b[m]
фактически является константой для кода, который вы показали.
Другим элементом планирования является определение индексации j
на основе k
, но это также довольно просто, это просто k%100
.
Последний элемент планирования - учет изменения знака c
на протяжении итераций цикла.Мы можем использовать k/100
, чтобы помочь нам определить это.
Поэтому мы создадим глобально уникальную переменную индекса потока и используем ее в качестве нашей k
переменной:
__global__ void my_kernel(T1 *a, T2 *b, T3 *c, int N){
int k = threadIdx.x+blockDim.x*blockIdx.x;
int m = <some value>;
int j = k%100;
if (k < N)
a[k] = b[m] + ((k/100)&1)?(-c[j]):c[j];
}
Убедитесь, чтозапустить его как минимум с 1600 нитями (в 1D сетке), передав 1600 для N
.T1
, T2
, T3
могут быть typedefs для любых типов, соответствующих a
, b
, c
.
Обратите внимание, что в вашем коде есть следующее:
for(int j=0; j<100; j+=)
^
This doesn't make sense to me.
Итак, я только что сделал вид, что это было вместо этого:
for(int j=0; j<100; j++)
с такой модификацией (не такой, как было отредактировано в вопросе):
for (int m=0; m<7; m++) // outer loop
^
Единственное изменение заключается в вычислении переменной m
:
__global__ void my_kernel(T1 *a, T2 *b, T3 *c, int N){
int k = threadIdx.x+blockDim.x*blockIdx.x;
int m = k/200; // change this line
int j = k%100;
if (k < N)
a[k] = b[m] + ((k/100)&1)?(-c[j]):c[j];
}