<commit_before>#ifndef _MY_CUDA_H_
#define _MY_CUDA_H_

#include <cmath>


#include <sys/time.h>


#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <string>
#include <fstream>
#include <cassert>
#include <vector>
#include <cuda_runtime.h>

#include <iostream>

//Some easy to use typedefs
typedef float4 real4;
typedef float real;

using namespace std;




#define cl_mem void*


#  define CU_SAFE_CALL_KERNEL( call , kernel )       CU_SAFE_CALL(call);

// This will output the proper CUDA error strings in the event that a CUDA host call returns an error                                      
#define CU_SAFE_CALL(err)  __checkCudaErrors (err, __FILE__, __LINE__)                                                                  
                                                                                                                                           
inline void __checkCudaErrors(cudaError err, const char *file, const int line )                                                            
{                                                                                                                                          
    if(cudaSuccess != err)                                                                                                                 
    {                                                                                                                                      
        fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.\n",file, line, (int)err, cudaGetErrorString( err ) );                     
        exit(-1);                                                                                                                          
    }                                                                                                                                      
}                                                                                                                                          
                                                                                                                                           
// This will output the proper error string when calling cudaGetLastError                                                                  
#define getLastCudaError(msg)      __getLastCudaError (msg, __FILE__, __LINE__)                                                            
                                                                                                                                           
inline void __getLastCudaError(const char *errorMessage, const char *file, const int line )                                                
{                                                                                                                                          
    cudaError_t err = cudaGetLastError();                                                                                                  
    if (cudaSuccess != err)                                                                                                                
    {                                                                                                                                      
        fprintf(stderr, "%s(%i) : getLastCudaError() CUDA error : %s : (%d) %s.\n",                                                        
        file, line, errorMessage, (int)err, cudaGetErrorString( err ) );                                                                   
        exit(-1);                                                                                                                          
    }                                                                                                                                      
}                                                                                                                                          
                                                                                                                                           
// end of CUDA Helper Functions  



//OpenCL to CUDA macro / functions
__inline__ cudaError_t clFinish(int param)
{
  #if CUDART_VERSION >= 4000                                                                                                                 
          return cudaDeviceSynchronize();                                                                                                    
  #else                                                                                                                                      
          return cudaThreadSynchronize();                                                                                                    
  #endif  
}

    
static int getNumberOfCUDADevices()
{
  // Get number of devices supporting CUDA
  int temp = 0;
  CU_SAFE_CALL(cudaGetDeviceCount(&temp));
  return temp;
}


namespace my_dev {

  class context {
  protected:
    size_t dev;
  
   
    int ciDeviceCount;   
    int ciErrNum;

    bool hContext_flag;
    bool hInit_flag;
    bool logfile_flag;
    bool disable_timing;
    
    ofstream *logFile;
    
    int logID;  //Unique ID to every log line
    
    
    //Events:
    cudaEvent_t start, stop;
    
    //Compute capabilty, important for default compilation mode
    int ccMajor;
    int ccMinor;
    int defaultComputeMode;   

    
    
  public:
    
     int multiProcessorCount;   //Required to configure parts of the code

    context() {
      hContext_flag     = false;     
      hInit_flag        = false;
      logfile_flag      = false;
      disable_timing    = false;
      
      hInit_flag        = true;                 
    }
    ~context() {
      if (hContext_flag )
      {
      }
    }
     
    
    int create(std::ofstream &log, bool disableTiming = false)
    {
      disable_timing = disableTiming;
      logfile_flag = true;
      logFile = &log;
      logID = 0;
      return create(disable_timing); 
    }


    int create(bool disableT = false) {
      assert(hInit_flag);     
      
      disable_timing = disableT;
      
      printf("Creating CUDA context \n");
            
      // Get number of devices supporting CUDA
      ciDeviceCount = 0;
      CU_SAFE_CALL(cudaGetDeviceCount(&ciDeviceCount));
            
      printf("Found %d suitable devices: \n",ciDeviceCount);
      for(int i=0; i < ciDeviceCount; i++)
      {
        cudaDeviceProp deviceProp;   
        cudaGetDeviceProperties(&deviceProp, i);     
        printf(" %d: %s\n",i, deviceProp.name);
      }

      return ciDeviceCount;
    }
    
