<commit_before>// Copyright <NAME> 2016

// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at

//     http://www.apache.org/licenses/LICENSE-2.0

// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.

#include <iostream>
#include <memory>
#include <vector>
#include <map>

#include "EasyCL.h"

#include "CL/cl.h"

using namespace std;

static size_t grid[3];
static size_t block[3];
static unique_ptr<CLKernel> kernel;

unique_ptr<EasyCL> cl;
static cl_context *ctx;
static cl_command_queue *queue;
static cl_int err;

static vector<cl_mem> clmems;
static vector<cl_mem> kernelArgsToBeReleased;
static map<void *, int> idxByAddr;

static bool initialized = false;

void hostside_opencl_funcs_init() {
    cout << "initialize cl context" << endl;
    cl.reset(EasyCL::createForFirstGpuOtherwiseCpu());
    ctx = cl->context;
    queue = cl->queue;
}

void assure_initialized() {
    // yes this is not threadsafe.  or anything safe really...
    if(!initialized) {
        hostside_opencl_funcs_init();
        initialized = true;
    }
}

struct cudaDeviceProp {
    char name[256];
    size_t totalGlobalMem;
    size_t sharedMemPerBlock;
    int regsPerBlock;
    int warpSize;
    size_t memPitch;
    int maxThreadsPerBlock;
    int maxThreadsDim[3];
    int maxGridSize[3];
    size_t totalConstMem;
    int major;
    int minor;
    int clockRate;
    size_t textureAlignment;
    int deviceOverlap;
    int multiProcessorCount;
    int kernelExecTimeoutEnabled;
    int integrated;
    int canMapHostMemory;
    int computeMode;
    int concurrentKernels;
    int ECCEnabled;
    int pciBusID;
    int pciDeviceID;
    int tccDriver;
};

extern "C" {
    size_t cudaMalloc(void **p_mem, size_t N);
    size_t cudaFree(void *mem);
    size_t cudaMemcpy(void *dst, const void *, size_t, size_t cudaMemcpyKind);
    size_t cudaSetDevice (int device);
    size_t cudaGetDeviceProperties (struct cudaDeviceProp *prop, int device);
    size_t cudaMemsetAsync(void *devPtr, int value, size_t count, void *stream = 0);
    const char *cudaGetErrorString (size_t error);
    size_t cudaGetDevice (int *device);
    size_t cudaGetDeviceCount (int *count);
    size_t cudaMemcpyAsync (void *dst, const void *src, size_t count, size_t kind, void *stream = 0);
    size_t cudaStreamSynchronize(void *stream);
    size_t cudaGetLastError();
    size_t cudaConfigureCall(
        unsigned long long grid_xy, unsigned int grid_z,
        unsigned long long block_xy, unsigned int block_z, size_t sharedMem=0, void *stream=0);
    void configureKernel(
        const char *kernelName, const char *clSourcecodeString);
}

size_t cudaGetLastError() {
    cout << "cudaGetLastError" << endl;
    return 0;
}

size_t cudaStreamSynchronize(void *stream) {
    cout << "cudaStreamSynchronize stream=" << stream << endl;
    cl->finish();
    return 0;
}

size_t cudaMemcpyAsync (void *dst, const void *src, size_t count, size_t kind, void *stream) {
    cout << "cudaMemcpyAsync count=" << count << " kind=" << kind << " stream=" << stream << endl;
    return 0;
}

const char *cudaGetErrorString (size_t error) {
    cout << "cudaGetErrorString error=" << error << endl;
    return "all was ok?";
}

size_t cudaGetDevice (int *device) {
    cout << "cudaGetDevice" << endl;
    *device = 0;
    return 0;
}

size_t cudaGetDeviceCount (int *count) {
    cout << "cudaGetDeviceCount" << endl;
    *count = 1;
    return 0;
}

size_t cudaMemsetAsync(void *devPtr, int value, size_t count, void *stream) {
    cout << "cudaMemsetAsync stub value=" << value << " count=" << count << " stream=" << stream << endl;
    return 0;
}

size_t cudaSetDevice (int device) {
    cout << "cudaSetDevice stub device=" << device << endl;
    return 0;
}
<commit_msg>add impllementation for cudaMemcpyAsync<commit_after>// Copyright <NAME> 2016

// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at

//     http://www.apache.org/licenses/LICENSE-2.0

// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.

#include <iostream>
#include <memory>
#include <vector>
#include <map>

#include "EasyCL.h"

#include "CL/cl.h"

using namespace std;

static size_t grid[3];
static size_t block[3];
static unique_ptr<CLKernel> kernel;

unique_ptr<EasyCL> cl;
static cl_context *ctx;
static cl_command_queue *queue;
static cl_int err;

static vector<cl_mem> clmems;
static vector<cl_mem> kernelArgsToBeReleased;
static map<void *, int> idxByAddr;

static bool initialized = false;

void hostside_opencl_funcs_init() {
    cout << "initialize cl context" << endl;
    cl.reset(EasyCL::createForFirstGpuOtherwiseCpu());
    ctx = cl->context;
    queue = cl->queue;
}

