<commit_before>#include <stdio.h>

// We assume that size are divisible
#define RADIUS        2
#define BLOCK_SIZE_X    16 
#define BLOCK_SIZE_Y    16 
#define MATRIX_WIDTH (128)
#define MATRIX_HEIGTH (128)
#define NUM_ELEMENTS (MATRIX_HEIGTH * MATRIX_WIDTH)

// CUDA API error checking macro
#define cudaCheck(error) \
  if (error != cudaSuccess) { \
    printf("Fatal error: %s at %s:%d\n", \
      cudaGetErrorString(error), \
      __FILE__, __LINE__); \
    exit(1); \
  }

__global__ void stencil_2d(int *in, int *out) 
{
    /*

      Fill kernel code!
    
    */
}

int main()
{
  unsigned int i, j;
  int h_in[NUM_ELEMENTS], h_out[NUM_ELEMENTS];
  int *d_in, *d_out;

  // To access element (i, j) of the matrix, use h_in[i + j * MATRIX_WIDTH]
  // Initialize host data
  for (i = 0; i < (NUM_ELEMENTS); ++i)
    h_in[i] = 1;

  // Allocate space on the device
  cudaCheck( cudaMalloc( &d_in, NUM_ELEMENTS * sizeof(int)) );
  cudaCheck( cudaMalloc( &d_out, NUM_ELEMENTS * sizeof(int)) );

  // Copy input data to device
  cudaCheck( cudaMemcpy( d_in, h_in, NUM_ELEMENTS * sizeof(int), cudaMemcpyHostToDevice) );

  dim3 blocks = dim3(MATRIX_WIDTH / BLOCK_SIZE_X, MATRIX_WIDTH / BLOCK_SIZE_Y, 1);
  dim3 threads = dim3(BLOCK_SIZE_X, BLOCK_SIZE_Y, 1);

  stencil_2d<<< blocks, threads >>> (d_in, d_out);

  cudaCheck(cudaPeekAtLastError());

  cudaCheck( cudaMemcpy( h_out, d_out, NUM_ELEMENTS * sizeof(int), cudaMemcpyDeviceToHost) );

  // Verify results (inclusion-exclusion principle)
  for (j = 0; j < MATRIX_WIDTH; ++j) {
    for (i = 0; i < MATRIX_HEIGTH; ++i) {
      int expected = (2 * RADIUS + 1) * (2 * RADIUS + 1) - 
        (j < RADIUS ? (RADIUS - j) * (2 * RADIUS + 1) : 0) -
        (i < RADIUS ? (RADIUS - i) * (2 * RADIUS + 1) : 0) +
        ((j < RADIUS && i < RADIUS) ? (RADIUS - j) * (RADIUS - i) : 0) -
        (j > MATRIX_WIDTH - RADIUS - 1 ? (j + RADIUS + 1 - MATRIX_WIDTH) * (2 * RADIUS + 1) : 0) -
        (i > MATRIX_HEIGTH - RADIUS - 1 ? (i + RADIUS + 1 - MATRIX_HEIGTH) * (2 * RADIUS + 1) : 0) +
        (j > MATRIX_WIDTH - RADIUS - 1 && i > MATRIX_HEIGTH - RADIUS - 1 ?
         (j + RADIUS + 1 - MATRIX_WIDTH) * (i + RADIUS + 1 - MATRIX_HEIGTH) : 0) + 
        (j < RADIUS && i > MATRIX_HEIGTH - RADIUS - 1 ? 
         (RADIUS - j) * (i + RADIUS + 1 - MATRIX_HEIGTH) : 0) + 
        (i < RADIUS && j > MATRIX_WIDTH - RADIUS - 1 ?
         (RADIUS - i) *  (j + RADIUS + 1 - MATRIX_WIDTH) : 0);
      if (h_out[j + i * MATRIX_WIDTH] != expected) {
        printf("Element h_out[%d + %d * MATRIX_WIDTH] == %d != %d\n", j, i, h_out[i], expected);
        break;
      }
    }
  }

  if (i * j == NUM_ELEMENTS)
    printf("SUCCESS!\n");
  else
    printf("FAILURE!\n");

  // Free out memory
  cudaFree(d_in);
  cudaFree(d_out);

  return 0;
}<commit_msg>fix HEIGHT<commit_after>#include <stdio.h>

