<commit_before>#include <cuda.h>
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>
#include <iostream>
#include <getopt.h>
#include <cstdio>
#define checkError(err)                                                                \
  if ((err) != cudaSuccess) {                                                          \
    printf("ERROR: %s in %s, line %d\n",cudaGetErrorString(err), __FILE__, __LINE__);  \
    exit(EXIT_FAILURE);                                                                \
  }

using namespace cv;
using namespace std;

__device__
bool inside_image(int row, int col, int width, int height) {
  return row >= 0 && row < height && col >= 0 && col < width;
}

__global__
void convolutionKernel(unsigned char* image, float* kernel, float* out_image, int kernel_n, int width, int height) {
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  
  if (row < height && col < width) {
    int n = kernel_n / 2;
    float accumulation = 0;
    for (int i = -n; i <= n; i++) {
      for (int j = -n; j <= n; j++) {
        if (inside_image(row + i, col + j, width, height)) {
          int image_idx = (row + i) * width + (col + j);
          int kernel_idx = (n + i) * kernel_n + (n + j);
          accumulation += image[image_idx] * kernel[kernel_idx];
        }
      }
    }
    out_image[row * width + col] = accumulation;
  }
}

__global__
void magnitudeKernel(float* x, float* y, unsigned char* r, int width, int height) {
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;

  if (row < height && col < width) {
    int idx = row * width + col;
    r[idx] = (unsigned char) hypotf(x[idx], y[idx]);
  }
}

void sobel(unsigned char *h_img, unsigned char *h_img_sobel, int width, int height, bool measure) {
  unsigned char *d_img, *d_img_sobel;
  float *d_img_sobel_x, *d_img_sobel_y;
  float *d_sobel_x, *d_sobel_y;
  long long size = width * height;
  cudaError_t err;
  cudaEvent_t start, stop;
<commit_msg>Using Fast inverse square root for magnitude calculation<commit_after>#include <cuda.h>
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>
#include <iostream>
#include <getopt.h>
#include <cstdio>
#define checkError(err)                                                                \
  if ((err) != cudaSuccess) {                                                          \
    printf("ERROR: %s in %s, line %d\n",cudaGetErrorString(err), __FILE__, __LINE__);  \
    exit(EXIT_FAILURE);                                                                \
  }

using namespace cv;
using namespace std;

__device__
bool inside_image(int row, int col, int width, int height) {
  return row >= 0 && row < height && col >= 0 && col < width;
}

__global__
void convolutionKernel(unsigned char* image, float* kernel, float* out_image, int kernel_n, int width, int height) {
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  
  if (row < height && col < width) {
    int n = kernel_n / 2;
    float accumulation = 0;
    for (int i = -n; i <= n; i++) {
      for (int j = -n; j <= n; j++) {
        if (inside_image(row + i, col + j, width, height)) {
          int image_idx = (row + i) * width + (col + j);
          int kernel_idx = (n + i) * kernel_n + (n + j);
          accumulation += image[image_idx] * kernel[kernel_idx];
        }
      }
    }
    out_image[row * width + col] = accumulation;
  }
}

__device__
float Q_rsqrt( float number )
{
	long i;
	float x2, y;
	const float threehalfs = 1.5F;

	x2 = number * 0.5F;
	y  = number;
	i  = * ( long * ) &y;                       // evil floating point bit level hacking
	i  = 0x5f3759df - ( i >> 1 );               // what the fuck? 
	y  = * ( float * ) &i;
	y  = y * ( threehalfs - ( x2 * y * y ) );   // 1st iteration
//	y  = y * ( threehalfs - ( x2 * y * y ) );   // 2nd iteration, this can be removed

	return y;
}

__device__
float Q_sqrt(float x) {
  return x * Q_rsqrt(x);
}

__global__
void magnitudeKernel(float* x, float* y, unsigned char* r, int width, int height) {
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;

  if (row < height && col < width) {
    int idx = row * width + col;
    r[idx] = (unsigned char) Q_sqrt(x[idx] * x[idx] + y[idx] * y[idx]);
  }
}

void sobel(unsigned char *h_img, unsigned char *h_img_sobel, int width, int height, bool measure) {
  unsigned char *d_img, *d_img_sobel;
  float *d_img_sobel_x, *d_img_sobel_y;
  float *d_sobel_x, *d_sobel_y;
  long long size = width * height;
  cudaError_t err;
  cudaEvent_t start, stop;

{'fraction_non_alphanumeric': 0.13313795268997278, 'fraction_numerical': 0.006489428511618171, 'mean_word_length': 2.2415196743554953, 'pattern_counts': {'":': 0, '<': 33, '<?xml version=': 0, '>': 23, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 12, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '874181', 'n_tokens_mistral': 1702, 'n_tokens_neox': 1659, 'n_words': 482}