sobel.cu

__device__ char *
get_pixel_uv(char *pDevPtr, int pitch, int row, int col) {
    return pDevPtr + row * pitch + col * 2;
}

__device__ char *
get_pixel_y(char *pDevPtr, const int pitch, const int row, const int col, const int width, const int height) {
    if (col >= 0 && row >= 0 && col < width && row < height) {
        return pDevPtr + row * pitch + col;
    } else {
        return NULL;
    }
}

__device__ int
get_pixel_val(const char *ptr) {
    return NULL != ptr ? (int)ptr[0] : 0;
}

__global__ void
sobel_grey(char *ptrYdata, char *ptrUVdata, const int pitch, const int width, const int height) {
    const int row = blockIdx.y * blockDim.y + threadIdx.y;
    const int col = blockIdx.x * blockDim.x + threadIdx.x;
    char *yOut    = get_pixel_y(ptrYdata, pitch, row, col, width, height);
    if (NULL == yOut) {
        return;
    }
    int sum_w = 0, sum_h = 0, tmp_a, tmp_b;
    tmp_a = get_pixel_val(get_pixel_y(ptrYdata, pitch, row - 1, col - 1, width, height));
    tmp_b = get_pixel_val(get_pixel_y(ptrYdata, pitch, row - 1, col + 1, width, height));
    sum_w += tmp_a - tmp_b;
    sum_h += tmp_a + tmp_b;
    sum_h += get_pixel_val(get_pixel_y(ptrYdata, pitch, row - 1, col, width, height)) * 2;
    tmp_a = get_pixel_val(get_pixel_y(ptrYdata, pitch, row, col - 1, width, height));
    tmp_b = get_pixel_val(get_pixel_y(ptrYdata, pitch, row, col + 1, width, height));
    sum_w += (tmp_a - tmp_b) * 2;
    tmp_a = get_pixel_val(get_pixel_y(ptrYdata, pitch, row + 1, col - 1, width, height));
    tmp_b = get_pixel_val(get_pixel_y(ptrYdata, pitch, row + 1, col + 1, width, height));
    sum_w += tmp_a - tmp_b;
    sum_h += -(tmp_a + tmp_b);
    sum_h += -get_pixel_val(get_pixel_y(ptrYdata, pitch, row + 1, col, width, height)) * 2;
    if (sum_w < 0) {
        sum_w = -sum_w;
    }
    if (sum_w > 255) {
        sum_w = 255;
    }
    if (sum_h < 0) {
        sum_h = -sum_h;
    }
    if (sum_h > 255) {
        sum_h = 255;
    }
    __syncthreads();
    yOut[0] = (char)((sum_w + sum_h) / 2);
}

static int
sobel_grey(CUdeviceptr ptrYdata, CUdeviceptr ptrUVdata, int pitch, int width, int height) {
    dim3 threadsPerBlock(32, 32);
    dim3 blocks(((width + 31) >> 5), ((height + 31) >> 5));
    sobel_grey<<<blocks, threadsPerBlock>>>((char *)ptrYdata, (char *)ptrUVdata, pitch, width, height);
    return 0;
}