Нормализуйте набор векторов с помощью библиотеки Thrust от Nvidia - PullRequest
Новая
       47

Нормализуйте набор векторов с помощью библиотеки Thrust от Nvidia

2 голосов
/ 20 апреля 2011

Я только что узнал о библиотеке тяги Nvidia.Просто, чтобы попробовать, он написал небольшой пример, который должен нормализовать кучу векторов. 

#include <cstdio>

#include <thrust/transform.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>

struct normalize_functor: public thrust::unary_function<double4, double4>
{
    __device__ __host__ double4 operator()(double4 v)
    {
        double len = sqrt(v.x*v.x + v.y*v.y + v.z*v.z);
        v.x /= len;
        v.y /= len;
        v.z /= len;
        printf("%f %f %f\n", v.x, v.y, v.z);
    }
};

int main()
{
    thrust::host_vector<double4> v(2);
    v[0].x = 1; v[0].y = 2; v[0].z = 3;
    v[1].x = 4; v[1].y = 5; v[1].z = 6;

    thrust::device_vector<double4> v_d = v; 
    thrust::for_each(v_d.begin(), v_d.end(), normalize_functor());

    // This doesn't seem to copy back
    v = v_d;

    // Neither this does..
    thrust::host_vector<double4> result = v_d;

    for(int i=0; i<v.size(); i++)
        printf("[ %f %f %f ]\n", result[i].x, result[i].y, result[i].z);

    return 0;
}

Кажется, что приведенный выше пример работает, однако я не могу скопировать данные обратно .. Я подумал, что простойназначение будет вызывать cudaMemcpy.Он работает для копирования данных с хоста на устройство, но не обратно ???

Во-вторых, я не уверен, правильно ли я это делаю.Документация for_each гласит:

for_each применяет функциональный объект f к каждому элементу в диапазоне [first, last);Возвращаемое значение f, если оно есть, игнорируется.

Но шаблон структуры unary_function ожидает два аргумента шаблона (один для возвращаемого значения) и заставляет operator () также возвращать значение, это приводит кпредупреждение при компиляции.Я не понимаю, как я должен написать унарный функтор без возвращаемого значения.

Далее следует расположение данных.Я просто выбрал double4, так как это приведет к двум инструкциям извлечения ld.v2.f64 и ld.f64 IIRC.Однако мне интересно, как Thrust извлекает данные изнутри (и сколько потоков / блоков cuda).Если бы я выбрал вместо этого структуру из 4 векторов, он мог бы извлекать данные объединенным способом.

Наконец, тяга обеспечивает кортежи.Как насчет массива кортежей?Как будут расположены данные в этом случае.

