Question

Я пытаюсь применить фильтр Гаусса к изображению, используя CUDA.

int main() {

    double GKernel[5][5];
    FilterCreation(GKernel);
    double * kernel = new double[25];
    int i,j,k = 0;
    for( int i = 0; i < 5; i++){
     for( int j = 0; j < 5; j++){
       kernel[k++] = GKernel[i][j];
     }
    }

    double * deviceKernel;

    cudaMalloc((void **)&deviceKernel, 25 * sizeof(double));
    cudaMemcpy(deviceKernel, kernel, 25 * sizeof(double), cudaMemcpyHostToDevice);

    Info dat = readBMP("konik.bmp");
    unsigned char * devPtr;
    unsigned char * devPtrFilter;
    size_t pitch;

    unsigned char * test= new unsigned char  [dat.size - dat.offset ];

    cudaMalloc (( void **)& devPtr , dat.size * sizeof ( unsigned char  ));
    cudaMalloc (( void **)& devPtrFilter , dat.size * sizeof ( unsigned char  ));

    cudaMemcpy ( devPtr , dat.data , sizeof ( unsigned char ) *  dat.size , cudaMemcpyHostToDevice );

        greyScale<<<dat.height,dat.width>>>(devPtr,dat.height,dat.width);
        CreateGaussFilter<<<dat.height,dat.width>>>(devPtr,devPtrFilter,5,5,deviceKernel,dat.height,dat.width);

    cudaMemcpy ( test, devPtrFilter , sizeof ( unsigned char )  *  dat.size ,cudaMemcpyDeviceToHost );
    cudaDeviceSynchronize ();

    ofstream fout;
    fout.open("output.bmp", ios::binary | ios::out);
    fout.write( reinterpret_cast<char *>(dat.info), dat.offset);

    fout.write( reinterpret_cast<char *>(test), dat.size - dat.offset );
    fout.close();
    return 0;
}

Сначала я вызываю ядро greyScale, которое работает как надо.

__global__ void greyScale( unsigned char * src , int rows, int cols){

        int i = blockDim.x * blockIdx.x * threadIdx.x;
        int j = blockDim.y * blockIdx.y * threadIdx.y;

        if( i >= rows || j >= cols ) {
                return;
        }
            unsigned char r = src[3 * (i * cols + j)];
            unsigned char g = src[3 * (i * cols + j) + 1];
            unsigned char b = src[3 * (i * cols + j) + 2];

            unsigned char linearIntensity = (unsigned char)(0.2126f * r + 0.7512f * g + 0);

            src[3 * (i * cols + j)] = linearIntensity;
            src[3 * (i * cols + j) + 1] = linearIntensity;
            src[3 * (i * cols + j) + 2] = linearIntensity;

}

Но меня беспокоит мояФункция ядра CreateGaussFilter:

__global__ void CreateGaussFilter(unsigned char * src, unsigned char * dst, int kernalHeight, int kernalWidth, double *kernalArray, int rows, int cols){
    int verticleImageBound=(kernalHeight-1)/2;
    int horizontalImageBound=(kernalWidth-1)/2;

    int row = ( blockDim.x * blockIdx.x * threadIdx.x ) + verticleImageBound;
    int col = ( blockDim.y * blockIdx.y * threadIdx.y ) + horizontalImageBound;

    if ( row >= rows - verticleImageBound || col >= cols - horizontalImageBound ) return;
            float  value=0;
            for(int kRow=0;kRow<kernalHeight;kRow++){
                  for(int kCol=0;kCol<kernalWidth;kCol++){
                  float pixel = src[ 3 * ((kRow+row-verticleImageBound ) * cols + (kCol+col-horizontalImageBound)) ] * kernalArray[kRow + kCol * kernalWidth];
                  value+=pixel;
             }
            }

            printf("value = %i\n",value);

            dst[3 * ( row * cols + col )] = round(value);
            dst[3 * ( row * cols + col ) + 1] = round(value);
            dst[3 * ( row * cols + col ) + 2] = round(value);

}

