float1 vs float w CUDA

Question

Zauważyłem, że istnieje typ struktury float1 w cuda. Czy jest jakaś korzyść z wydajności w porównaniu z prostym float, na przykład w przypadku korzystania z float array kontra float1 array?

struct __device_builtin__ float1 
{ 
    float x; 
}; 

W float4 nie ma korzyści wydajności, w zależności od okazji, ponieważ ustawienie jest 4x4bytes = 16bytes. Czy jest to tylko do specjalnych zastosowań w __device__ funkcji z parametrami float1?

Z góry dziękuję.

Answer 1

Po komentarzu @talonmies do posta CUDA Thrust reduction with double2 arrays, porównałem obliczenia normy wektora używając CUDA Thrust i przełączając między float i float1. Rozważałem szereg elementów N=1000000 na karcie GT210 (cc 1.2). Wydaje się, że obliczenie normy zajmuje dokładnie tyle samo czasu w obu przypadkach, a mianowicie około 3.4s, więc nie ma poprawy wydajności. Jak wynika z poniższego kodu, być może float jest nieco wygodniejszy w użyciu niż float1.

Na koniec zauważ, że korzyść z float4 wynika z wyrównania __builtin__align__, a nie z __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(); 

}