0

在我的unary_op.operator中，我需要创建一个临时数组。
我想cudaMalloc这是要走的路。
但是，它是性能高效还是有更好的设计？

struct my_unary_op
{
    __host__ __device__ int operator()(const int& index) const
    {
        int* array;
        cudaMalloc((void**)&array, 10*sizeof(int));

        for(int i = 0; i < 10; i++)
            array[i] = index;

        int sum=0;
        for(int i=0; i < 10 ; i++)
            sum += array[i];

        return sum;
    };

};
int main()
{
    thrust::counting_iterator<int> first(0);
    thrust::counting_iterator<int> last = first+100;

    my_unary_op unary_op = my_unary_op();

    thrust::plus<int> binary_op;

    int init = 0;
    int sum = thrust::transform_reduce(first, last, unary_op, init, binary_op);

    return 0;
};
4

1 回答 1

2

您将无法cudaMalloc()在__device__函数中编译，因为它是仅主机函数。但是，您可以使用普通malloc()或new（在计算能力 >= 2.0 的设备上），但在设备上运行时这些不是很有效。有两个原因。第一个是并发运行的线程在内存分配调用期间被序列化。第二个是调用以块的形式分配全局内存，这些块的排列方式使得当内存加载和存储指令由一个扭曲中的 32 个线程运行时，它们并不相邻，因此您无法获得正确合并的内存访问。

您可以通过在__device__函数中使用固定大小的 C 样式数组（即int array[10];）来解决这两个问题。小型、固定大小的数组有时可以由编译器优化，以便将它们存储在寄存器文件中，以实现极快的访问。如果编译器将它们存储在全局内存中，它将使用本地内存。本地内存存储在全局内存中，但它以这样的方式交错，当一个 warp 中的 32 个线程运行加载或存储指令时，每个线程访问内存中的相邻位置，从而使事务能够完全合并。

如果您在运行时不知道 C 数组的大小，请在数组中分配一个最大大小并保留其中一些未使用的大小。

我认为固定大小的数组使用的内存总量将取决于 GPU 上并发处理的线程总数，而不是内核启动的线程总数。在这个答案中，@mharris 展示了如何计算最大可能的并发线程数，对于 GTX580，它是 24,576。因此，如果固定大小的数组是 16 个 32 位值，则该数组可能使用的最大内存量为 1536KiB。

如果您需要范围广泛的数组大小，您可以使用模板来编译具有多种不同大小的内核。然后，在运行时，选择一个能够容纳您需要的大小的。但是，如果您只是分配可能需要的最大值，那么内存使用量将不会成为您可以启动的线程数的限制因素。

于 2012-12-30T22:10:32.173 回答