    void createQueue(size_t dev = 0, int ctxCreateFlags = 0) 
    {
      //use CU_CTX_MAP_HOST as flag for zero-copy memory
      //Here we finally create and assign the context to a device
      assert(!hContext_flag);
      assert(hInit_flag);
      this->dev = dev;
      assert((int)dev < ciDeviceCount);
                  
      printf("Trying to use device: %d ...", (int)dev);
      //Faster and async kernel launches when using large size arrays of local memory
      //ctxCreateFlags |= CU_CTX_LMEM_RESIZE_TO_MAX;

      //Create the context for this device handle
      //CU_SAFE_CALL(cuCtxCreate(&hContext, ctxCreateFlags, hDevice));
      
      int res = cudaSetDevice(dev);
      if(res != cudaSuccess)
      {
       printf("failed (error #: %d), now trying all devices starting at 0 \n", res);

	for(int i=0; i < ciDeviceCount; i++)
	{
		printf("Trying device: %d  ...", i);
                if(cudaSetDevice(i) != cudaSuccess)
		{
			printf("failed!\n");
			if(i+1 == ciDeviceCount)
			{
  			  printf("All devices failed, exit! \n");
			  exit(0);
			}
		}
		else
		{
			printf("success! \n");
                        this->dev = i;
			break;
		}
	}
      }
      else
      {
	      printf("success!\n");
      }
      
      cudaDeviceProp deviceProp;                                                                                                     
      CU_SAFE_CALL(cudaGetDeviceProperties(&deviceProp, dev));
      //Get the number of multiprocessors of the device
      multiProcessorCount = deviceProp.multiProcessorCount;
      
      hContext_flag = true;
    }
    
    
    void startTiming(cudaStream_t stream=0)
    {
      if(disable_timing) return;
      int eventflags =  cudaEventDefault;      
      CU_SAFE_CALL(cudaEventCreateWithFlags(&start, eventflags));  
      CU_SAFE_CALL(cudaEventCreateWithFlags(&stop, eventflags));
      CU_SAFE_CALL(cudaEventRecord(start, stream));
    }
    
    //Text and ID to be printed with the log message on screen / in the file
    void stopTiming(const char *text, int type = -1, cudaStream_t stream=0)
    {
      if(disable_timing) return;
      
      CU_SAFE_CALL(cudaEventRecord(stop, stream));
      CU_SAFE_CALL(cudaEventSynchronize(stop));
      float time;
      CU_SAFE_CALL(cudaEventElapsedTime(&time, start, stop));      
      CU_SAFE_CALL(cudaEventDestroy(start));
      CU_SAFE_CALL(cudaEventDestroy(stop));      
      
      printf("%s took:\t%f\t millisecond\n", text, time);
      
      if(logfile_flag)
      {
        (*logFile) << logID++ << "\t"  << type << "\t" << text << "\t" << time << endl;
      }
    }
    
    /////////////
    //Kept for compatability
    int              get_command_queue() {return 0;}
    //////////       
    
  };
  
  
  ////////////////////////////////////////
  
    //Class to handle streams / queues
  
  class dev_stream
  {
    private:
      cudaStream_t stream;
      
    public:     
      dev_stream(unsigned int flags = 0)
      {
        createStream(flags);
      }
            
      
      void createStream(unsigned int flags = 0)
      {
         CU_SAFE_CALL(cudaStreamCreate(&stream));
      }   
      
      void destroyStream()
      {
        CU_SAFE_CALL(cudaStreamDestroy(stream));
      }
      
      void sync()
      {
        CU_SAFE_CALL(cudaStreamSynchronize(stream));        
      }
      
      bool isFinished()
      {
        cudaError_t res = cudaStreamQuery(stream);
        if(res == cudaSuccess) return true;
        if(res == cudaErrorNotReady) return false;
        
        //Some other result
        CU_SAFE_CALL(res);
        return false;
      }
      
      cudaStream_t s()
      {
        return stream;
      }
      
      
      ~dev_stream() {
      destroyStream();
    }
  };
  
    
  ///////////////////////
  
