Характеризуя вычисления, которые получают выгоду от загрузки / хранения cuda с использованием смещений по сравнению с использованием векторов
Следующий фрагмент кода иллюстрирует два способа загрузки и хранения данных, один с использованием смещений, а другой с использованием вектора. На иллюстрации вычисление просто умножить на 2.
__global__ foo( ..., int *d_i, int *d_o, ...){
#ifdef LOAD_STORE_USING_OFFSETS
d_o[x + (0*offset)] = 2*d_i[x + (0*offset)];
d_o[x + (1*offset)] = 2*d_i[x + (1*offset)];
d_o[x + (2*offset)] = 2*d_i[x + (2*offset)];
d_o[x + (3*offset)] = 2*d_i[x + (3*offset)];
#endif
#ifdef LOAD_STORE_USING_VECTORS
int4 v_4 = ((int4 *)d_i)[x];
v_4.x = 2 * v_4.x;
v_4.y = 2 * v_4.y;
v_4.z = 2 * v_4.z;
v_4.w = 2 * v_4.w;
((int4 *)d_o)[x] = v_4;
#endif
}
Для вопросов ниже, предположим, что каждое значение обрабатывается одинаково ("смущающе параллельно").
Вопросы:
1) Я сталкивался с людьми, которые утверждают, что векторная версия должна быть предпочтительной всегда, так как она требует только одну инструкцию для загрузки / сохранения - но даже если для офсетной версии требуется 4 инструкции на загрузку / хранение, не будет трех дополнительных инструкций быть скрытым под задержкой завершения загрузки / сохранения для первой инструкции (т. е. варп сможет выполнить все 4 загрузки / сохранения перед блокировкой)?
2) Существуют ли определенные типы вычислений, для которых один метод будет работать лучше, чем другой?
3) Если да, то каковы характеристики вычислений, для которых лучше подходит каждый метод?