<commit_before>// System includes
#include <iostream>
#include <string>

// CUDA runtime
#include <cuda_runtime.h>

#include "vvms_cuda.h"
#include "cuda_helper.h"

namespace himan
{

namespace plugin
{

namespace vvms_cuda
{

__global__ void Calculate(const double* __restrict__ dT,
							const double* __restrict__ dVV,
							const double* __restrict__ dP,
							double* __restrict__ dVVOut,
							vvms_cuda_options opts,
							int* dMissingValuesCount);

} // namespace vvms_cuda
} // namespace plugin
} // namespace himan


__device__ void himan::plugin::vvms_cuda::Calculate(const double* __restrict__ dT,
														const double* __restrict__ dVV,
														const double* __restrict__ dP,
														double* __restrict__ VVMS,
														vvms_cuda_options opts,
														int* dMissingValuesCount)
{

	int idx = blockIdx.x * blockDim.x + threadIdx.x;

	if (idx < opts.N)
	{
		double P = (opts.isConstantPressure) ? opts.PConst : dP[idx];

		if (dT[idx] == kFloatMissing || dVV[idx] == kFloatMissing || P == kFloatMissing)
		{
			atomicAdd(dMissingValuesCount, 1);
			VVMS[idx] = kFloatMissing;
		}
		else
		{
			VVMS[idx] = 287 * -dVV[idx] * (opts.TBase + dT[idx]) / (9.81 * P);
		}<commit_msg>device --> global<commit_after>// System includes
#include <iostream>
#include <string>

// CUDA runtime
#include <cuda_runtime.h>

#include "vvms_cuda.h"
#include "cuda_helper.h"

namespace himan
{

namespace plugin
{

namespace vvms_cuda
{

__global__ void Calculate(const double* __restrict__ dT,
							const double* __restrict__ dVV,
							const double* __restrict__ dP,
							double* __restrict__ dVVOut,
							vvms_cuda_options opts,
							int* dMissingValuesCount);

} // namespace vvms_cuda
} // namespace plugin
} // namespace himan


__global__ void himan::plugin::vvms_cuda::Calculate(const double* __restrict__ dT,
														const double* __restrict__ dVV,
														const double* __restrict__ dP,
														double* __restrict__ VVMS,
														vvms_cuda_options opts,
														int* dMissingValuesCount)
{

	int idx = blockIdx.x * blockDim.x + threadIdx.x;

	if (idx < opts.N)
	{
		double P = (opts.isConstantPressure) ? opts.PConst : dP[idx];

		if (dT[idx] == kFloatMissing || dVV[idx] == kFloatMissing || P == kFloatMissing)
		{
			atomicAdd(dMissingValuesCount, 1);
			VVMS[idx] = kFloatMissing;
		}
		else
		{
			VVMS[idx] = 287 * -dVV[idx] * (opts.TBase + dT[idx]) / (9.81 * P);
		}
{'fraction_non_alphanumeric': 0.1147875816993464, 'fraction_numerical': 0.005718954248366013, 'mean_word_length': 3.0345963756177925, 'pattern_counts': {'":': 0, '<': 11, '<?xml version=': 0, '>': 10, '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': '6592287', 'n_tokens_mistral': 1125, 'n_tokens_neox': 879, 'n_words': 203}