  class base_mem
  {
    public:     
    //Memory usage counters
    static long long currentMemUsage;
    static long long maxMemUsage;  
    
    void increaseMemUsage(int bytes)
    {
      currentMemUsage +=  bytes;   
      
      if(currentMemUsage > maxMemUsage)
        maxMemUsage = currentMemUsage;
    }
    
    void decreaseMemUsage(int bytes)
    {
      currentMemUsage -=  bytes;
    }
    
    static void printMemUsage()
    {      
      printf("Current usage: %lld bytes ( %lld MB) \n", currentMemUsage, currentMemUsage / (1024*1024));
      printf("Maximum usage: %lld bytes ( %lld MB) \n", maxMemUsage, maxMemUsage / (1024*1024));     
      
      size_t free, total;
      cudaMemGetInfo(&free, &total); 
      printf("Build-in usage: free: %ld bytes ( %ld MB , total: %ld) \n", free, free / (1024*1024), total / (1024*1024));     
      
    }  
    
    static long long getMaxMemUsage()
    {      
      return maxMemUsage;
    }   
    
  };


  template<class T>
  class dev_mem : base_mem {
  protected:
    
    typedef struct textureInfo
    {
      const struct textureReference *texture;
      int      texOffset; //The possible extra offset when using textures and combined memory
      int      texSize;
    } textureInfo;        
        
 
    vector<textureInfo> textures;
    
    int size;
    T           *hDeviceMem;
    T           *host_ptr;
    void        *DeviceMemPtr;
    void        *tempDeviceMemPtr;
        
    bool pinned_mem, context_flag, flags;
    bool hDeviceMem_flag;
    bool childMemory; //Indicates that this is a shared buffer that will be freed by a parent
    
    void cuda_free() {      
      if(childMemory) //Only free if we are NOT a child
      {
        return;
      }

      
//       assert(context_flag);
      if (hDeviceMem_flag)
      {
	assert(size > 0);
	CU_SAFE_CALL(cudaFree(hDeviceMem));
        decreaseMemUsage(size*sizeof(T));
        
        if(pinned_mem){
          CU_SAFE_CALL(cudaFreeHost((void*)host_ptr));}
        else{
          free(host_ptr);}
          hDeviceMem_flag = false;
      }
    } //cuda_free

  public:
    

    ///////// Constructors

    dev_mem() {
      size              = 0;
      pinned_mem        = false;
      hDeviceMem_flag   = false;
      context_flag      = false;
      host_ptr          = NULL;
      childMemory       = false;
    }

    dev_mem(class context &c) {
      size              = 0;      
      pinned_mem        = false;
      context_flag      = false;
      hDeviceMem_flag   = false;
      host_ptr          = NULL;
      childMemory       = false;
      setContext(c);
    }
    
    //CUDA has no memory flags like opencl
    //so just put it to 0 and keep function format same for 
    //compatability
    dev_mem(class context &c, int n, bool zero = false,
	    int flags = 0, bool pinned = false) {
      context_flag      = false;
      childMemory       = false;      
      hDeviceMem_flag   = false;
      pinned_mem        = pinned;
      size              = 0;
      setContext(c);
      if (zero) this->ccalloc(n, pinned, flags);
      else      this->cmalloc(n, pinned, flags);
    }
    
//     dev_mem(class context &c, std::vector<T> data,
// 	    int flags = 0,  bool pinned = false) {
//       context_flag    = false;
//       hDeviceMem_flag = false;
//       childMemory       = false;      
//       pinned_mem      = pinned;
//       size            = 0;
//       setContext(c);
//       this->cmalloc(data, flags);
//     }
    
    void free_mem()
    {
      cuda_free();
    }

    //////// Destructor
    
    ~dev_mem() {      
        cuda_free();
    }
    
    ///////////

    void setContext(class context &c) {      
      context_flag     = true;
    }


    ///////////
    
    //Get the reference of memory allocated by another piece of memory
    
