许多GPU上设置交叉点

我正在研究GPU（CUDA）对我的大量计算任务的适合程度。

我有250万1024个元素的整数排序向量。 我需要比较所有这些。 比较是由C ++标准库的std :: set_intersection执行的标准set_intersection，并且比较是对称的。

这意味着我需要做的比较次数是（250万×250万-250万）/ 2 =〜3.125×10 12比较。

所有可用于CUDA的标准优化库只适用于两对不同大小的向量，这不是我的情况，因为我的1024向量大小不能改变，但我有大量的对比对。

到目前为止，我的所有测试都表明，在GTX 1060 GPU上，与使用SSE指令加速set_intersection的英特尔（R）Xeon®CPU E5-1620 0 @ 3.60GHz的单核相比，我只能实现10倍的加速算法。

这似乎主要是因为我的数据集很大，每个计算任务都需要8KB（2×1024个int向量）来计算结果。

我似乎无法以任何有用的方式使用CUDA共享内存，因为我的计算需要大量数据才能用于每个线程，并且很快超过了任何可用设备的CUDA共享内存。

我将不胜感激关于如何在现代GPU上处理这类问题的一些想法，特别是与大数据集相关的问题。

x10与单个CPU核心的比率是否通常可以观察到这种问题？

#include <stdio.h>
#include <iostream>
#include <chrono>

typedef struct hash_type {
int32_t data[1024];
} hash_type;

__global__
void test_kernel_without_shared_memory(int job_size, hash_type* dev_full_hash_list_ptr, uint64_t nb_of_hashes, uint8_t* result_ptr)
{
    int job_id = blockIdx.x*blockDim.x + threadIdx.x;

    if (job_id < job_size) {
        uint32_t i = (uint32_t)(nb_of_hashes+0.5-sqrtf((nb_of_hashes-0.5)*(nb_of_hashes-0.5) - 2*job_id));
        uint32_t j = i + job_id - ((i-1)*(nb_of_hashes-(double)i/2));
        i--;

        size_t a = 0, b = 0;
        size_t counter = 0;

        while(a < 1024 && b < 1024) {
            if((dev_full_hash_list_ptr+i)->data[a] < (dev_full_hash_list_ptr+j)->data[b]) {
                a++;
            } else if((dev_full_hash_list_ptr+i)->data[a] > (dev_full_hash_list_ptr+j)->data[b]) {
                b++;
            } else {
                counter++;
                a++; b++;
            }
        }
        result_ptr[job_id] = (uint8_t)(0.5 + 100*counter / (double) (2 * 1024 - counter));
    }
}

#define threads_per_block 1024
#define max_grid_dimension 255*255*255*255
uint8_t* GPU_compute_distances(hash_type* full_hash_list, uint64_t nb_hashes) {
    hash_type* dev_full_hash_list;
    cudaMalloc(&dev_full_hash_list, nb_hashes*sizeof(hash_type)); 
    cudaMemcpy(dev_full_hash_list, full_hash_list, nb_hashes*sizeof(hash_type), cudaMemcpyHostToDevice);

    uint64_t total_job_size = (nb_hashes * nb_hashes - nb_hashes)/2;
    uint8_t* dev_result_ptr;
    cudaMalloc(&dev_result_ptr, total_job_size*sizeof(uint8_t)); 

    printf("total_job_size=%lun", total_job_size);
    printf("nb_hashes=%lun", nb_hashes);

    test_kernel_without_shared_memory<<<total_job_size/threads_per_block+1, threads_per_block>>>(total_job_size, dev_full_hash_list, nb_hashes, dev_result_ptr);

    uint8_t* result_ptr;
    result_ptr = (uint8_t*)malloc(total_job_size*sizeof(uint8_t));

    cudaMemcpy(result_ptr, dev_result_ptr, total_job_size*sizeof(uint8_t), cudaMemcpyDeviceToHost);

    cudaFree(dev_result_ptr); 
    cudaFree(dev_full_hash_list); 

    return result_ptr;
}

#define test_size 10000
int main() {
    hash_type* full_hash_list;
    full_hash_list = (hash_type*)malloc(test_size*sizeof(hash_type));
    for (int i = 0; i< test_size; i++) {
        for (int j = 0; j < 100; j++) {
            full_hash_list[i].data[j] = -1;
        }
        for (int j = 100; j < 1024; j++) {
            full_hash_list[i].data[j] = i*1024+j;
        }
    }

    auto started = std::chrono::high_resolution_clock::now();

    uint8_t* result = GPU_compute_distances(full_hash_list, test_size);

    auto done = std::chrono::high_resolution_clock::now();
    std::cout << "Task duration (msec):" << std::chrono::duration_cast<std::chrono::milliseconds>(done-started).count() << "n";

    for (int i = 0; i< 10; i++) {
        std::cout << (uint32_t)result[i] << "n";
    }

    free(full_hash_list);
    free(result);

    return 0;
}

哪里

uint32_t i = (uint32_t)(nb_of_hashes+0.5-sqrtf((nb_of_hashes-0.5)*(nb_of_hashes-0.5) - 2*job_id));
uint32_t j = i + job_id - ((i-1)*(nb_of_hashes-(double)i/2));
i--;

是一个公式，它允许将位于数据上半部分矩阵中的比较job_id转换为指针偏移量。 例如，在5个矢量的示例列表中：

x 0 1 2 3
x x 4 5 6
x x x 7 8
x x x x 9
x x x x x

0 - > [0] [1]，1 - > [0] [2]等。对我的性能问题并不重要。

我第一次尝试使用（10000x10000-10000）/ 2调用thrust :: set_intersection传递哈希值。

nvprof给了我很多cudaMalloc / cudaFree调用，并且执行速度非常慢（与示例中纯粹的CUDA代码相比，速度要慢很多）。

我的第二次尝试是通过一次调用thrust :: for_each_n将指针传递给要比较的向量对。 这个实现与样本中的实现非常接近，但它没有给我实验不同内核参数的可能性。

最后和最快的是示例代码中提供的。

我还尝试将每个块的线程数减少到4，并使用共享装饰器将2个向量复制到本地向量和向量数组[4]，接着进行__syncthreads（）调用。 我只能每块启动4个线程的事实是一个没有意外的性能杀手。

因此，当前具有最佳性能的实现是将缓存管理完全留给GPU及其二级缓存。

我也尝试使用L1 + L2缓存进行编译，与L2缓存相比，性能下降约25％。

我还调查了CUDA纹理内存的可能性，但没有将它推入工作示例中，以便从我阅读的在线文档中进行测试，因为我无法找到合适的寻址模式，所以对我来说看起来不太合适。 但我可能在这里完全错了。

链接地址: http://www.djcxy.com/p/25401.html

上一篇: Many set intersections on GPU

下一篇: Distance to region