Thrust :: sort慢，GTX960M中大小为300k的结构数组

问题描述 投票：-1回答：1

我目前正在研究GPU渲染算法，在该算法中，我需要对此结构的数组进行排序：

struct RadiosityData {
    vec4 emission;
    vec4 radiosity;
    float nPixLight;
    float nPixCam;
    float __padding[2];
};

我正在使用以下代码对数组进行排序：

thrust::device_ptr<RadiosityData> dev_ptr = thrust::device_pointer_cast(GPUpointer_ssbo);
thrust::sort(dev_ptr, dev_ptr + N);

其中GPUpointer_ssbo是来自cudaOpenGL互操作的GPU指针，N等于〜300k。比较是通过：

__host__ __device__ bool operator<(const RadiosityData& lhs, const RadiosityData& rhs) { return (lhs.nPixCam > rhs.nPixCam); };

我的GTX960M上的排序非常慢：不进行排序，我的应用每帧执行〜10ms，而进行排序大约需要35ms。这意味着排序大约需要25ms。我正在使用VS-NSIGHT测量执行时间

我知道这个问题可能是GPU同步问题，因为我在调用推力之前正在执行OpenGL操作。但是，我不相信该参数，因为如果我使用未排序的数组在OpenGL中显示数据，则仍然需要10ms的时间，这意味着OpenGL代码本身没有同步问题。

这种性能对于这种“小型”阵列是否值得期待？是否有更好的GPU排序算法可用于此类问题？

------------编辑：我正在使用默认的VS2019 CUDA命令在发行版中进行编译，该命令是：

Driver API（NVCC编译类型为.cubin，.gpu或.ptx）设置CUDAFE_FLAGS =-sdk_dir“ C：\ Program Files（x86）\ Windows Kits \ 10 \”“ C：\ Program Files \ NVIDIA GPU计算工具包\ CUDA \ v10.2 \ bin \ nvcc.exe” --use-local-env -ccbin“ C：\ Program Files（x86）\ Microsoft Visual Studio \ 2019 \ Community \ VC \ Tools \ MSVC \ 14.26.28801 \ bin \ HostX86 \ x64“ -x cu --keep-dir x64 \ Release -maxrregcount = 0 --machine 64 --compile -cudart static -o x64 \ Release \ sortBufferCUDA。 cu.obj“ C：\ Users \ Jose \ Desktop \ RealTimeDiffuseIlumination \ OpenGL-avanzado \ sortBufferCUDA.cu”

运行时API（NVCC编译类型是混合对象或.c文件）设置CUDAFE_FLAGS =-sdk_dir“ C：\ Program Files（x86）\ Windows Kits \ 10 \”“ C：\ Program Files \ NVIDIA GPU计算工具包\ CUDA \ v10.2 \ bin \ nvcc.exe” --use-local-env -ccbin“ C：\ Program Files（x86）\ Microsoft Visual Studio \ 2019 \ Community \ VC \ Tools \ MSVC \ 14.26.28801 \ bin \ HostX86 \ x64“ -x cu --keep-dir x64 \ Release -maxrregcount = 0 --machine 64 --compile -cudart static -Xcompiler” / EHsc / nologo / Fd / FS / Zi“ -o x64 \ Release \ sortBufferCUDA.cu.obj” C：\ Users \ Jose \ Desktop \ RealTimeDiffuseIlumination \ OpenGL-avanzado \ sortBufferCUDA.cu“

--------------编辑2：

以下是一个最小的工作示例：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#include <thrust/device_ptr.h>
#include <thrust/sort.h>
#include <thrust/execution_policy.h>
#include <thrust/extrema.h>
#include <cuda_runtime_api.h>
#include <cuda.h>
#include <thrust/device_vector.h>


struct RadiosityData {
    float emission[4];
    float radiosity[4];
    float nPixLight;
    float nPixCam;
    float __padding[2];
};

extern "C" void CUDAsort();

__host__ __device__ bool operator<(const RadiosityData& lhs, const RadiosityData& rhs) { return (lhs.nPixCam > rhs.nPixCam); };

int pri = 1;

thrust::device_vector<RadiosityData> dev;

void CUDAsort() {
    if (pri == 1) {
        pri = 0;
        dev.resize(300000);

    }
    thrust::sort(dev.begin(), dev.end());

}

int main()
{
    float time;
    cudaEvent_t start, stop;


    while (true) {
        cudaEventCreate(&start);
        cudaEventCreate(&stop);
        cudaEventRecord(start, 0);
        CUDAsort();

        cudaEventRecord(stop, 0);
        cudaEventSynchronize(stop);
        cudaEventElapsedTime(&time, start, stop);
        printf("Time to generate:  %3.1f ms \n", time);
    }
    return 0;
}
opengl cuda thrust
1个回答
0
投票

肯定会尝试在排序时四处移动48字节结构，但可能不是最有效的方法。