    void cmalloc_copy(bool pinned, bool flags, void *cudaMem, 
                      void* ParentHost_ptr, int offset, int n,
                      int allignOffset)
    {
      assert(context_flag);
   
      this->pinned_mem  = pinned;
      this->flags       = flags;
      this->childMemory = true;

      size = n;
        
      //Dont forget to add the allignment values      
      //The following line has a bug, it first casts and then adds offset*sizeof ELEMENTS
      //host_ptr = (T*)ParentHost_ptr +  allignOffset*sizeof(uint); 
      //This line correctly increases the memory location with number of bytes before casting
      host_ptr = (T*) ((char*)ParentHost_ptr +  allignOffset*sizeof(uint));

#if 0 /* egaburov: to fix void* pointer arithmetic warrning */
      hDeviceMem   = (T*)(cudaMem + offset*sizeof(uint) + allignOffset*sizeof(uint));
#else
      hDeviceMem   = (T*)cudaMem + offset*sizeof(uint) + allignOffset*sizeof(uint);
#endif
      DeviceMemPtr = (void*)(size_t)(hDeviceMem);

      hDeviceMem_flag = true;      
    }

    void cmalloc(int n, bool pinned = false, int flags = 0) 
    {
      assert(context_flag);
//       assert(!hDeviceMem_flag);
      this->pinned_mem = pinned;      
      this->flags = flags;
      if (size > 0) cuda_free();
      size = n;
            
      if(pinned_mem){    
        CU_SAFE_CALL(cudaMallocHost((T**)&host_ptr, size*sizeof(T)));}
      else{
        host_ptr = (T*)malloc(size*sizeof(T));}
        
      CU_SAFE_CALL(cudaMalloc((T**)&hDeviceMem, size*sizeof(T)));
      increaseMemUsage(size*sizeof(T));
      DeviceMemPtr = (void*)(size_t)hDeviceMem;    

      hDeviceMem_flag = true;
    }