фактически выводит черное изображение (это работает, когда я запускаю его на CPU без CUDA), после отладки с printf кажется, что код не будет выполняться после цикла for, я неконечно, границы работают не в CUDA-версии, а размер блока одинаков для работающего ядра greyScale. Почему код не выполняется после цикла for? Произошла ли ошибка где-то и закрыла поток? Я в замешательстве.

Код для загрузки изображений:

struct Info{
    int width;
    int height;
    int offset;
    unsigned char * info;
    unsigned char * data;

    int size;
};

Info readBMP(char* filename)
{
    int i;
    std::ifstream is(filename, std::ifstream::binary);
    is.seekg(0, is.end);
    i = is.tellg();
    is.seekg(0);
    unsigned char *info = new unsigned char[i];
    is.read((char *)info,i);

    int width = *(int*)&info[18];
    int height = *(int*)&info[22];
    int offset = *(int*)&info[10];

    unsigned char a[offset];
    unsigned char *b = new unsigned char[i - offset];
    std::copy(info,
              info + offset,
              a);

    std::copy(info + offset,
              info + i,
              b + 0);

    Info dat;
    dat.width = width;
    dat.height = height;
    dat.offset = offset;
    dat.size = i;
    dat.info = new unsigned char[offset - 1];
    dat.data = new unsigned char[i - offset + 1];

    for( int j = 0; j < offset ; j++ ){
        dat.info[j] = a[j];
    }

    for( int j = 0; j < i - offset; j++ ){
        dat.data[j] = b[j];
    }
    return dat;

}

Спасибо за помощь!

Robert Crovella · Answer 1 · 01 декабря 2019

Я сомневаюсь, что любое из ваших ядер работает правильно. У вас есть как минимум 2 проблемы:

Это не то, как создать правильный индекс:
```
blockDim.x * blockIdx.x * threadIdx.x
```
правильный индекс выглядит следующим образом:
```
blockDim.x * blockIdx.x + threadIdx.x
```
эта ошибка очевидна для обоих индексов .x и .y в обоих ядрах.
Вы не уверены в синтаксисе запуска ядра CUDA, например, здесь: CreateGaussFilter<<<dat.height,dat.width>>>,Первый аргумент <<<...>>> - это количество блоков в сетке. Второе - это количество потоков в блоке. Если вы передадите скалярные величины для обоих этих аргументов (что вы делаете), вы получите одномерную сетку одномерных потоковых блоков. 1D здесь означает, что в ядре ваши значения индекса .y всегда будут равны нулю, поэтому это утверждение: int col = ( blockDim.y * blockIdx.y * threadIdx.y ) + horizontalImageBound; приведет к тому, что каждый поток в вашем ядре будет иметь значение col horizontalImageBound
в вашем утверждении в ядре printf, %i не является правильным параметром формата для количества float.

Вы должны изучить любой CUDA 2Dкод ядра для правильного использования. Необходимо внести изменения как в код вашего хоста, так и в код ядра.

Несколько других замечаний.

Хорошо, если вы не снимаете заголовки с вашегокод использует. Некоторые люди, пытающиеся помочь вам, захотят запустить ваш код . Сделать это легко для них (если вы хотите от них помощи). Просто мое предложение, как и весь этот мой пост.
Ожидается, что вы предоставите полный пример. Смотрите пункт 1 здесь . Например, ваш опубликованный код нигде не указывает определение FilterCreation. И у меня нет вашего konik.bmp, поэтому либо укажите, как я могу его получить, либо, что еще лучше, напишите свой код , который вы публикуете здесь таким образом, чтобы он не зависел от внешнего файла,Например, создайте фиктивное изображение в коде и пропустите процесс загрузки файла.
Это не должно быть так сложно. Возьмите то, что вы опубликовали, и создайте новый проект только с этим кодом. Это компилируется? Если нет, продолжайте добавлять к своей публикации, пока она не скомпилируется. Тогда ваш опубликованный код воспроизводит проблему? Если нет, продолжайте корректировать, пока не произойдет. Другими словами, поставьте себя на место тех, кто пытается вам помочь. Опять же, просто предложения.

