#include <cuda.h>
#include <cutil.h>
#include <stdio.h>
#define BLOCKS 32
#define IND(x,y,w) (y * w + x)
typedef unsigned char uchar;
const int conv_r = 2;
float conv_m[5][5] = {
{ 0.01, 0.01, 0.01, 0.01, 0.01 },
{ 0.01, 0.1, 0.1, 0.1, 0.01 },
{ 0.01, 0.1, 0, 0.1, 0.01 },
{ 0.01, 0.1, 0.1, 0.1, 0.01 },
{ 0.01, 0.01, 0.01, 0.01, 0.01 }
};
__global__ void imageConvoluter ( uchar* in, uchar* out, int w, int h, float conv_m[5][5] ) {
int idx = blockIdx.x*blockDim.x + threadIdx.x;
int idy = blockIdx.y*blockDim.y + threadIdx.y;
int index = idx * w + idy;
out[index] = 0;
for ( int idx_ = -conv_r; idx_ <= conv_r; idx_++ ) {
for ( int idy_ = -conv_r; idy_ <= conv_r; idy_ ++ ) {
int local_index_x = idx + idx_;
int local_index_y = idy + idy_;
if ( (local_index_x > 0) && (local_index_x < w) && (local_index_y > 0) && (local_index_y < h) )
out[index] += in[IND(local_index_x, local_index_y, w)] * conv_m[idy_+conv_r][idx_+conv_r];
}
}
}
int main () {
unsigned int image_w, image_h;
uchar *data_original = NULL,
*data_original_dev = NULL,
*data_convoluted_dev = NULL;
CUT_SAFE_CALL( cutLoadPGMub( "pic.pgm", &data_original, &image_w, &image_h ) );
printf( "W=%d, H=%d\n", image_w, image_h );
int image_size = image_w * image_h * sizeof(uchar);
uchar *data_convoluted = new uchar[image_w * image_h],
*data_convoluted_cpu = new uchar[image_w * image_h];
cudaMalloc( &data_original_dev, image_size );
cudaMalloc( &data_convoluted_dev, image_size );
CUDA_SAFE_CALL( cudaMemcpy( data_original_dev, data_original, image_size, cudaMemcpyHostToDevice ) );
dim3 blocks = dim3( BLOCKS, BLOCKS );
dim3 threads = dim3( image_w / BLOCKS, image_h / BLOCKS );
imageConvoluter <<< blocks, threads >>> ( data_original_dev, data_convoluted_dev, image_w, image_h, conv_m );
for ( int idx = conv_r; idx < image_h-conv_r; idx++ ) {
for ( int idy = conv_r; idy < image_w-conv_r; idy++ ) {
int index = IND( idx, idy, image_w );
data_convoluted_cpu[index] = 0;
for ( int idx_ = -conv_r; idx_ <= conv_r; idx_++ ) {
for ( int idy_ = -conv_r; idy_ <= conv_r; idy_ ++ ) {
int index_around = IND((idx + idx_), (idy + idy_), image_w);
data_convoluted_cpu[index] += data_original[index_around] * conv_m[idy_+conv_r][idx_+conv_r];
}
}
}
}
CUDA_SAFE_CALL( cudaMemcpy( data_convoluted, data_convoluted_dev, image_size, cudaMemcpyDeviceToHost ) );
// CUT_SAFE_CALL( cutCompareub( data_convoluted, data_convoluted_cpu, image_w * image_h ) );
CUT_SAFE_CALL( cutSavePGMub( "pic_convoluted.pgm", data_convoluted, image_w, image_h ) );
CUT_SAFE_CALL( cutSavePGMub( "pic_convoluted_cpu.pgm", data_convoluted_cpu, image_w, image_h ) );
cutFree( data_original );
cudaFree( data_original_dev );
delete data_convoluted;
delete data_convoluted_cpu;
}
Изображение считывается в массив, производится свёртка изображения в соседний массив. На этапе копирования памяти с девайса обратно на хост случается некоторый конфуз (выловлено в дебаггере):
При выполнении программы без дебаггера на этом вызове процесс встаёт, и далее ничего не происходит.
Если я правильно понимаю, `unspecified launch failure' возникает из-за неправильного обращения к памяти и скорее всего является простым сегфолтом.
Проблема в том, что после выполнения этого кода кудовский девайс помирает полностью и заново не поднимается:
Хотелось бы понять, что я делаю не так, и что вообще происходит.
