0

cuda で Sauvola Binarization を実装しようとしています。このために、ホストの 2D 配列でイメージを読み取り、ピッチを使用してデバイスで 2D 配列のメモリを割り当てました。メモリを割り当てた後、ホストの 2D 配列をデバイス 2d にコピーしようとしています。 cudaMemcpy2D を使用した配列、コンパイルは正常に実行されますが、実行時にここでクラッシュします。どこが欠けているのか理解できません。親切に何か提案してください。私が書いたコードは次のとおりです。

#include "BinMain.h"
#include "Binarization.h"
#include <stdlib.h>
#include <stdio.h>
#include <conio.h>
#include <cuda.h>
#include <cuda_runtime.h>

void printDevProp(cudaDeviceProp);
void CUDA_SAFE_CALL( cudaError_t);


int main()
{
    //Read an IplImage in imgOriginal as grayscale
    IplImage * imgOriginal = cvLoadImage("E:\\1.tiff",CV_LOAD_IMAGE_GRAYSCALE);

    //Create a size variable of type CvSize for cvCreateImage Parameter
    CvSize size = cvSize(imgOriginal->width,imgOriginal->height);

    //create an image for storing the result image with same height and width as     imgOriginal
    IplImage * imgResult = cvCreateImage(size,imgOriginal->depth,imgOriginal-    >nChannels);

    //Create a 2D array for storing the pixels value of each of the pixel of imgOriginal grayscale image
    int ** arrOriginal = (int **)malloc(imgOriginal->height * sizeof(int *));
    for (int i = 0; i < imgOriginal->height; i++) 
{
    arrOriginal[i] = (int*)malloc(imgOriginal->width * sizeof(int));
}

//Create a 2D array for storing the returned device array
int ** arrReturn = (int **)malloc(imgOriginal->height * sizeof(int *));
for (int i = 0; i < imgOriginal->height; i++) 
{
    arrReturn[i] = (int*)malloc(imgOriginal->width * sizeof(int));
}

//Create a CvScalar variable to copy pixel values in 2D array (arrOriginal)
CvScalar s;

//Copying the pixl values
for(int j = 0;j<imgOriginal->height;j++)
{
    for(int k =0;k<imgOriginal->width;k++)
    {
        s = cvGet2D(imgOriginal,j,k);
        arrOriginal[j][k] = s.val[0];
    }
}

//Cuda Device Property
int devCount;
cudaGetDeviceCount(&devCount);
printf("CUDA Device Query...\n");
printf("There are %d CUDA devices.\n", devCount);

// Iterate through devices
for (int i = 0; i < devCount; ++i)
{
    // Get device properties
    printf("\nCUDA Device #%d\n", i);
    cudaDeviceProp devProp;
    cudaGetDeviceProperties(&devProp, i);
    printDevProp(devProp);
}

//Start the clock
clock_t start = clock();

//Allocating Device memory for 2D array using pitch
size_t host_orig_pitch = imgOriginal->width * sizeof(int)* imgOriginal->height; //host original array pitch in bytes
size_t dev_pitch;   //device array pitch in bytes which will be used in cudaMallocPitch
size_t dev_pitchReturn; //device return array pitch in bytes
size_t host_ret_pitch = imgOriginal->width * sizeof(int)* imgOriginal->height;  //host return array pitch in bytes

int * devArrOriginal;   //device 2d array of original image
int * result;   //device 2d array for returned array
int dynmicRange = 128; //Dynamic Range for calculating the threshold from sauvola's formula

//Allocating memory by using cudaMallocPitch
CUDA_SAFE_CALL(cudaMallocPitch((void**)&devArrOriginal,&dev_pitch,imgOriginal->width * sizeof(int),imgOriginal->height * sizeof(int)));

//Allocating memory for returned array
CUDA_SAFE_CALL(cudaMallocPitch((void**)&result,&dev_pitchReturn,imgOriginal->width * sizeof(int),imgOriginal->height * sizeof(int)));

//Copying 2D array from host memory to device mempry by using cudaMemCpy2D
CUDA_SAFE_CALL(cudaMemcpy2D((void*)devArrOriginal,dev_pitch,(void*)arrOriginal,host_orig_pitch,imgOriginal->width * sizeof(float),imgOriginal->height,cudaMemcpyHostToDevice));
    int windowSize = 19;    //Size of the window for calculating mean and variance
    //Launching the kernel by calling myKernelLauncher function.
    myKernelLauncher(devArrOriginal,result,windowSize,imgOriginal->width,imgOriginal-    >height,dev_pitch,dynmicRange);
    //Calling the sauvola binarization function by passing the parameters as 
    //1.arrOriginal 2D array 2.Original image height 3.Original image width
    //int ** result = AdaptiveBinarization(arrOriginal,imgOriginal->height,imgOriginal-    >width);//binarization(arrOriginal,imgOriginal->width,imgOriginal->height);
    //
CUDA_SAFE_CALL(cudaMemcpy2D(arrReturn,host_ret_pitch,result,dev_pitchReturn,imgOriginal->width * sizeof(int),imgOriginal->height * sizeof(int),cudaMemcpyDeviceToHost));
//create a CvScalar variable to set the data in imgResult
CvScalar ss;

//Copy the pixel values from returned array to imgResult
for(int i=0;i<imgOriginal->height;i++)
{
    for(int j=0;j<imgOriginal->width;j++)
    {
        ss = cvScalar(arrReturn[i][j]*255);
        cvSet2D(imgResult,i,j,ss);
        //k++; //No need for k if returned array is 2D
    }
}

printf("Done \n");
//calculate and print the time elapsed
printf("Time elapsed: %f\n", ((double)clock() - start) / CLOCKS_PER_SEC);

//Create a windoe and show the resule image
cvNamedWindow("Result",CV_WINDOW_AUTOSIZE);
cvShowImage("Result",imgResult);
cvWaitKey(0);
getch();

//Release the various resources
cvReleaseImage(&imgResult);
cvReleaseImage(&imgOriginal);
cvDestroyWindow("Result");
for(int i = 0; i < imgOriginal->height; i++)
    free(arrOriginal[i]);

free(arrOriginal);
free(result);
cudaFree(&devArrOriginal);
cudaFree(&result);

}

