ядро cuda - регистры

Question

Переменные Float4, определенные в ядре, должны храниться в регистрах! Я сделал простой тест. В первом ядре я использую регистры для оптимизации трафика памяти, во втором я читаю напрямую из глобальной памяти.

__global__ void kernel(float4 *arg1, float4 *arg2, float4 *arg3)
{
     int x = blockIdx.x * blockDim.x + threadIdx.x;

     float4 temp1 = arg2[x];
     float4 temp2 = arg3[x];
     //some computations using temp1 and temp2
     arg2[x] = temp1;
     arg3[x] = temp2;

     arg1[x] = make_float4(temp1.x, temp1.y, temp1.z, temp1.w);
}



 __global__ void kernel(float4 *arg1, float4 *arg2, float4 *arg3)
{
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     //some computations using a direct access to global memory
     //for example arg2[x].x
     arg1[x] = make_float4(arg2[x].x, arg2[x].y, arg2[x].z, arg2[x].w);
}

Первое ядро ​​на 9-10% быстрее. Разница не так уж и велика. При использовании регистров может принести больше пользы?

Answer 1

Во-первых, вы не можете сказать, что будет и не будет в регистрах исключительно на основе кода C. Это, конечно, не является источником разницы в производительности между двумя кодами. Фактически, оба ядра используют регистры для переменных float4, и код, с которым они компилируются, практически идентичен.

Первое ядро:

ld.param.u64    %rd3, [__cudaparm__Z7kernel0P6float4S0_S0__arg2];
add.u64         %rd4, %rd3, %rd2;
ld.global.v4.f32        {%f1,%f2,%f3,%f4}, [%rd4+0];
.loc    16      21      0
ld.param.u64    %rd5, [__cudaparm__Z7kernel0P6float4S0_S0__arg3];
add.u64         %rd6, %rd5, %rd2;
ld.global.v4.f32        {%f5,%f6,%f7,%f8}, [%rd6+0];
st.global.v4.f32        [%rd4+0], {%f1,%f2,%f3,%f4};
st.global.v4.f32        [%rd6+0], {%f5,%f6,%f7,%f8};
.loc    16      24      0
ld.param.u64    %rd7, [__cudaparm__Z7kernel0P6float4S0_S0__arg1];
add.u64         %rd8, %rd7, %rd2;
st.global.v4.f32        [%rd8+0], {%f1,%f2,%f3,%f4};

второе ядро:

ld.param.u64    %rd3, [__cudaparm__Z7kernel1P6float4S0_S0__arg2];
add.u64         %rd4, %rd3, %rd2;
ld.global.v4.f32        {%f1,%f2,%f3,%f4}, [%rd4+0];
ld.param.u64    %rd5, [__cudaparm__Z7kernel1P6float4S0_S0__arg1];
add.u64         %rd6, %rd5, %rd2;
st.global.v4.f32        [%rd6+0], {%f1,%f2,%f3,%f4};

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