Ниже приведен код, который я попытался построить на основе того, что вы показали, избегая проблем, о которых я упоминал выше. Я не утверждаю, что он выдает правильный вывод, но должен дать вам представление о том, как исправить некоторые из ошибок, указанных выше.

#include <iostream>
#include <fstream>
struct Info{
    int width;
    int height;
    int offset;
    unsigned char * info;
    unsigned char * data;

    int size;
};

Info readBMP(char* filename)
{
    int i;
    std::ifstream is(filename, std::ifstream::binary);
    is.seekg(0, is.end);
    i = is.tellg();
    is.seekg(0);
    unsigned char *info = new unsigned char[i];
    is.read((char *)info,i);

    int width = *(int*)&info[18];
    int height = *(int*)&info[22];
    int offset = *(int*)&info[10];

    unsigned char a[offset];
    unsigned char *b = new unsigned char[i - offset];
    std::copy(info,
              info + offset,
              a);

    std::copy(info + offset,
              info + i,
              b + 0);

    Info dat;
    dat.width = width;
    dat.height = height;
    dat.offset = offset;
    dat.size = i;
    dat.info = new unsigned char[offset - 1];
    dat.data = new unsigned char[i - offset + 1];

    for( int j = 0; j < offset ; j++ ){
        dat.info[j] = a[j];
    }

    for( int j = 0; j < i - offset; j++ ){
        dat.data[j] = b[j];
    }
    return dat;

}

__global__ void CreateGaussFilter(unsigned char * src, unsigned char * dst, int kernalHeight, int kernalWidth, double *kernalArray, int rows, int cols){
    int verticleImageBound=(kernalHeight-1)/2;
    int horizontalImageBound=(kernalWidth-1)/2;

    int row = ( blockDim.x * blockIdx.x + threadIdx.x ) + verticleImageBound;
    int col = ( blockDim.y * blockIdx.y + threadIdx.y ) + horizontalImageBound;

    if ( row >= rows - verticleImageBound || col >= cols - horizontalImageBound ) return;
            float  value=0;
            for(int kRow=0;kRow<kernalHeight;kRow++){
                  for(int kCol=0;kCol<kernalWidth;kCol++){
                  float pixel = src[ 3 * ((kRow+row-verticleImageBound ) * cols + (kCol+col-horizontalImageBound)) ] * kernalArray[kRow + kCol * kernalWidth];
                  value+=pixel;
             }
            }

            printf("value = %f\n",round(value));

            dst[3 * ( row * cols + col )] = round(value);
            dst[3 * ( row * cols + col ) + 1] = round(value);
            dst[3 * ( row * cols + col ) + 2] = round(value);

}

__global__ void greyScale( unsigned char * src , int rows, int cols){

        int i = blockDim.x * blockIdx.x + threadIdx.x;
        int j = blockDim.y * blockIdx.y + threadIdx.y;

        if( i >= rows || j >= cols ) {
                return;
        }
            unsigned char r = src[3 * (i * cols + j)];
            unsigned char g = src[3 * (i * cols + j) + 1];
            unsigned char b = src[3 * (i * cols + j) + 2];

            unsigned char linearIntensity = (unsigned char)(0.2126f * r + 0.7512f * g + 0);

            src[3 * (i * cols + j)] = linearIntensity;
            src[3 * (i * cols + j) + 1] = linearIntensity;
            src[3 * (i * cols + j) + 2] = linearIntensity;

}

