Size and alignment of 3-component built-in vector types are incompatible with CUDA

[ex-rzr]

HIP 3-component vectors have in fact same sizes and alignments as 4-component vectors.

CUDA 3-component vectors are packed:
sizeof(T3) = 3 * sizeof(T) and alignof(T3) = alignof(T)

See https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#built-in-vector-types

    std::cout << "sizeof(float1): " << sizeof(float1) << '\n';
    std::cout << "sizeof(float2): " << sizeof(float2) << '\n';
    std::cout << "sizeof(float3): " << sizeof(float3) << '\n';
    std::cout << "sizeof(float4): " << sizeof(float4) << '\n';

    std::cout << "alignof(float1): " << alignof(float1) << '\n';
    std::cout << "alignof(float2): " << alignof(float2) << '\n';
    std::cout << "alignof(float3): " << alignof(float3) << '\n';
    std::cout << "alignof(float4): " << alignof(float4) << '\n';

HIP (hcc):

sizeof(float1): 4
sizeof(float2): 8
sizeof(float3): 16 <<<<<<<<<<<<<<<
sizeof(float4): 16

alignof(float1): 4
alignof(float2): 8
alignof(float3): 16 <<<<<<<<<<<<<<<
alignof(float4): 16

CUDA

sizeof(float1): 4
sizeof(float2): 8
sizeof(float3): 12 <<<<<<<<<<<<<<<
sizeof(float4): 16

alignof(float1): 4
alignof(float2): 8
alignof(float3): 4 <<<<<<<<<<<<<<<
alignof(float4): 16

If this is correct behavior then it should be mentioned in the docs (https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_kernel_language.md#short-vector-types).

[JosephGeoBenjamin]

Hi,
This issue causes serious bug in OpenCV library; Since CV uses short vectors of uchar for 3-channel images, the results are not as expected, on moving data from one data structure to other on HIP-backend on AMD.

Can anyone point out workarounds or fixes related to this issue ?

[ex-rzr]

I use this workaround:

template <> __forceinline__ __device__
void gather_force_store<float3>(const float fx, const float fy, const float fz,
        const int stride, const int pos,
        float3* force) {
  // Store into non-strided "float3" array
#if defined(__HIP_PLATFORM_HCC__)
  // Workaround: unlike CUDA, HIP-hcc has sizeof(float3) != sizeof(CudaForce) (and == sizeof(float4))
  // TODO-HIP: Remove when https://github.com/ROCm-Developer-Tools/HIP/issues/706 is fixed
  reinterpret_cast<float*>(force)[pos * 3 + 0] = fx;
  reinterpret_cast<float*>(force)[pos * 3 + 1] = fy;
  reinterpret_cast<float*>(force)[pos * 3 + 2] = fz;
#else
  force[pos].x = fx;
  force[pos].y = fy;
  force[pos].z = fz;
#endif
}

(not sure that it can be applied to your situation)

[JosephGeoBenjamin]

@ex-rzr Thanks for your response but I guess I don't have access to that repo. (gives 404 error)

[ex-rzr]

oh, sorry. The previous message has been updated.

[NEELMCW]

@ex-rzr Hi Anton

could you provide us access to view NAMD repository ?

[ex-rzr]

@NEELMCW I don't have permissions to change the repo's settings.
But I posted the code with the workaround I use there.

[xuhuisheng]

This issue had solved on ROCm-5.1.0.
ROCm-Developer-Tools/hipamd@e27f645