我们可以尝试的是：

  1. 将用于排序的浮点值从结构数组（AoS）中拉出到float数组中
  2. 创建索引数组以与此0 1 2 3 ...一起使用
  3. sort_by_key浮点值，带有整数索引
  4. 使用重新排列的索引数组，从输入到输出执行AoS的单个置换副本
  5. 将输出数组复制回输入数组，以模拟“就地”排序

这看起来需要很多工作，但是根据我的测试，它实际上要快一些：

$ cat t30.cu
#include <thrust/sort.h>
#include <thrust/device_vector.h>
#include <iostream>
#include <thrust/execution_policy.h>
#include <time.h>
#include <sys/time.h>
#include <cstdlib>
#define USECPSEC 1000000ULL
long long dtime_usec(unsigned long long start){

  timeval tv;
  gettimeofday(&tv, 0);
  return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}

struct RadiosityData {
#ifdef USE_VEC
 float4 emission;
 float4 radiosity;
#else
 float emission[4];
 float radiosity[4];
#endif
        float nPixLight;
 float nPixCam;
        float __padding[2];
};

__global__ void copyKernel(RadiosityData *d, float *f, int *i, int n){
 int idx=threadIdx.x+blockDim.x*blockIdx.x;
 if (idx < n){
  f[idx] = d[idx].nPixCam;
  i[idx] = idx;}
}

__host__ __device__ bool operator<(const RadiosityData &lhs, const RadiosityData  &rhs) { return (lhs.nPixCam > rhs.nPixCam); };
struct my_sort_functor
{
 template <typename T1, typename T2>
 __host__ __device__ bool operator()(T1 lhs, T2 rhs) { return (lhs.nPixCam > rhs.nPixCam); };
};
const int N = 300000;
int main(){
 RadiosityData *GPUpointer_ssbo, *o;
 int sz = N*sizeof(RadiosityData);
 thrust::device_vector<RadiosityData> ii(N);
 GPUpointer_ssbo = thrust::raw_pointer_cast(ii.data());
 thrust::device_ptr<RadiosityData> dev_ptr = thrust::device_pointer_cast(GPUpointer_ssbo);
//method 1: ordinary thrust sort
 long long dt = dtime_usec(0);
 thrust::sort(dev_ptr, dev_ptr+N);
 cudaDeviceSynchronize();
 dt = dtime_usec(dt);
 std::cout << "ordinary sort time: " << dt/(float)USECPSEC << "s" << std::endl;
//method 2: reduced sort-and-copy
 cudaMalloc(&o, sz);
 thrust::device_ptr<RadiosityData> dev_optr = thrust::device_pointer_cast(o);
 for (int i = 0; i < N; i++) {RadiosityData q{0}; q.nPixCam = rand(); ii[i] = q;}
 float *d;
 int *k;
 cudaMalloc(&d, N*sizeof(float));
 cudaMalloc(&k, N*sizeof(int));
 thrust::device_ptr<int> dev_kptr = thrust::device_pointer_cast(k);
 cudaDeviceSynchronize();
 dt = dtime_usec(0);
 copyKernel<<<(N+511)/512, 512>>>(GPUpointer_ssbo, d, k, N);
 thrust::sort_by_key(thrust::device, d, d+N, k);
 thrust::copy(thrust::make_permutation_iterator(dev_ptr, dev_kptr), thrust::make_permutation_iterator(dev_ptr, dev_kptr+N), dev_optr);
 cudaMemcpy(GPUpointer_ssbo, o, sz, cudaMemcpyDeviceToDevice);
 cudaDeviceSynchronize();
 dt = dtime_usec(dt);
 std::cout << "sort+copy time: " << dt/(float)USECPSEC << "s" << std::endl;
}

$ nvcc -o t30 t30.cu -arch=sm_52
$ ./t30
ordinary sort time: 0.009527s
sort+copy time: 0.003143s
$ nvcc -o t30 t30.cu -arch=sm_52 -DUSE_VEC
$ ./t30
ordinary sort time: 0.004409s
sort+copy time: 0.002859s
$

（（CUDA 10.1.105，GTX960，fedora核心29）

因此，使用改进的方法，我们观察到加速了大约50％。

[还观察到，当用矢量类型（float）而不是4元素数组表示4- float4数量时，通过推力可以更有效地处理AoS。我怀疑这会使编译器识别出更有效的结构复制方法。

[另外请注意，根据我的测试，为正确的GPU架构（在我的情况下为sm_52）进行编译似乎有所改进。 YMMV。

我不主张此代码或我发布的任何其他代码的正确性。任何使用我发布的代码的人均需自担风险。我仅声称自己已尝试解决原始帖子中的问题，并提供了一些解释。我并不是说我的代码没有缺陷，也不适合任何特定目的。使用（或不使用）后果自负。

最新问题
© www.soinside.com 2019 - 2024. All rights reserved.