int main() {

    double GKernel[5][5] = {0.1};
    //FilterCreation(GKernel);
    double * kernel = new double[25];
    int i,j,k = 0;
    for( int i = 0; i < 5; i++){
     for( int j = 0; j < 5; j++){
       kernel[k++] = GKernel[i][j];
     }
    }

    double * deviceKernel;

    cudaMalloc((void **)&deviceKernel, 25 * sizeof(double));
    cudaMemcpy(deviceKernel, kernel, 25 * sizeof(double), cudaMemcpyHostToDevice);

    Info dat; // = readBMP("konik.bmp");
    dat.width = 766;
    dat.height = 511;
    dat.size = dat.width*dat.height*3;
    dat.offset = 0;
    dat.data = new unsigned char[dat.size];
    unsigned char * devPtr;
    unsigned char * devPtrFilter;
    size_t pitch;

    unsigned char * test= new unsigned char  [dat.size - dat.offset ];

    cudaMalloc (( void **)& devPtr , dat.size * sizeof ( unsigned char  ));
    cudaMalloc (( void **)& devPtrFilter , dat.size * sizeof ( unsigned char  ));

    cudaMemcpy ( devPtr , dat.data , sizeof ( unsigned char ) *  dat.size , cudaMemcpyHostToDevice );

    dim3 block(32,32);
    dim3 grid((dat.height+31)/32, (dat.width+31)/32);

        greyScale<<<grid,block>>>(devPtr,dat.height,dat.width);
        CreateGaussFilter<<<grid,block>>>(devPtr,devPtrFilter,5,5,deviceKernel,dat.height,dat.width);

    cudaMemcpy ( test, devPtrFilter , sizeof ( unsigned char )  *  dat.size ,cudaMemcpyDeviceToHost );
    cudaDeviceSynchronize ();
#if 0
    std::ofstream fout;
    fout.open("output.bmp", std::ios::binary | std::ios::out);
    fout.write( reinterpret_cast<char *>(dat.info), dat.offset);

    fout.write( reinterpret_cast<char *>(test), dat.size - dat.offset );
    fout.close();
#endif
    return 0;
}

Были некоторые другие проблемы с вашим кодом.

ваш RGB в градациях серого использует 0 вместо b.
ваше ядро гауссов не записывает во все выходные точки, поэтому мы сначала заполним вывод 0.
Я предоставил свои собственные коэффициенты гауссовского ядра 5x5.

С этими дополнительными изменениями:

$ cat t8.cu
#include <iostream>
#include <fstream>
#include <stdio.h>

struct Info{
    int width;
    int height;
    int offset;
    unsigned char * info;
    unsigned char * data;

    int size;
};

Info readBMP(const char* filename)
{
    int i;
    std::ifstream is(filename, std::ifstream::binary);
    is.seekg(0, is.end);
    i = is.tellg();
    is.seekg(0);
    unsigned char *info = new unsigned char[i];
    is.read((char *)info,i);

    int width = *(int*)&info[18];
    int height = *(int*)&info[22];
    int offset = *(int*)&info[10];

    unsigned char a[offset];
    unsigned char *b = new unsigned char[i - offset];
    std::copy(info,
              info + offset,
              a);

    std::copy(info + offset,
              info + i,
              b + 0);

    Info dat;
    dat.width = width;
    dat.height = height;
    dat.offset = offset;
    dat.size = i;
    dat.info = new unsigned char[offset - 1];
    dat.data = new unsigned char[i - offset + 1];

    for( int j = 0; j < offset ; j++ ){
        dat.info[j] = a[j];
    }

    for( int j = 0; j < i - offset; j++ ){
        dat.data[j] = b[j];
    }
    return dat;

}

__global__ void CreateGaussFilter(unsigned char * src, unsigned char * dst, int kernalHeight, int kernalWidth, double *kernalArray, int rows, int cols){
    int verticleImageBound=(kernalHeight-1)/2;
    int horizontalImageBound=(kernalWidth-1)/2;

    int row = ( blockDim.x * blockIdx.x + threadIdx.x ) + verticleImageBound;
    int col = ( blockDim.y * blockIdx.y + threadIdx.y ) + horizontalImageBound;

    if ( row >= rows - verticleImageBound || col >= cols - horizontalImageBound ) return;
            float  value=0;
            for(int kRow=0;kRow<kernalHeight;kRow++){
                  for(int kCol=0;kCol<kernalWidth;kCol++){
                  float pixel = src[ 3 * ((kRow+row-verticleImageBound ) * cols + (kCol+col-horizontalImageBound)) ] * kernalArray[kRow + kCol * kernalWidth];
                  value+=pixel;
             }
            }

//            printf("value = %f\n",round(value));

            dst[3 * ( row * cols + col )] = round(value);
            dst[3 * ( row * cols + col ) + 1] = round(value);
            dst[3 * ( row * cols + col ) + 2] = round(value);

}

