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个线程，编号为0..255的tid试图执行每一行代码。例如，在这种情况下：

shared[tid] = values[tid]; 

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

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

对于这类琐碎问题，最简单的解决方案是使数组大小与块大小匹配

BLOCK_SIZE = N

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

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

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