在CUDA中float1 vs float

Question

我注意到float1cuda 中有一个结构类型.是否有任何性能优势超过简单float,例如,在使用float arrayvs float1 array？

struct __device_builtin__ float1
{
    float x;
};

在float4有性能优势,根据场合,由于取向是4x4bytes = 16字节.它只是用于__device__带float1参数的函数的特殊用途吗？

提前致谢.

Answer 1

根据 @talonmies 对CUDA Thrust reduction with double2 arrays帖子的评论，我比较了使用 CUDA Thrust 的向量范数的计算以及在float和之间切换float1。N=1000000我考虑了GT210 卡 (cc 1.2) 上的一系列元素。看起来这两种情况下范数的计算花费了完全相同的时间，即大约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();

}