Я просмотрел примеры, но я не нашел пример, который объясняет, какую структуру данных выбрать для набора векторов (пример dot_products_with_zip.cu что-то говорито "структуре массивов" вместо "массивов структур", но я не вижу структур, используемых в примере.

Обновление

Я исправил приведенный выше код и попыталсязапустим более крупный пример, на этот раз нормализуя векторы по 10 тыс. 

#include <cstdio>

#include <thrust/transform.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>

struct normalize_functor
{
    __device__ __host__ void operator()(double4& v)
    {
        double len = sqrt(v.x*v.x + v.y*v.y + v.z*v.z);
        v.x /= len;
        v.y /= len;
        v.z /= len;
    }
};

int main()
{
    int n = 10000;
    thrust::host_vector<double4> v(n);
    for(int i=0; i<n; i++) {
        v[i].x = rand();
        v[i].y = rand();
        v[i].z = rand();
    }

    thrust::device_vector<double4> v_d = v;

    thrust::for_each(v_d.begin(), v_d.end(), normalize_functor());

    v = v_d;

    return 0;
}

Профилирование с помощью computeprof показывает низкую занятость и не слитый доступ к памяти: 

Kernel Occupancy Analysis

Kernel details : Grid size: 23 x 1 x 1, Block size: 448 x 1 x 1
Register Ratio      = 0.984375  ( 32256 / 32768 ) [24 registers per thread] 
Shared Memory Ratio     = 0 ( 0 / 49152 ) [0 bytes per Block] 
Active Blocks per SM        = 3 / 8
Active threads per SM       = 1344 / 1536
Potential Occupancy     = 0.875  ( 42 / 48 )
Max achieved occupancy  = 0.583333  (on 9 SMs)
Min achieved occupancy  = 0.291667  (on 5 SMs)
Occupancy limiting factor   = Block-Size

Memory Throughput Analysis for kernel launch_closure_by_value on device GeForce GTX 470

Kernel requested global memory read throughput(GB/s): 29.21
Kernel requested global memory write throughput(GB/s): 17.52
Kernel requested global memory throughput(GB/s): 46.73
L1 cache read throughput(GB/s): 100.40
L1 cache global hit ratio (%): 48.15
Texture cache memory throughput(GB/s): 0.00
Texture cache hit rate(%): 0.00
L2 cache texture memory read throughput(GB/s): 0.00
L2 cache global memory read throughput(GB/s): 42.44
L2 cache global memory write throughput(GB/s): 46.73
L2 cache global memory throughput(GB/s): 89.17
L2 cache read hit ratio(%): 88.86
L2 cache write hit ratio(%): 3.09
Local memory bus traffic(%): 0.00
Global memory excess load(%): 31.18
Global memory excess store(%): 62.50
Achieved global memory read throughput(GB/s): 4.73
Achieved global memory write throughput(GB/s): 45.29
Achieved global memory throughput(GB/s): 50.01
Peak global memory throughput(GB/s): 133.92

Интересно, как я могу оптимизироватьэто?

Ответы [ 2 ]

4 голосов
/ 21 апреля 2011

Если вы хотите изменить последовательность на месте с помощью for_each, вам нужно взять аргумент по ссылке в функторе: 

struct normalize_functor
{
    __device__ __host__ void operator()(double4& ref)
    {
        double v = ref;
        double len = sqrt(v.x*v.x + v.y*v.y + v.z*v.z);
        v.x /= len;
        v.y /= len;
        v.z /= len;
        printf("%f %f %f\n", v.x, v.y, v.z);
        ref = v;
    }
};

В качестве альтернативы вы можете использовать определение normalize_functor с алгоритмом transform, указав v_d в качестве исходного и целевого диапазона: 

thrust::transform(v_d.begin(), v_d.end(), v_d.begin(), normalize_functor());

Мое личное предпочтение - использовать transform в этой ситуации, но производительность должна быть одинаковой в любом случае.

1 голос
/ 21 апреля 2011

Что касается оптимизации, с Thrust мало что можно сделать - это не совсем задумано библиотеками.Не желая говорить за Натана Белла, который является одним из авторов Thrust и который уже писал в этой теме, цель состоит в том, чтобы сделать ряд параллельных алгоритмов данных для GPU доступным простым, интуитивно понятным способом без необходимости писать многоесли есть, код CUDA.И это, на мой взгляд, впечатляет.Производительность ядра многих базовых ядер близка к уровню техники, но всегда есть оптимизации, которые могут быть выполнены в конкретных случаях, которые нелегко сделать в общем шаблонном коде.Это часть цены, которую вы платите за простоту использования и гибкость, которые предоставляет Thrust.

Сказав это, я подозреваю, что есть несколько настроек, которые можно попробовать в функции оператора, которые могут улучшить ситуацию.Я обычно писал бы что-то вроде этого: 

struct normalize_functor
{
    __device__ __host__ void operator()(double4& v)
    {
        double4 nv = v;
        double len = sqrt(nv.x*nv.x + nv.y*nv.y + nv.z*nv.z);
        nv.x /= len;
        nv.y /= len;
        nv.z /= len;
        (void)nv.h;
        v = nv;
    };
};

Теперь, хотя это не так красиво, как оригинал, он должен гарантировать, что компилятор испускает векторизованные инструкции загрузки и сохранения.В прошлом я видел случаи, когда компилятор оптимизировал загрузку и хранение неиспользуемых членов векторных типов, что приводило к генерации PTX, генерирующей скалярную нагрузку и запоминающие устройства, и в результате нарушало объединение.Имея чистую загрузку и сохранение float4 и убедившись, что используется каждый элемент структуры, он может обойти эту нежелательную «оптимизацию», которая присутствовала как минимум в выпусках 2.x и 3.x nvcc.Я не уверен, так ли это до сих пор с компилятором 4.0.

...