Cuda - 大阵列初始化

Question

为gpu初始化大量整数的最佳方法(有效)是什么？我需要为前两个元素分配1,为其他元素分配0(对于Eratosthenes的Sieve).

1. cudaMemcpy
2. cudaMemset +在内核中设置2个第一个元素的值
3. 初始化直接在内核中
4. 别的

注意:数组大小是动态的(n作为参数传递).

我目前的版本:

int array = (int*) malloc(array_size);
array[0] = 1;
array[1] = 1;
for (int i = 2; i < n; i++) {
    array[i] = 0;
}
HANDLE_ERROR(cudaMemcpy(dev_array, array, array_size, cudaMemcpyHostToDevice));
kernel<<<10, 10>>>(dev_array);

我要感谢一个例子.

Answer 1

一种可能性是__device__通过在文件范围(即任何函数之外)添加以下声明来直接初始化GPU上的数组(如果它具有恒定大小):

__device__ int dev_array[SIZE] = {1, 1};

其余的元素将用零初始化(您可以检查PTX程序集以确保这一点).

那么,它可以在内核中使用:

__global__ void kernel(void)
{
    int tid = ...;
    int elem = dev_array[tid];
    ...
}

如果尺寸可变,您可以结合cudaMalloc()使用cudaMemset():

int array_size = ...;
int *dev_array;

cudaMalloc((void **) &dev_array, array_size * sizeof(int));
cudaMemset(dev_array, 0, array_size * sizeof(int));

然后将前两个元素设置为:

int helper_array[2] = {1, 1};
cudaMemcpy(dev_array, helper_array, 2 * sizeof(int), cudaMemcpyHostToDevice);

从计算能力2.0开始,您还可以通过malloc()设备功能直接在内核中分配整个数组:

__global__ void kernel(int array_size)
{
    int *dev_array;
    int tid = ...;

    if (tid == 0) {
        dev_array = (int *) malloc(array_size * sizeof(int));
        if (dev_array == NULL) {
            ...
        }
        memset(dev_array, 0, array_size * sizeof(int));
        dev_array[0] = dev_array[1] = 1;  
    }
    __syncthreads();

    ...
}

请注意,来自不同块的线程不知道屏障同步.

从CUDA C编程指南:

CUDA内核malloc()函数至少size从设备堆中分配字节,并返回指向已分配内存的指针,或者NULL如果存在的内存 不足以满足请求.保证返回的指针与16字节边界对齐.

不幸的是,该calloc()功能没有实现,因此无论如何你都需要设置它.分配的内存具有CUDA上下文的生命周期,但您可以随时free()从此内核或后续内核中明确调用:

由给定CUDA线程分配的内存malloc()仍然在CUDA上下文的生命周期内分配,或直到通过调用显式释放free().它可以被任何其他CUDA线程使用,甚至可以在后续内核启动时使用.

尽管如此,我并不介意补充cudaMemcpy(),因为它只需要复制两个元素,它可能需要不到总执行时间的0.01%(很容易分析).选择使您的代码清晰的任何方式.否则,这是一个不成熟的优化.