Ошибка доступа к памяти Cuda: CudaIllegalAddress, обработка изображений (стереофокусное изображение)

Question

I'm using cuda to deal with image proccessing. but my result is always get 'cudaErrorIllegalAddress : an illegal memory access was encountered' 

Что я сделал, это ниже.

Во-первых, нагрузка преобразуется изображение (RGB к серому) к устройству, я использовать 'cudaMallocPitch' и 'cudaMemcpy2D'

unsigned char *dev_srcleft; 
size_t dev_srcleftPitch 
cudaMallocPitch((void**)&dev_srcleft, &dev_srcleftPitch, COLS * sizeof(int), ROWS)); 
cudaMemcpy2D(dev_srcleft, dev_srcleftPitch, host_srcConvertL.data, host_srcConvertL.step, 
    COLS, ROWS, cudaMemcpyHostToDevice); 

И, 2D массив Выделение для хранения результата. значение результата описывается как 27 бит, поэтому я пытаюсь использовать «int», который составляет 4 байта = 32 бит, а не только для больших размеров, для производительности требуется атомная операция (atomicOr, atomicXor). , и мое устройство не поддерживает 64-битную атомную операцию.

int *dev_leftTrans; 
cudaMallocPitch((void**)&dev_leftTrans, &dev_leftTransPitch, COLS * sizeof(int), ROWS); 
cudaMemset2D(dev_leftTrans, dev_leftTransPitch, 0, COLS, ROWS); 

распределения памяти и memcpy2D прекрасно работает, и я проверить на

Mat temp_output(ROWS, COLS, 0); 
cudaMemcpy2D(temp_output.data, temp_output.step, dev_srcleft, dev_srcleftPitch, COLS, ROWS, cudaMemcpyDeviceToHost); 
imshow("temp", temp_output); 

Затем ли код ядра.

__global__ void TestKernel(unsigned char *src, size_t src_pitch, 
            int *dst, size_t dst_pitch, 
          unsigned int COLS, unsigned int ROWS) 
{ 
    const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; 
    const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; 

    unsigned char src_val = src[x + y * src_pitch]; 
    dst[x + y * dst_pitch] = src_val; 
} 

dim3 dimblock(3, 3); 
dim3 dimGrid(ceil((float)COLS/dimblock.x), ceil((float)ROWS/dimblock.y)); 
TestKernel << <dimGrid, dimblock, dimblock.x * dimblock.y * sizeof(char) >> > 
    (dev_srcleft, dev_srcleftPitch, dev_leftTrans, dev_leftTransPitch, COLS, ROWS); 

Параметр COLS and ROWS - размер изображения. Я думаю, что ошибка возникает здесь: TestKerenl.

src_val, чтение из глобальной памяти работает хорошо, но когда я пытаюсь получить доступ к DST, она взрывается с cudaErrorIllegalAddress

Я не знаю, что это неправильно, и я sufferd в течение 4-х дней. пожалуйста, помогите мне

ниже мой полный код

#include <cuda.h> 
#include <cuda_runtime.h> 
#include <cuda_runtime_api.h> 
#include <device_functions.h> 
#include <cuda_device_runtime_api.h> 
#include <device_launch_parameters.h> 
#include <math.h> 
#include <iostream> 
#include <opencv2\opencv.hpp> 
#include<string> 



#define HANDLE_ERROR(err)(HandleError(err, __FILE__, __LINE__)) 
static void HandleError(cudaError_t err, const char*file, int line) 
{ 
    if (err != cudaSuccess) 
    { 
     printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line); 
     exit(EXIT_FAILURE); 
    } 
} 
using namespace std; 
using namespace cv; 

string imagePath = "Ted"; 
string imagePathL = imagePath + "imL.png"; 
string imagePathR = imagePath + "imR.png"; 


__global__ void TestKernel(unsigned char*src, size_t src_pitch, 
         int *dst, size_t dst_pitch, 
         unsigned int COLS, unsigned int ROWS) 
{ 
    const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; 
    const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; 
    if ((COLS< x) && (ROWS < y)) return; 

    unsigned char src_val = src[x + y * src_pitch]; 
    dst[x + y * dst_pitch] = src_val; 
} 