// We assume that size are divisible
#define RADIUS        2
#define BLOCK_SIZE_X    16 
#define BLOCK_SIZE_Y    16 
#define MATRIX_WIDTH (128)
#define MATRIX_HEIGHT (128)
#define NUM_ELEMENTS (MATRIX_HEIGHT * MATRIX_WIDTH)

// CUDA API error checking macro
#define cudaCheck(error) \
  if (error != cudaSuccess) { \
    printf("Fatal error: %s at %s:%d\n", \
      cudaGetErrorString(error), \
      __FILE__, __LINE__); \
    exit(1); \
  }

__global__ void stencil_2d(int *in, int *out) 
{
    /*

      Fill kernel code!
    
    */
}

int main()
{
  unsigned int i, j;
  int h_in[NUM_ELEMENTS], h_out[NUM_ELEMENTS];
  int *d_in, *d_out;

  // To access element (i, j) of the matrix, use h_in[i + j * MATRIX_WIDTH]
  // Initialize host data
  for (i = 0; i < (NUM_ELEMENTS); ++i)
    h_in[i] = 1;

  // Allocate space on the device
  cudaCheck( cudaMalloc( &d_in, NUM_ELEMENTS * sizeof(int)) );
  cudaCheck( cudaMalloc( &d_out, NUM_ELEMENTS * sizeof(int)) );

  // Copy input data to device
  cudaCheck( cudaMemcpy( d_in, h_in, NUM_ELEMENTS * sizeof(int), cudaMemcpyHostToDevice) );

  dim3 blocks = dim3(MATRIX_WIDTH / BLOCK_SIZE_X, MATRIX_HEIGHT / BLOCK_SIZE_Y, 1);
  dim3 threads = dim3(BLOCK_SIZE_X, BLOCK_SIZE_Y, 1);

  stencil_2d<<< blocks, threads >>> (d_in, d_out);

  cudaCheck(cudaPeekAtLastError());

  cudaCheck( cudaMemcpy( h_out, d_out, NUM_ELEMENTS * sizeof(int), cudaMemcpyDeviceToHost) );

  // Verify results (inclusion-exclusion principle)
  for (j = 0; j < MATRIX_WIDTH; ++j) {
    for (i = 0; i < MATRIX_HEIGHT; ++i) {
      int expected = (2 * RADIUS + 1) * (2 * RADIUS + 1) - 
        (j < RADIUS ? (RADIUS - j) * (2 * RADIUS + 1) : 0) -
        (i < RADIUS ? (RADIUS - i) * (2 * RADIUS + 1) : 0) +
        ((j < RADIUS && i < RADIUS) ? (RADIUS - j) * (RADIUS - i) : 0) -
        (j > MATRIX_WIDTH - RADIUS - 1 ? (j + RADIUS + 1 - MATRIX_WIDTH) * (2 * RADIUS + 1) : 0) -
        (i > MATRIX_HEIGHT - RADIUS - 1 ? (i + RADIUS + 1 - MATRIX_HEIGHT) * (2 * RADIUS + 1) : 0) +
        (j > MATRIX_WIDTH - RADIUS - 1 && i > MATRIX_HEIGHT - RADIUS - 1 ?
         (j + RADIUS + 1 - MATRIX_WIDTH) * (i + RADIUS + 1 - MATRIX_HEIGHT) : 0) + 
        (j < RADIUS && i > MATRIX_HEIGHT - RADIUS - 1 ? 
         (RADIUS - j) * (i + RADIUS + 1 - MATRIX_HEIGHT) : 0) + 
        (i < RADIUS && j > MATRIX_WIDTH - RADIUS - 1 ?
         (RADIUS - i) *  (j + RADIUS + 1 - MATRIX_WIDTH) : 0);
      if (h_out[j + i * MATRIX_WIDTH] != expected) {
        printf("Element h_out[%d + %d * MATRIX_WIDTH] == %d != %d\n", j, i, h_out[i], expected);
        break;
      }
    }
  }

  if (i * j == NUM_ELEMENTS)
    printf("SUCCESS!\n");
  else
    printf("FAILURE!\n");

  // Free out memory
  cudaFree(d_in);
  cudaFree(d_out);

  return 0;
}
{'fraction_non_alphanumeric': 0.15530846484935437, 'fraction_numerical': 0.020444763271162122, 'mean_word_length': 2.565856777493606, 'pattern_counts': {'":': 0, '<': 29, '<?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': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '727498', 'n_tokens_mistral': 2583, 'n_tokens_neox': 2327, 'n_words': 577}