c++ - float1 与 CUDA 中的 float

Question

我注意到 cuda 中有一个 float1 结构类型。与简单的 float 相比，是否有任何性能优势，例如，在使用 float 数组 与 float1 数组 的情况下？

struct __device_builtin__ float1
{
    float x;
};

在 float4 中有一个性能优势，这取决于场合，因为对齐是 4x4bytes = 16bytes。是否仅用于带有 float1 参数的 __device__ 函数？

提前致谢。

Answer 1

关注@talonmies 对帖子的评论 CUDA Thrust reduction with double2 arrays ，我比较了使用 CUDA Thrust 计算 vector 的范数以及在 float 和 float1 之间切换。我考虑过 GT210 卡 (cc 1.2) 上的 N=1000000 元素数组。似乎两种情况的范数计算时间完全相同，即大约 3.4s，因此没有性能提升。从下面的代码可以看出，float 可能比 float1 使用起来更舒适。

最后，请注意 float4 的优势源于对齐 __builtin__align__，而不是 __device_builtin__。

#include <thrust\device_vector.h>
#include <thrust\transform_reduce.h>

struct square
{
    __host__ __device__ float operator()(float x)
    {
        return x * x;
    }
};

struct square1
{
    __host__ __device__ float operator()(float1 x)
    {
        return x.x * x.x;
    }
};

void main() {

    const int N = 1000000;

    float time;
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    thrust::device_vector<float> d_vec(N,3.f);

    cudaEventRecord(start, 0);
    float reduction = sqrt(thrust::transform_reduce(d_vec.begin(), d_vec.end(), square(), 0.0f, thrust::plus<float>()));
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Elapsed time reduction:  %3.1f ms \n", time);

    printf("Result of reduction = %f\n",reduction);

    thrust::host_vector<float1>   h_vec1(N);
    for (int i=0; i<N; i++) h_vec1[i].x = 3.f;
    thrust::device_vector<float1> d_vec1=h_vec1;

    cudaEventRecord(start, 0);
    float reduction1 = sqrt(thrust::transform_reduce(d_vec1.begin(), d_vec1.end(), square1(), 0.0f, thrust::plus<float>()));
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Elapsed time reduction1:  %3.1f ms \n", time);

    printf("Result of reduction1 = %f\n",reduction1);

    getchar();

}