int main(void) 
{ 
    //Print_DeviceProperty(); 
    //Left Image Load 
    Mat host_srcImgL = imread(imagePathL, CV_LOAD_IMAGE_UNCHANGED); 
    if (host_srcImgL.empty()){ cout << "Left Image Load Fail!" << endl;  return; } 
    Mat host_srcConvertL; 
    cvtColor(host_srcImgL, host_srcConvertL, CV_BGR2GRAY); 

    //Right Image Load 
    Mat host_srcImgR = imread(imagePathR, CV_LOAD_IMAGE_UNCHANGED); 
    if (host_srcImgL.empty()){ cout << "Right Image Load Fail!" << endl; return; } 
    Mat host_srcConvertR; 
    cvtColor(host_srcImgR, host_srcConvertR, CV_BGR2GRAY); 

    //Create parameters 
    unsigned int COLS = host_srcConvertL.cols; 
    unsigned int ROWS = host_srcConvertR.rows; 
    unsigned int SIZE = COLS * ROWS; 
    imshow("Left source image", host_srcConvertL); 
    imshow("Right source image", host_srcConvertR); 

    unsigned char *dev_srcleft, *dev_srcright, *dev_disp; 
    int *dev_leftTrans, *dev_rightTrans; 
    size_t dev_srcleftPitch, dev_srcrightPitch, dev_dispPitch, dev_leftTransPitch, dev_rightTransPitch; 
    cudaMallocPitch((void**)&dev_srcleft, &dev_srcleftPitch, COLS, ROWS); 
    cudaMallocPitch((void**)&dev_srcright, &dev_srcrightPitch, COLS, ROWS); 
    cudaMallocPitch((void**)&dev_disp, &dev_dispPitch, COLS, ROWS); 
    cudaMallocPitch((void**)&dev_leftTrans, &dev_leftTransPitch, COLS * sizeof(int), ROWS); 
    cudaMallocPitch((void**)&dev_rightTrans, &dev_rightTransPitch, COLS * sizeof(int), ROWS); 

    cudaMemcpy2D(dev_srcleft, dev_srcleftPitch, host_srcConvertL.data, host_srcConvertL.step, 
    COLS, ROWS, cudaMemcpyHostToDevice); 
    cudaMemcpy2D(dev_srcright, dev_srcrightPitch, host_srcConvertR.data, host_srcConvertR.step, 
    COLS, ROWS, cudaMemcpyHostToDevice); 
    cudaMemset(dev_disp, 255, dev_dispPitch * ROWS); 

    dim3 dimblock(3, 3); 
    dim3 dimGrid(ceil((float)COLS/dimblock.x), ceil((float)ROWS/dimblock.y)); 




    cudaEvent_t start, stop; 
    float elapsedtime; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 
    cudaEventRecord(start, 0); 

    TestKernel << <dimGrid, dimblock, dimblock.x * dimblock.y * sizeof(char) >> > 
    (dev_srcleft, dev_srcleftPitch, dev_leftTrans, dev_leftTransPitch, COLS, ROWS); 
    /*TestKernel << <dimGrid, dimblock, dimblock.x * dimblock.y * sizeof(char) >> > 
    (dev_srcright, dev_srcrightPitch, dev_rightTrans, dev_rightTransPitch, COLS, ROWS);*/ 
    cudaThreadSynchronize(); 

    cudaError_t res = cudaGetLastError(); 
    if (res != cudaSuccess) 
    printf("%s : %s\n", cudaGetErrorName(res), cudaGetErrorString(res)); 

    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    cudaEventElapsedTime(&elapsedtime, start, stop); 
    cudaEventDestroy(start); 
    cudaEventDestroy(stop); 
    cout << elapsedtime << "msec" << endl; 


    Mat temp_output(ROWS, COLS, 0); 
    cudaMemcpy2D((int*)temp_output.data, temp_output.step, dev_leftTrans, dev_leftTransPitch, COLS, ROWS, cudaMemcpyDeviceToHost); 
    imshow("temp", temp_output); 
    waitKey(0); 
    return 0; 
} 

И это моя среда vs2013, v6.5 CUDA Device 'отеле находится ниже

Major revision number:   3 
Minor revision number:   0 
Name:       GeForce GTX 760 (192-bit) 
Total global memory:   1610612736 
Total shared memory per block: 49152 
Total registers per block:  65536 
Warp size:      32 
Maximum memory pitch:   2147483647 
Maximum threads per block:  1024 
Maximum dimension 0 of block: 1024 
Maximum dimension 1 of block: 1024 
Maximum dimension 2 of block: 64 
Maximum dimension 0 of grid: 2147483647 
Maximum dimension 1 of grid: 65535 
Maximum dimension 2 of grid: 65535 
Clock rate:     888500 
Total constant memory:   65536 
Texture alignment:    512 
Concurrent copy and execution: Yes 
Number of multiprocessors:  6 
Kernel execution timeout:  Yes 

Answer 1

Другая проблема в вашем коде: pitch Использование для массива dst. Pitch всегда в байтах, поэтому сначала вам нужно бросить dst указатель на char*, вычислить строку смещения, а затем бросить его обратно int*:

int* dst_row = (int*)(((char*)dst) + y * dst_pitch); 
dst_row[x] = src_val; 

Answer 2

Одна проблема заключается в том, что ядро ​​Безразлично' t выполните любую проверку потока.

При определении сетки блоков, как это:

dim3 dimGrid(ceil((float)COLS/dimblock.x), ceil((float)ROWS/dimblock.y)); 

вы будете часто запуская дополнительные блоки. Причина в том, что если COLS или ROW не равномерно делятся на размеры блока (в этом случае 3), тогда вы получите дополнительные блоки, чтобы покрыть остаток в каждом случае.

Эти дополнительные блоки будут иметь потоки, которые приносят полезную работу, а некоторые из них получат доступ за пределы. Для защиты от этого, это принято ставить потокобезопасности проверку в ядре, чтобы предотвратить недоступный доступ:

const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; 
const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; 

if ((x < COLS) && (y < ROWS)) { // add this 
    unsigned char src_val = src[x + y * src_pitch]; 
    dst[x + y * dst_pitch] = src_val; 
    } // add this 

Это означает, что только нити, которые имеют действительный (в пределах поля) x и y фактически сделают доступ.

В качестве альтернативы, (3,3) не может быть особенно хорошим выбором размеров блока по соображениям производительности. Обычно рекомендуется создавать размеры блоков, продукт которых кратен 32, поэтому (32,4) или (16,16) могут быть примерами лучшего выбора.