// Print device properties
void printDevProp(cudaDeviceProp devProp)
{
printf("Major revision number:         %d\n",  devProp.major);
printf("Minor revision number:         %d\n",  devProp.minor);
printf("Name:                          %s\n",  devProp.name);
printf("Total global memory:           %u\n",  devProp.totalGlobalMem);
printf("Total shared memory per block: %u\n",  devProp.sharedMemPerBlock);
printf("Total registers per block:     %d\n",  devProp.regsPerBlock);
printf("Warp size:                     %d\n",  devProp.warpSize);
printf("Maximum memory pitch:          %u\n",  devProp.memPitch);
printf("Maximum threads per block:     %d\n",  devProp.maxThreadsPerBlock);
for (int i = 0; i < 3; ++i)
printf("Maximum dimension %d of block:  %d\n", i, devProp.maxThreadsDim[i]);
for (int i = 0; i < 3; ++i)
printf("Maximum dimension %d of grid:   %d\n", i, devProp.maxGridSize[i]);
printf("Clock rate:                    %d\n",  devProp.clockRate);
printf("Total constant memory:         %u\n",  devProp.totalConstMem);
printf("Texture alignment:             %u\n",  devProp.textureAlignment);
printf("Concurrent copy and execution: %s\n",  (devProp.deviceOverlap ? "Yes" : "No"));
printf("Number of multiprocessors:     %d\n",  devProp.multiProcessorCount);
printf("Kernel execution timeout:      %s\n",  (devProp.kernelExecTimeoutEnabled ? "Yes" : "No"));
return;
}

/* Utility Macro : CUDA SAFE CALL */ 
void CUDA_SAFE_CALL( cudaError_t call) 
{ 

cudaError_t ret = call; 
switch(ret) 
{ 
    case cudaSuccess: 
         break; 
    default : 
            {
                printf(" ERROR at line :%i.%d' ' %s\n", 
                __LINE__,ret,cudaGetErrorString(ret)); 
                exit(-1); 
                break; 
            } 
} 
}

コードの流れは次のとおりです。 1. イメージからホストに 2D 配列を作成し、カーネルから返された配列用に別の配列を作成します。2. CudaMallocPitch を使用してデバイスの 2D 配列にメモリを割り当てます。 3. カーネルによって返される 2D 配列にメモリを割り当てます。4. cudaMemcpy2d を使用して、元の 2 次元配列をホストからデバイス配列にコピーします。5. カーネルを起動します。6. cudaMemcpy2D を使用して、返されたデバイス配列をホスト配列にコピーします。

4 番目のポイントに達すると、プログラムがクラッシュします。これは、「SauvolaBinarization_CUDA_OpenCV.exe の 0x773415de で未処理の例外が発生しました: 0xC0000005: アクセス違反の読み取り場所 0x01611778」という未処理の例外です。

問題はメモリの割り当て中にあるに違いないと思いますが、関数を初めて使用していて、どのように機能するかわかりません。親切に提案してください。

4

1 に答える 1

2

まず、「cudaMallocPitch」を適切に呼び出していません。「高さ」パラメーターは行数を表す必要があるため、代わりに:

imgOriginal->height * sizeof(int)

あなたは単に使用する必要があります：

imgOriginal->height

行ごとのバイト数はす​​でに「ピッチ」プロパティに含まれているため、これで問題ありません。ただし、主な問題は、ホスト イメージにメモリを割り当てる方法にあります。あなたが書くとき：

//Create a 2D array for storing the pixels value of each of the pixel of imgOriginal grayscale image
    int ** arrOriginal = (int **)malloc(imgOriginal->height * sizeof(int *));
    for (int i = 0; i < imgOriginal->height; i++) 
{
    arrOriginal[i] = (int*)malloc(imgOriginal->width * sizeof(int));
}

配列へのポインターを持つ配列を効果的に作成しています。作成している CUDA API 呼び出し:

CUDA_SAFE_CALL(cudaMemcpy2D((void*)devArrOriginal,dev_pitch,(void*)arrOriginal,host_orig_pitch,imgOriginal->width * sizeof(float),imgOriginal->height,cudaMemcpyHostToDevice));

は、入力メモリ バッファが連続していることを期待しています。入力画像の最初の行 (合計 "imgOriginal->width * sizeof(float)" バイト) がアドレスから読み取られます。

 (void*)arrOriginal

ただし、そのアドレスから始まる有効なデータの量は、「imgOriginal->height * sizeof(int *)」バイトのみです。2 つのバイト カウントが異なる可能性が非常に高く、不明な場所から読み取ることになるため、クラッシュにつながります。

これを解決するには、次のように「arrOriginal」を 1 つの連続したブロックとして割り当てることを検討してください。

int * arrOriginal = (int *)malloc(imgOriginal->height * imgOriginal->width * sizeof(int));

また、この場合、ピッチは次のようになります。

 "imgOriginal->width * sizeof(int)"
于 2012-08-28T14:08:40.343 に答える