<commit_before>// -*- C++ -*-

#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <string>
#include <algorithm>

// yucky, but for asking the system how many cores we have
#include <unistd.h>

// header file for openmp
#include <omp.h>

// header files for kokkos
#include <Kokkos_Core.hpp>

#include <cuda_runtime.h>

using std::string;
using std::vector;

#define BLOCK_SIZE 64;

//Pre-C++11 timing (thanks jeff)
double getElapsedTime(const timespec start, const timespec end) {
	timespec temp;
	if ((end.tv_nsec-start.tv_nsec)<0) {
		temp.tv_sec = end.tv_sec-start.tv_sec-1;
		temp.tv_nsec = 1000000000+end.tv_nsec-start.tv_nsec;
	} else {
		temp.tv_sec = end.tv_sec-start.tv_sec;
		temp.tv_nsec = end.tv_nsec-start.tv_nsec;
	}
	return double(temp.tv_sec) + double(temp.tv_nsec) / 1e9;
}

//Hacky random double... Right now it's just ints casted that way
double random_double() {
  return (double) rand();
}

__global__
void
cudaDocontractFieldFieldScalar_kernel(const double * const __restrict__ d_left, const double * const __restrict__ d_right,
    double * d_out,
    int numCells,
    int numLeftFields,
    int numRightFields,
    int numPoints) {

  int myID = (blockIdx.x * blockDim.x) + threadIdx.x;
  if(myID < (numCells * numLeftFields * numRightFields)) {

    int myMatrix = myID / (numLeftFields * numRightFields);
    int matrixIndex = myID % (numLeftFields * numRightFields);
    int matrixRow = matrixIndex / numRightFields;
    int matrixCol = matrixIndex % numRightFields;
    double temp = 0;
    for (int qp = 0; qp < numPoints; qp++) {

      temp += d_left[myMatrix*numPoints*numLeftFields + numPoints*matrixRow + qp] *
        d_right[myMatrix*numPoints*numRightFields + qp*numRightFields + matrixCol];

    }
    d_out[myID]= temp;

  }

}

void
cudaDocontractFieldFieldScalar(double * h_out,
    double * h_inLeft,
    double * h_inRight,
    int numCells,
    int numLeftFields,
    int numRightFields,
    int numPoints,
    timespec * tic = NULL,
    timespec * toc = NULL) {

  double * d_right;
  double * d_left;
  double * d_out;

  cudaMalloc(&d_right, sizeof(double) * numCells * numPoints * numRightFields);
  cudaMalloc(&d_left, sizeof(double) * numCells * numPoints * numLeftFields);
  cudaMalloc(&d_out, sizeof(double) * numCells * numRightFields * numLeftFields);
  cudaMemset(d_out, 0, sizeof(double) * numCells * numRightFields * numLeftFields);

  cudaMemcpy(d_right, h_inRight,
      sizeof(double) * numCells * numPoints * numRightFields, cudaMemcpyHostToDevice);
  cudaMemcpy(d_left, h_inLeft,
      sizeof(double) * numCells * numPoints * numLeftFields, cudaMemcpyHostToDevice);

  dim3 blockSize(1024);
  dim3 gridSize((numCells * numLeftFields * numRightFields / 1024) + 1);

  if (tic != NULL)
    clock_gettime(CLOCK_MONOTONIC, tic);

  cudaDocontractFieldFieldScalar_kernel<<<gridSize, blockSize>>>(d_left,
      d_right, d_out, numCells, numLeftFields, numRightFields, numPoints);
  cudaDeviceSynchronize();

  if (toc != NULL)
    clock_gettime(CLOCK_MONOTONIC, toc);

  cudaMemcpy(h_out, d_out, sizeof(double) * numCells * numLeftFields * numRightFields, cudaMemcpyDeviceToHost);

}

template<class DeviceType, class LeftViewType, class RightViewType, class OutputViewType>
struct contractFieldFieldScalarFunctor {
	typedef DeviceType device_type;
	LeftViewType _leftFields;
	RightViewType _rightFields;
	OutputViewType _outputFields;
	int _numPoints;
	int _numLeftFields;
	int _numRightFields;

