Размер данных в инструкции на отношения Деформации в CUDA

Question

Я пытался увидеть количество инструкций, выполняемых в ядре, когда размер типа данных изменен

Для того, чтобы получить структуру данных пользовательского размера я создал-структуру, как следующее,

#define DATABYTES 40 

__host__ __device__ 
struct floatArray 
{ 
    float a[DATABYTES/4]; 
}; 

, а затем создали ядро ​​просто скопировать выше массив типа данных из одного массива в другую

__global__ 
void copy_large_data(floatArray * d_in, floatArray * d_out) 
{ 
    d_out[threadIdx.x] = d_in[threadIdx.x]; 
} 

затем вызывается вышеупомянутое ядро ​​для о олько 32 нить с одного блоком

floatArray * d_in; 
floatArray * d_out; 

cudaMalloc(&d_in, 32 * sizeof(floatArray)); 
cudaMalloc(&d_out, 32 * sizeof(floatArray)); 

copy_large_data<<<1, 32>>>(d_in, d_out); 

Когда я профилировать программу с помощью nvprof и проверен на instructions per warp я мог видеть, что значение параметра изменяется с изменением величины DATABYTES.

Мои вопросы: является ли причиной увеличения количества этих команд из-за массива внутри структуры floatArray. Поскольку, когда мы вызываем копию в ядре, она фактически расширяет и копирует каждый элемент массива a внутри структуры floatArray, создавая дополнительные инструкции.

Есть ли способ скопировать настраиваемую структурную переменную в ядро ​​с использованием одной инструкции?

Answer 1

Вы правы в своем предположении, что количество команд копирования увеличивается при изменении размера массива. Вы можете проверить это в коде PTX и в сборке, как я покажу ниже.

Максимальный размер инструкции загрузки/хранения составляет 128 бит, см., Например, here. Это означает, что для вашего дела вы все равно можете улучшить в 4 раза, используя float4 вместо float.

В качестве альтернативы вы можете указать выравнивание вашей структуры данных в явном виде, как объяснено в programming guide:

#define DATABYTES 32 
struct __align__(16) floatArray 
{ 
    float a[DATABYTES/4]; 
}; 

---

Чтобы увидеть код PTX сгенерировать объектный файл nvcc -c ... и использовать cubobjdump --dump-ptx objfile.o. Для примера соответствующая часть выглядит следующим образом:

ld.global.f32 %f1, [%rd7]; 
ld.global.f32 %f2, [%rd7+4]; 
ld.global.f32 %f3, [%rd7+8]; 
ld.global.f32 %f4, [%rd7+12]; 
ld.global.f32 %f5, [%rd7+16]; 
ld.global.f32 %f6, [%rd7+20]; 
ld.global.f32 %f7, [%rd7+24]; 
ld.global.f32 %f8, [%rd7+28]; 
ld.global.f32 %f9, [%rd7+32]; 
ld.global.f32 %f10, [%rd7+36]; 
st.global.f32 [%rd6+36], %f10; 
st.global.f32 [%rd6+32], %f9; 
st.global.f32 [%rd6+28], %f8; 
st.global.f32 [%rd6+24], %f7; 
st.global.f32 [%rd6+20], %f6; 
st.global.f32 [%rd6+16], %f5; 
st.global.f32 [%rd6+12], %f4; 
st.global.f32 [%rd6+8], %f3; 
st.global.f32 [%rd6+4], %f2; 
st.global.f32 [%rd6], %f1; 

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

Соответственно вы можете проверить сборку с cubobjdump --dump-sass objfile.o