    void ccalloc(int n, bool pinned = false, int flags = 0) {
      assert(context_flag);<commit_msg>Fixed the void* warning fix<commit_after>#ifndef _MY_CUDA_H_
#define _MY_CUDA_H_

#include <cmath>


#include <sys/time.h>


#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <string>
#include <fstream>
#include <cassert>
#include <vector>
#include <cuda_runtime.h>

#include <iostream>

//Some easy to use typedefs
typedef float4 real4;
typedef float real;

using namespace std;




#define cl_mem void*


#  define CU_SAFE_CALL_KERNEL( call , kernel )       CU_SAFE_CALL(call);

// This will output the proper CUDA error strings in the event that a CUDA host call returns an error                                      
#define CU_SAFE_CALL(err)  __checkCudaErrors (err, __FILE__, __LINE__)                                                                  
                                                                                                                                           
inline void __checkCudaErrors(cudaError err, const char *file, const int line )                                                            
{                                                                                                                                          
    if(cudaSuccess != err)                                                                                                                 
    {                                                                                                                                      
        fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.\n",file, line, (int)err, cudaGetErrorString( err ) );                     
        exit(-1);                                                                                                                          
    }                                                                                                                                      
}                                                                                                                                          
                                                                                                                                           
// This will output the proper error string when calling cudaGetLastError                                                                  
#define getLastCudaError(msg)      __getLastCudaError (msg, __FILE__, __LINE__)                                                            
                                                                                                                                           
inline void __getLastCudaError(const char *errorMessage, const char *file, const int line )                                                
{                                                                                                                                          
    cudaError_t err = cudaGetLastError();                                                                                                  
    if (cudaSuccess != err)                                                                                                                
    {                                                                                                                                      
        fprintf(stderr, "%s(%i) : getLastCudaError() CUDA error : %s : (%d) %s.\n",                                                        
        file, line, errorMessage, (int)err, cudaGetErrorString( err ) );                                                                   
        exit(-1);                                                                                                                          
    }                                                                                                                                      
}                                                                                                                                          
                                                                                                                                           
// end of CUDA Helper Functions  



//OpenCL to CUDA macro / functions
__inline__ cudaError_t clFinish(int param)
{
  #if CUDART_VERSION >= 4000                                                                                                                 
          return cudaDeviceSynchronize();                                                                                                    
  #else                                                                                                                                      
          return cudaThreadSynchronize();                                                                                                    
  #endif  
}

    
static int getNumberOfCUDADevices()
{
  // Get number of devices supporting CUDA
  int temp = 0;
  CU_SAFE_CALL(cudaGetDeviceCount(&temp));
  return temp;
}


namespace my_dev {

  class context {
  protected:
    size_t dev;
  
   
    int ciDeviceCount;   
    int ciErrNum;

    bool hContext_flag;
    bool hInit_flag;
    bool logfile_flag;
    bool disable_timing;
    
    ofstream *logFile;
    
    int logID;  //Unique ID to every log line
    
    
    //Events:
    cudaEvent_t start, stop;
    
    //Compute capabilty, important for default compilation mode
    int ccMajor;
    int ccMinor;
    int defaultComputeMode;   

    
    
  public:
    
     int multiProcessorCount;   //Required to configure parts of the code

    context() {
      hContext_flag     = false;     
      hInit_flag        = false;
      logfile_flag      = false;
      disable_timing    = false;
      
      hInit_flag        = true;                 
    }
    ~context() {
      if (hContext_flag )
      {
      }
    }
     
    
    int create(std::ofstream &log, bool disableTiming = false)
    {
      disable_timing = disableTiming;
      logfile_flag = true;
      logFile = &log;
      logID = 0;
      return create(disable_timing); 
    }


    int create(bool disableT = false) {
      assert(hInit_flag);     
      
      disable_timing = disableT;
      
      printf("Creating CUDA context \n");
            
      // Get number of devices supporting CUDA
      ciDeviceCount = 0;
      CU_SAFE_CALL(cudaGetDeviceCount(&ciDeviceCount));
            
      printf("Found %d suitable devices: \n",ciDeviceCount);
      for(int i=0; i < ciDeviceCount; i++)
      {
        cudaDeviceProp deviceProp;   
        cudaGetDeviceProperties(&deviceProp, i);     
        printf(" %d: %s\n",i, deviceProp.name);
      }

      return ciDeviceCount;
    }
    
    void createQueue(size_t dev = 0, int ctxCreateFlags = 0) 
    {
      //use CU_CTX_MAP_HOST as flag for zero-copy memory
      //Here we finally create and assign the context to a device
      assert(!hContext_flag);
      assert(hInit_flag);
      this->dev = dev;
      assert((int)dev < ciDeviceCount);
                  
      printf("Trying to use device: %d ...", (int)dev);
      //Faster and async kernel launches when using large size arrays of local memory
      //ctxCreateFlags |= CU_CTX_LMEM_RESIZE_TO_MAX;

      //Create the context for this device handle
      //CU_SAFE_CALL(cuCtxCreate(&hContext, ctxCreateFlags, hDevice));
      
      int res = cudaSetDevice(dev);
      if(res != cudaSuccess)
      {
       printf("failed (error #: %d), now trying all devices starting at 0 \n", res);

	for(int i=0; i < ciDeviceCount; i++)
	{
		printf("Trying device: %d  ...", i);
                if(cudaSetDevice(i) != cudaSuccess)
		{
			printf("failed!\n");
			if(i+1 == ciDeviceCount)
			{
  			  printf("All devices failed, exit! \n");
			  exit(0);
			}
		}
		else
		{
			printf("success! \n");
                        this->dev = i;
			break;
		}
	}
      }
      else
      {
	      printf("success!\n");
      }
      
      cudaDeviceProp deviceProp;                                                                                                     
      CU_SAFE_CALL(cudaGetDeviceProperties(&deviceProp, dev));
      //Get the number of multiprocessors of the device
      multiProcessorCount = deviceProp.multiProcessorCount;
      
      hContext_flag = true;
    }
    
    
    void startTiming(cudaStream_t stream=0)
    {
      if(disable_timing) return;
      int eventflags =  cudaEventDefault;      
      CU_SAFE_CALL(cudaEventCreateWithFlags(&start, eventflags));  
      CU_SAFE_CALL(cudaEventCreateWithFlags(&stop, eventflags));
      CU_SAFE_CALL(cudaEventRecord(start, stream));
    }
    
    //Text and ID to be printed with the log message on screen / in the file
    void stopTiming(const char *text, int type = -1, cudaStream_t stream=0)
    {
      if(disable_timing) return;
      
      CU_SAFE_CALL(cudaEventRecord(stop, stream));
      CU_SAFE_CALL(cudaEventSynchronize(stop));
      float time;
      CU_SAFE_CALL(cudaEventElapsedTime(&time, start, stop));      
      CU_SAFE_CALL(cudaEventDestroy(start));
      CU_SAFE_CALL(cudaEventDestroy(stop));      
      
      printf("%s took:\t%f\t millisecond\n", text, time);
      
      if(logfile_flag)
      {
        (*logFile) << logID++ << "\t"  << type << "\t" << text << "\t" << time << endl;
      }
    }
    
    /////////////
    //Kept for compatability
    int              get_command_queue() {return 0;}
    //////////       
    
  };
  
  
  ////////////////////////////////////////
  
    //Class to handle streams / queues
  
  class dev_stream
  {
    private:
      cudaStream_t stream;
      
    public:     
      dev_stream(unsigned int flags = 0)
      {
        createStream(flags);
      }
            
      
      void createStream(unsigned int flags = 0)
      {
         CU_SAFE_CALL(cudaStreamCreate(&stream));
      }   
      
      void destroyStream()
      {
        CU_SAFE_CALL(cudaStreamDestroy(stream));
      }
      
      void sync()
      {
        CU_SAFE_CALL(cudaStreamSynchronize(stream));        
      }
      
      bool isFinished()
      {
        cudaError_t res = cudaStreamQuery(stream);
        if(res == cudaSuccess) return true;
        if(res == cudaErrorNotReady) return false;
        
        //Some other result
        CU_SAFE_CALL(res);
        return false;
      }
      
      cudaStream_t s()
      {
        return stream;
      }
      
      
      ~dev_stream() {
      destroyStream();
    }
  };
  
    
  ///////////////////////
  
  class base_mem
  {
    public:     
    //Memory usage counters
    static long long currentMemUsage;
    static long long maxMemUsage;  
    
    void increaseMemUsage(int bytes)
    {
      currentMemUsage +=  bytes;   
      
      if(currentMemUsage > maxMemUsage)
        maxMemUsage = currentMemUsage;
    }
    
    void decreaseMemUsage(int bytes)
    {
      currentMemUsage -=  bytes;
    }
    
    static void printMemUsage()
    {      
      printf("Current usage: %lld bytes ( %lld MB) \n", currentMemUsage, currentMemUsage / (1024*1024));
      printf("Maximum usage: %lld bytes ( %lld MB) \n", maxMemUsage, maxMemUsage / (1024*1024));     
      
      size_t free, total;
      cudaMemGetInfo(&free, &total); 
      printf("Build-in usage: free: %ld bytes ( %ld MB , total: %ld) \n", free, free / (1024*1024), total / (1024*1024));     
      
    }  
    
    static long long getMaxMemUsage()
    {      
      return maxMemUsage;
    }   
    
  };


  template<class T>
  class dev_mem : base_mem {
  protected:
    
    typedef struct textureInfo
    {
      const struct textureReference *texture;
      int      texOffset; //The possible extra offset when using textures and combined memory
      int      texSize;
    } textureInfo;        
        
 
    vector<textureInfo> textures;
    
    int size;
    T           *hDeviceMem;
    T           *host_ptr;
    void        *DeviceMemPtr;
    void        *tempDeviceMemPtr;
        
    bool pinned_mem, context_flag, flags;
    bool hDeviceMem_flag;
    bool childMemory; //Indicates that this is a shared buffer that will be freed by a parent
    
    void cuda_free() {      
      if(childMemory) //Only free if we are NOT a child
      {
        return;
      }

      
//       assert(context_flag);
      if (hDeviceMem_flag)
      {
	assert(size > 0);
	CU_SAFE_CALL(cudaFree(hDeviceMem));
        decreaseMemUsage(size*sizeof(T));
        
        if(pinned_mem){
          CU_SAFE_CALL(cudaFreeHost((void*)host_ptr));}
        else{
          free(host_ptr);}
          hDeviceMem_flag = false;
      }
    } //cuda_free

  public:
    

    ///////// Constructors

    dev_mem() {
      size              = 0;
      pinned_mem        = false;
      hDeviceMem_flag   = false;
      context_flag      = false;
      host_ptr          = NULL;
      childMemory       = false;
    }

    dev_mem(class context &c) {
      size              = 0;      
      pinned_mem        = false;
      context_flag      = false;
      hDeviceMem_flag   = false;
      host_ptr          = NULL;
      childMemory       = false;
      setContext(c);
    }
    
    //CUDA has no memory flags like opencl
    //so just put it to 0 and keep function format same for 
    //compatability
    dev_mem(class context &c, int n, bool zero = false,
	    int flags = 0, bool pinned = false) {
      context_flag      = false;
      childMemory       = false;      
      hDeviceMem_flag   = false;
      pinned_mem        = pinned;
      size              = 0;
      setContext(c);
      if (zero) this->ccalloc(n, pinned, flags);
      else      this->cmalloc(n, pinned, flags);
    }
    
//     dev_mem(class context &c, std::vector<T> data,
// 	    int flags = 0,  bool pinned = false) {
//       context_flag    = false;
//       hDeviceMem_flag = false;
//       childMemory       = false;      
//       pinned_mem      = pinned;
//       size            = 0;
//       setContext(c);
//       this->cmalloc(data, flags);
//     }
    
    void free_mem()
    {
      cuda_free();
    }

    //////// Destructor
    
    ~dev_mem() {      
        cuda_free();
    }
    
    ///////////

    void setContext(class context &c) {      
      context_flag     = true;
    }


    ///////////
    
    //Get the reference of memory allocated by another piece of memory
    
    void cmalloc_copy(bool pinned, bool flags, void *cudaMem, 
                      void* ParentHost_ptr, int offset, int n,
                      int allignOffset)
    {
      assert(context_flag);
   
      this->pinned_mem  = pinned;
      this->flags       = flags;
      this->childMemory = true;

      size = n;
        
      //Dont forget to add the allignment values      
      //The following line has a bug, it first casts and then adds offset*sizeof ELEMENTS
      //host_ptr = (T*)ParentHost_ptr +  allignOffset*sizeof(uint); 
      //This line correctly increases the memory location with number of bytes before casting
      host_ptr = (T*) ((char*)ParentHost_ptr +  allignOffset*sizeof(uint));

      
#if 1 /* egaburov: to fix void* pointer arithmetic warrning .  */
      /* jbedorf: fixed so address is correct, casting before adding gives a different
       * result. So divide by size of object */
      //       hDeviceMem   = (T*)(cudaMem + offset*sizeof(uint) + allignOffset*sizeof(uint));
      hDeviceMem   = (T*)cudaMem + ((offset*sizeof(uint) + allignOffset*sizeof(uint)) / sizeof(T));
#else
      hDeviceMem   = (T*)cudaMem + offset*sizeof(uint) + allignOffset*sizeof(uint);
      
#endif
      DeviceMemPtr = (void*)(size_t)(hDeviceMem);

      hDeviceMem_flag = true;      
    }

    void cmalloc(int n, bool pinned = false, int flags = 0) 
    {
      assert(context_flag);
//       assert(!hDeviceMem_flag);
      this->pinned_mem = pinned;      
      this->flags = flags;
      if (size > 0) cuda_free();
      size = n;
            
      if(pinned_mem){    
        CU_SAFE_CALL(cudaMallocHost((T**)&host_ptr, size*sizeof(T)));}
      else{
        host_ptr = (T*)malloc(size*sizeof(T));}
        
      CU_SAFE_CALL(cudaMalloc((T**)&hDeviceMem, size*sizeof(T)));
      increaseMemUsage(size*sizeof(T));
      DeviceMemPtr = (void*)(size_t)hDeviceMem;    

      hDeviceMem_flag = true;
    }

    void ccalloc(int n, bool pinned = false, int flags = 0) {
      assert(context_flag);
{'fraction_non_alphanumeric': 0.08884647354000554, 'fraction_numerical': 0.004186805405904627, 'mean_word_length': 1.1127804878048781, 'pattern_counts': {'":': 0, '<': 69, '<?xml version=': 0, '>': 59, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 266, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '2201897', 'n_tokens_mistral': 8606, 'n_tokens_neox': 8186, 'n_words': 2001}