void assure_initialized() {
    // yes this is not threadsafe.  or anything safe really...
    if(!initialized) {
        hostside_opencl_funcs_init();
        initialized = true;
    }
}

struct cudaDeviceProp {
    char name[256];
    size_t totalGlobalMem;
    size_t sharedMemPerBlock;
    int regsPerBlock;
    int warpSize;
    size_t memPitch;
    int maxThreadsPerBlock;
    int maxThreadsDim[3];
    int maxGridSize[3];
    size_t totalConstMem;
    int major;
    int minor;
    int clockRate;
    size_t textureAlignment;
    int deviceOverlap;
    int multiProcessorCount;
    int kernelExecTimeoutEnabled;
    int integrated;
    int canMapHostMemory;
    int computeMode;
    int concurrentKernels;
    int ECCEnabled;
    int pciBusID;
    int pciDeviceID;
    int tccDriver;
};

extern "C" {
    size_t cudaMalloc(void **p_mem, size_t N);
    size_t cudaFree(void *mem);
    size_t cudaMemcpy(void *dst, const void *, size_t, size_t cudaMemcpyKind);
    size_t cudaSetDevice (int device);
    size_t cudaGetDeviceProperties (struct cudaDeviceProp *prop, int device);
    size_t cudaMemsetAsync(void *devPtr, int value, size_t count, void *stream = 0);
    const char *cudaGetErrorString (size_t error);
    size_t cudaGetDevice (int *device);
    size_t cudaGetDeviceCount (int *count);
    size_t cudaMemcpyAsync (void *dst, const void *src, size_t count, size_t kind, void *stream = 0);
    size_t cudaStreamSynchronize(void *stream);
    size_t cudaGetLastError();
    size_t cudaConfigureCall(
        unsigned long long grid_xy, unsigned int grid_z,
        unsigned long long block_xy, unsigned int block_z, size_t sharedMem=0, void *stream=0);
    void configureKernel(
        const char *kernelName, const char *clSourcecodeString);
}

size_t cudaGetLastError() {
    cout << "cudaGetLastError" << endl;
    return 0;
}

size_t cudaStreamSynchronize(void *stream) {
    cout << "cudaStreamSynchronize stream=" << stream << endl;

    assert(stream == 0);

    cl->finish();
    return 0;
}

size_t cudaMemcpyAsync (void *dst, const void *src, size_t count, size_t cudaMemcpyKind, void *stream) {
    cout << "cudaMemcpyAsync count=" << count << " cudaMemcpyKind=" << cudaMemcpyKind << " stream=" << stream << endl;

    assert(stream == 0);

    cout << "cudamempcy using opencl cudaMemcpyKind " << cudaMemcpyKind << " count=" << count << endl;
    if(cudaMemcpyKind == 2) {
        // device => host
        int srcidx = idxByAddr[(void *)src];
        cl_mem srcclmem = clmems[srcidx];
        err = clEnqueueReadBuffer(*queue, srcclmem, CL_TRUE, 0,
                                         count, dst, 0, NULL, NULL);
        cl->checkError(err);
        // cl->finish();
    } else if(cudaMemcpyKind == 1) {
        // host => device
        int dstidx = idxByAddr[(void *)dst];
        cl_mem dstclmem = clmems[dstidx];
        err = clEnqueueWriteBuffer(*queue, dstclmem, CL_TRUE, 0,
                                          count, src, 0, NULL, NULL);
        cl->checkError(err);
    } else {
        cout << "cudaMemcpyKind using opencl " << cudaMemcpyKind << endl;
        throw runtime_error("unhandled cudaMemcpyKind");
    }

    return 0;
}

const char *cudaGetErrorString (size_t error) {
    cout << "cudaGetErrorString error=" << error << endl;
    return "all was ok?";
}

size_t cudaGetDevice (int *device) {
    cout << "cudaGetDevice" << endl;
    *device = 0;
    return 0;
}

size_t cudaGetDeviceCount (int *count) {
    cout << "cudaGetDeviceCount" << endl;
    *count = 1;
    return 0;
}

size_t cudaMemsetAsync(void *devPtr, int value, size_t count, void *stream) {
    cout << "cudaMemsetAsync stub value=" << value << " count=" << count << " stream=" << stream << endl;
    assert(stream == 0);

    return 0;
}

size_t cudaSetDevice (int device) {
    cout << "cudaSetDevice stub device=" << device << endl;
    return 0;
}

{'fraction_non_alphanumeric': 0.10833512814990308, 'fraction_numerical': 0.006999784622011631, 'mean_word_length': 3.4435406698564592, 'pattern_counts': {'":': 0, '<': 163, '<?xml version=': 0, '>': 34, 'https://': 0, 'lorem ipsum': 0, 'www.': 2, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 4, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '4985719', 'n_tokens_mistral': 3053, 'n_tokens_neox': 2997, 'n_words': 914}