CUDA中的float1与float有什么区别？

Question

CUDA中的float1与float有什么区别？

7

我注意到在cuda中有一个float1结构体类型。例如，在使用float array和float1 array时，相对于简单的float，是否存在性能优势？

struct __device_builtin__ float1
{
    float x;
};

在 float4 中，根据场合不同，存在性能优势，因为对齐方式为 4x4 字节 = 16 字节。这只是在带有 float1 参数的 __device__ 函数中特殊使用吗？

提前致谢。

- BugShotGG

1

我记得在 StackOverflow 的一篇帖子的评论中提到 __device_builtin__ 没有性能影响，但我再也找不到那篇帖子了。 - Vitality

1

我找到了这篇文章：CUDA Thrust使用double2数组进行归约。 - Vitality

2

我认为它只是为了支持开发人员进行编译器技巧，以在为不同元组大小生成多个可执行代码集时节省源代码。 - ArchaeaSoftware

@ArchaeaSoftware，您是否想在我的评论基础上进一步发表自己的看法，形成一个新的答案？这对未来的用户可能会很有用，我会点赞的。 - Vitality

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Vitality · Accepted Answer

根据@talonmies在CUDA Thrust reduction with double2 arrays的评论，我比较了使用CUDA Thrust和在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();

}