	contractFieldFieldScalarFunctor(LeftViewType leftFields,
			RightViewType rightFields,
			OutputViewType outputFields,
			int numLeftFields,
			int numRightFields,
			int numPoints) :<commit_msg>Added cuda frees<commit_after>// -*- C++ -*-

#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <string>
#include <algorithm>

// yucky, but for asking the system how many cores we have
#include <unistd.h>

// header file for openmp
#include <omp.h>

// header files for kokkos
#include <Kokkos_Core.hpp>

#include <cuda_runtime.h>

using std::string;
using std::vector;

#define BLOCK_SIZE 64;

//Pre-C++11 timing (thanks jeff)
double getElapsedTime(const timespec start, const timespec end) {
	timespec temp;
	if ((end.tv_nsec-start.tv_nsec)<0) {
		temp.tv_sec = end.tv_sec-start.tv_sec-1;
		temp.tv_nsec = 1000000000+end.tv_nsec-start.tv_nsec;
	} else {
		temp.tv_sec = end.tv_sec-start.tv_sec;
		temp.tv_nsec = end.tv_nsec-start.tv_nsec;
	}
	return double(temp.tv_sec) + double(temp.tv_nsec) / 1e9;
}

//Hacky random double... Right now it's just ints casted that way
double random_double() {
  return (double) rand();
}

__global__
void
cudaDocontractFieldFieldScalar_kernel(const double * const __restrict__ d_left, const double * const __restrict__ d_right,
    double * d_out,
    int numCells,
    int numLeftFields,
    int numRightFields,
    int numPoints) {

  int myID = (blockIdx.x * blockDim.x) + threadIdx.x;
  if(myID < (numCells * numLeftFields * numRightFields)) {

    int myMatrix = myID / (numLeftFields * numRightFields);
    int matrixIndex = myID % (numLeftFields * numRightFields);
    int matrixRow = matrixIndex / numRightFields;
    int matrixCol = matrixIndex % numRightFields;
    double temp = 0;
    for (int qp = 0; qp < numPoints; qp++) {

      temp += d_left[myMatrix*numPoints*numLeftFields + numPoints*matrixRow + qp] *
        d_right[myMatrix*numPoints*numRightFields + qp*numRightFields + matrixCol];

    }
    d_out[myID]= temp;

  }

}

void
cudaDocontractFieldFieldScalar(double * h_out,
    double * h_inLeft,
    double * h_inRight,
    int numCells,
    int numLeftFields,
    int numRightFields,
    int numPoints,
    timespec * tic = NULL,
    timespec * toc = NULL) {

  double * d_right;
  double * d_left;
  double * d_out;

  cudaMalloc(&d_right, sizeof(double) * numCells * numPoints * numRightFields);
  cudaMalloc(&d_left, sizeof(double) * numCells * numPoints * numLeftFields);
  cudaMalloc(&d_out, sizeof(double) * numCells * numRightFields * numLeftFields);
  cudaMemset(d_out, 0, sizeof(double) * numCells * numRightFields * numLeftFields);

  cudaMemcpy(d_right, h_inRight,
      sizeof(double) * numCells * numPoints * numRightFields, cudaMemcpyHostToDevice);
  cudaMemcpy(d_left, h_inLeft,
      sizeof(double) * numCells * numPoints * numLeftFields, cudaMemcpyHostToDevice);

  dim3 blockSize(1024);
  dim3 gridSize((numCells * numLeftFields * numRightFields / 1024) + 1);

  if (tic != NULL)
    clock_gettime(CLOCK_MONOTONIC, tic);

  cudaDocontractFieldFieldScalar_kernel<<<gridSize, blockSize>>>(d_left,
      d_right, d_out, numCells, numLeftFields, numRightFields, numPoints);
  cudaDeviceSynchronize();

  if (toc != NULL)
    clock_gettime(CLOCK_MONOTONIC, toc);

  cudaFree(d_right);
  cudaFree(d_left);
  cudaFree(d_out);

  cudaMemcpy(h_out, d_out, sizeof(double) * numCells * numLeftFields * numRightFields, cudaMemcpyDeviceToHost);

}

template<class DeviceType, class LeftViewType, class RightViewType, class OutputViewType>
struct contractFieldFieldScalarFunctor {
	typedef DeviceType device_type;
	LeftViewType _leftFields;
	RightViewType _rightFields;
	OutputViewType _outputFields;
	int _numPoints;
	int _numLeftFields;
	int _numRightFields;

	contractFieldFieldScalarFunctor(LeftViewType leftFields,
			RightViewType rightFields,
			OutputViewType outputFields,
			int numLeftFields,
			int numRightFields,
			int numPoints) :
{'fraction_non_alphanumeric': 0.10519935344827586, 'fraction_numerical': 0.008620689655172414, 'mean_word_length': 4.512249443207127, 'pattern_counts': {'":': 0, '<': 37, '<?xml version=': 0, '>': 31, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 2, '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': '6195065', 'n_tokens_mistral': 2754, 'n_tokens_neox': 2598, 'n_words': 591}