__global__ void greyScale( unsigned char * src , int rows, int cols){

        int i = blockDim.x * blockIdx.x + threadIdx.x;
        int j = blockDim.y * blockIdx.y + threadIdx.y;

        if( i >= rows || j >= cols ) {
                return;
        }
            unsigned char r = src[3 * (i * cols + j)];
            unsigned char g = src[3 * (i * cols + j) + 1];
            unsigned char b = src[3 * (i * cols + j) + 2];

            unsigned char linearIntensity = (unsigned char)(0.21f * r + 0.72f * g + 0.07 * b);

            src[3 * (i * cols + j)] = linearIntensity;
            src[3 * (i * cols + j) + 1] = linearIntensity;
            src[3 * (i * cols + j) + 2] = linearIntensity;

}
int main() {

    double GKernel[5][5] = {{1,4,7,4,1},{4,16,26,16,4},{7,26,41,26,7},{4,16,26,16,4},{1,4,7,4,1}};
    //FilterCreation(GKernel);
    double * kernel = new double[25];
    int k = 0;
    for( int i = 0; i < 5; i++){
     for( int j = 0; j < 5; j++){
       kernel[k++] = GKernel[i][j]/273;
     }
    }

    double * deviceKernel;

    cudaMalloc((void **)&deviceKernel, 25 * sizeof(double));
    cudaMemcpy(deviceKernel, kernel, 25 * sizeof(double), cudaMemcpyHostToDevice);

    Info dat = readBMP("input.bmp");
#if 0
    dat.width = 766;
    dat.height = 511;
    dat.size = dat.width*dat.height*3;
    dat.offset = 0;
    dat.data = new unsigned char[dat.size];
    for (int i = 0; i<dat.size; i++) dat.data[i] = (i%dat.width)%255;
#endif
    unsigned char * devPtr;
    unsigned char * devPtrFilter;

    unsigned char * test= new unsigned char  [dat.size - dat.offset ];

    cudaMalloc (( void **)& devPtr , dat.size * sizeof ( unsigned char  ));
    cudaMalloc (( void **)& devPtrFilter , dat.size * sizeof ( unsigned char  ));
    cudaMemset(devPtrFilter, 0, dat.size);
    cudaMemcpy ( devPtr , dat.data , sizeof ( unsigned char ) *  dat.size , cudaMemcpyHostToDevice );

    dim3 block(32,32);
    dim3 grid((dat.height+31)/32, (dat.width+31)/32);

        greyScale<<<grid,block>>>(devPtr,dat.height,dat.width);
        CreateGaussFilter<<<grid,block>>>(devPtr,devPtrFilter,5,5,deviceKernel,dat.height,dat.width);

    cudaMemcpy ( test, devPtrFilter , sizeof ( unsigned char )  *  (dat.size - dat.offset) ,cudaMemcpyDeviceToHost );
    cudaDeviceSynchronize ();
    std::ofstream fout;
    fout.open("output.bmp", std::ios::binary | std::ios::out);
    fout.write( reinterpret_cast<char *>(dat.info), dat.offset);

    fout.write( reinterpret_cast<char *>(test), dat.size - dat.offset );
    fout.close();
    return 0;
}
$ nvcc -o t8 t8.cu
$ cuda-memcheck ./t8

и начиная с доступного файла BarbaraBlocks1.bmpна /usr/local/cuda/samples/3_Imaging/dct8x8/data на стандартной текущей установке Linux CUDA, которая выглядит следующим образом:

Это производит вывод, как это:

CUDA не выполняет код после вложенного цикла в функции ядра

Пожалуйста, войдите или зарегистрируйтесь чтобы ответить на этот вопрос.

1 Ответ

Пожалуйста, войдите или зарегистрируйтесь что бы добавить комментарий.

CUDA не выполняет код после вложенного цикла в функции ядра

Пожалуйста, войдите или зарегистрируйтесь чтобы ответить на этот вопрос.

1 Ответ

Пожалуйста, войдите или зарегистрируйтесь что бы добавить комментарий.

Нет похожих вопросов