c++ - float1 vs CUDA中的float

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

struct __device_builtin__ float1
{
    float x;
};

在float4中，由于对齐方式为4x4bytes = 16bytes，因此根据情况有性能上的好处。
它仅用于带有__device__参数的float1函数中的特殊用途吗？

提前致谢。

最佳答案

在@talonmies对CUDA Thrust reduction with double2 arrays发表评论之后，我比较了使用CUDA Thrust以及在float和float1之间切换时 vector 范数的计算。我已经考虑过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();

}

关于c++ - float1 vs CUDA中的float，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/24185811/