PyCUDA LogicError：cuModuleLoadDataEx 失败：遇到非法内存访问

Question

我正在尝试使用 pycuda 并行化双音排序。为此，我使用 SourceModule 和并行双音排序的 C 代码。对于内存副本管理，我使用 pycuda.driver 的 InOut 来简化一些内存传输

import pycuda.autoinit
import pycuda.driver as drv
from pycuda.compiler import SourceModule
from pycuda import gpuarray
import numpy as np
from time import time

ker = SourceModule(
    """
    __device__ void swap(int & a, int & b){
        int tmp = a;
        a = b;
        b = tmp;
    }

    __global__ void bitonicSort(int * values, int N){
        extern __shared__ int shared[];
        int tid = threadIdx.x + blockDim.x * blockIdx.x;
        // Copy input to shared mem.
        shared[tid] = values[tid];
        __syncthreads();
        // Parallel bitonic sort.
        for (int k = 2; k <= N; k *= 2){
            // Bitonic merge:
            for (int j = k / 2; j>0; j /= 2){
                int ixj = tid ^ j;
                if (ixj > tid){
                    if ((tid & k) == 0){
                        //Sort ascending
                        if (shared[tid] > shared[ixj]){
                            swap(shared[tid], shared[ixj]);
                        }
                    }
                    else{
                        //Sort descending
                        if (shared[tid] < shared[ixj]){
                            swap(shared[tid], shared[ixj]);
                        }
                    }
                }
                __syncthreads();
            }
        }
        values[tid] = shared[tid];
    }
    """
)

N = 8 #lenght of A
A = np.int32(np.random.randint(1, 20, N)) #random numbers in A
BLOCK_SIZE = 256
NUM_BLOCKS = (N + BLOCK_SIZE-1)//BLOCK_SIZE
bitonicSort = ker.get_function("bitonicSort")
t1 = time()
bitonicSort(drv.InOut(A), np.int32(N), block=(BLOCK_SIZE,1,1), grid=(NUM_BLOCKS,1), shared=4*N)
t2 = time()
print("Execution Time {0}".format(t2 - t1))
print(A)

在内核中我使用extern __shared__，在pycuda 中我使用带有相应4*N 的共享参数。也可以尝试在内核中使用__shared__ int shared[N]，但它也不起作用（在这里查看：Getting started with shared memory on PyCUDA）

在 Google Collab 中运行我收到以下错误：

/usr/local/lib/python3.6/dist-packages/pycuda/compiler.py in __init__(self, source, nvcc, options, keep, no_extern_c, arch, code, cache_dir, include_dirs)
    292 
    293         from pycuda.driver import module_from_buffer
--> 294         self.module = module_from_buffer(cubin)
    295 
    296         self._bind_module()

LogicError: cuModuleLoadDataEx failed: an illegal memory access was encountered

有谁知道是什么导致了这个错误？

Answer 1

您的设备代码未正确考虑数组的大小。

您将在一个块中启动 256 个线程。这意味着您将有 256 个线程，tid 编号为 0..255，试图执行每一行代码。例如，在这种情况下：

shared[tid] = values[tid]; 

例如，您将有一个线程尝试执行shared[255] = values[255];

您的shared 和values 数组都没有那么大。这就是非法内存访问错误的原因。

这类琐碎问题的最简单解决方案是让您的数组大小与您的块大小相匹配。

BLOCK_SIZE = N

根据我的测试，该更改会清除所有错误并生成正确排序的数组。

它不适用于大于 1024 的 N 或多块使用，但无论如何您的代码都必须修改为多块排序。

如果您在进行更改后仍然遇到问题，我建议您重新启动您的 python 会话或 colab 会话。