I want to know how to make this cuda kernel more precise or atleast consistent.

__global__ void adam_kernel(float* value, float* m, float* v, float* grad, int size) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) {
        m[idx] = 0.95f * m[idx] + (1 - 0.95f) * grad[idx];
        v[idx] = 0.999f * v[idx] + (1 - 0.999f) * grad[idx] * grad[idx];
        value[idx] -= 0.01f * m[idx] / (sqrtf(v[idx]) + 1e-8f);
    }
}

float randomFrom(float min, float max) {
    static std::random_device rdev;
    static std::default_random_engine re(rdev());
    std::uniform_real_distribution<float> dist(min, max);
    return dist(re);
}

int main()
{
    const int size = 764 * 64;

    // host variables
    float* value = new float[size];
    float* m = new float[size];
    float* v = new float[size];
    float* grad = new float[size];

    // init with host variables
    for (int i = 0; i < size; ++i) {
        value[i] = randomFrom(0, 0.00001f);
        m[i] = 0;
        v[i] = 0;
        grad[i] = randomFrom(0, 0.00001f);
    }

    // device variables
    float* d_value;
    float* d_m;
    float* d_v;
    float* d_grad;

    // alloc device memory
    cudaMalloc(&d_value, size * sizeof(float));
    cudaMalloc(&d_m, size * sizeof(float));
    cudaMalloc(&d_v, size * sizeof(float));
    cudaMalloc(&d_grad, size * sizeof(float));

    // copy host to device
    cudaMemcpy(d_value, value, size * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_m, m, size * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_v, v, size * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_grad, grad, size * sizeof(float), cudaMemcpyHostToDevice);

    // run kernel
    adam_kernel<<<1, size>>>(d_value, d_m, d_v, d_grad, size);

    // copy device to host
    cudaMemcpy(value, d_value, size * sizeof(float), cudaMemcpyDeviceToHost);

    // print results
    float sum = 0;
    for (int i = 0; i < size; ++i) {
        sum += value[i];
    }

    cout << "Sum: " << sum << endl;

    return 0;
}

I know this gives me these different results all the time since I am using float but is there no way to atleast make the results consistent? I don’t want to use double so is there another way?

Thanks in advance!