COMPASS  5.0.0
End-to-end AO simulation tool using GPU acceleration
CarmaObj Class Reference

this class provides wrappers to the generic carma object More...

#include <carma_host_obj.h>

Inheritance diagram for CarmaObj:
Collaboration diagram for CarmaObj:

Public Member Functions

 CarmaObj (const CarmaObj< T_data > *obj)
 
 CarmaObj (CarmaContext *current_context, const long *dims_data)
 
 CarmaObj (CarmaContext *current_context, const CarmaObj< T_data > *obj)
 
 CarmaObj (CarmaContext *current_context, const long *dims_data, const T_data *data)
 
 CarmaObj (CarmaContext *current_context, const long *dims_data, int nb_streams)
 
 CarmaObj (CarmaContext *current_context, const CarmaObj< T_data > *obj, int nb_streams)
 
 CarmaObj (CarmaContext *current_context, const long *dims_data, const T_data *data, int nb_streams)
 
 ~CarmaObj ()
 
void sync_h_data ()
 
T_data * get_h_data ()
 
int get_nb_streams () const
 
int add_stream ()
 
int add_stream (int nb)
 
int del_stream ()
 
int del_stream (int nb)
 
cudaStream_t get_cuda_stream (int stream)
 
int wait_stream (int stream)
 
int wait_all_streams ()
 
void swap_ptr (T_data *ptr)
 
void dealloc ()
 
 operator T_data * ()
 
std::string to_string ()
 
 operator std::string ()
 
char const * c_str ()
 
const T_data operator[] (int index) const
 
T_data * get_data ()
 
T_data * get_data_at (int index)
 
T_data * get_o_data ()
 
const T_data get_o_data_value () const
 
const long * get_dims ()
 
long get_dims (int i)
 
int get_nb_elements ()
 
CarmaContextget_context ()
 
int get_device ()
 
bool is_rng_init ()
 
template<typename T_dest >
int host2device (const T_dest *data)
 
template<typename T_dest >
int device2host (T_dest *data)
 
int host2device_async (const T_data *data, cudaStream_t stream)
 
int device2host_async (T_data *data, cudaStream_t stream)
 
int device2host_opt (T_data *data)
 
int host2device_vect (const T_data *data, int incx, int incy)
 
int device2host_vect (T_data *data, int incx, int incy)
 
int host2device_mat (const T_data *data, int lda, int ldb)
 
int device2host_mat (T_data *data, int lda, int ldb)
 
int copy_into (T_data *data, int nb_elem)
 
int copy_from (const T_data *data, int nb_elem)
 
int reset ()
 
int memset (T_data value)
 
cufftHandle * get_plan ()
 FFT plan. More...
 
cufftType get_type_plan ()
 FFT plan type. More...
 
unsigned int * get_values ()
 optional data (used for sort) More...
 
T_data sum ()
 
void init_reduceCub ()
 
void reduceCub ()
 
void clip (T_data min, T_data max)
 
int transpose (CarmaObj< T_data > *source)
 
int aimax (int incx)
 
int aimin (int incx)
 
T_data asum (int incx)
 
T_data nrm2 (int incx)
 
T_data dot (CarmaObj< T_data > *source, int incx, int incy)
 
void scale (T_data alpha, int incx)
 
void swap (CarmaObj< T_data > *source, int incx, int incy)
 
void copy (CarmaObj< T_data > *source, int incx, int incy)
 
void axpy (T_data alpha, CarmaObj< T_data > *source, int incx, int incy, int offset=0)
 
void rot (CarmaObj< T_data > *source, int incx, int incy, T_data sc, T_data ss)
 
void gemv (char trans, T_data alpha, CarmaObj< T_data > *matA, int lda, CarmaObj< T_data > *vectx, int incx, T_data beta, int incy)
 
void ger (T_data alpha, CarmaObj< T_data > *vectx, int incx, CarmaObj< T_data > *vecty, int incy, int lda)
 
void symv (char uplo, T_data alpha, CarmaObj< T_data > *matA, int lda, CarmaObj< T_data > *vectx, int incx, T_data beta, int incy)
 
void gemm (char transa, char transb, T_data alpha, CarmaObj< T_data > *matA, int lda, CarmaObj< T_data > *matB, int ldb, T_data beta, int ldc)
 
void symm (char side, char uplo, T_data alpha, CarmaObj< T_data > *matA, int lda, CarmaObj< T_data > *matB, int ldb, T_data beta, int ldc)
 
void syrk (char uplo, char transa, T_data alpha, CarmaObj< T_data > *matA, int lda, T_data beta, int ldc)
 
void syrkx (char uplo, char transa, T_data alpha, CarmaObj< T_data > *matA, int lda, CarmaObj< T_data > *matB, int ldb, T_data beta, int ldc)
 
void geam (char transa, char transb, T_data alpha, CarmaObj< T_data > *matA, int lda, T_data beta, CarmaObj< T_data > *matB, int ldb, int ldc)
 
void dgmm (char side, CarmaObj< T_data > *matA, int lda, CarmaObj< T_data > *vectx, int incx, int ldc)
 
int init_prng ()
 
int init_prng (long seed)
 
int destroy_prng ()
 
int prng (T_data *output, char gtype, float alpha, float beta)
 
int prng (T_data *output, char gtype, float alpha)
 
int prng (char gtype, float alpha, float beta)
 
int prng (char gtype, float alpha)
 
int prng (char gtype)
 
int prng_montagn (float init_montagn)
 
int init_prng_host (int seed)
 
int prng_host (char gtype)
 
int prng_host (char gtype, T_data stddev)
 
int prng_host (char gtype, T_data stddev, T_data alpha)
 
int destroy_prng_host ()
 

Protected Member Functions

void init (CarmaContext *current_context, const long *dims_data, const T_data *data, bool fromHost, int nb_streams)
 

Protected Attributes

T_data * d_data
 Input data => change to vector. More...
 
std::vector< T_data > h_data
 
T_data * o_data
 optional data (used for scan / reduction) More...
 
T_data * cub_data
 optional data (used for scan / reduction) More...
 
size_t cub_data_size
 
int ndim
 
long * dims_data
 dimensions of the array More...
 
int nb_elem
 number of elements in the array More...
 
int device
 device where the CarmaObj is allocate More...
 
CarmaContextcurrent_context
 
curandGenerator_t gen
 
curandState * d_states
 
int nb_threads
 
int nb_blocks
 
bool keys_only
 
bool owner = true
 
unsigned int * values
 optional data (used for sort) More...
 
size_t * d_num_valid
 used for compact More...
 
cufftHandle plan
 FFT plan. More...
 
cufftType type_plan
 FFT plan type. More...
 
CarmaStreamsstreams
 

Detailed Description

this class provides wrappers to the generic carma object

Author
COMPASS Team https://github.com/ANR-COMPASS
Version
5.0.0
Date
2011/01/28

Definition at line 65 of file carma_host_obj.h.

Constructor & Destructor Documentation

◆ CarmaObj() [1/7]

CarmaObj::CarmaObj ( const CarmaObj< T_data > *  obj)

◆ CarmaObj() [2/7]

CarmaObj::CarmaObj ( CarmaContext current_context,
const long *  dims_data 
)

◆ CarmaObj() [3/7]

CarmaObj::CarmaObj ( CarmaContext current_context,
const CarmaObj< T_data > *  obj 
)

◆ CarmaObj() [4/7]

CarmaObj::CarmaObj ( CarmaContext current_context,
const long *  dims_data,
const T_data *  data 
)

◆ CarmaObj() [5/7]

CarmaObj::CarmaObj ( CarmaContext current_context,
const long *  dims_data,
int  nb_streams 
)

◆ CarmaObj() [6/7]

CarmaObj::CarmaObj ( CarmaContext current_context,
const CarmaObj< T_data > *  obj,
int  nb_streams 
)

◆ CarmaObj() [7/7]

CarmaObj::CarmaObj ( CarmaContext current_context,
const long *  dims_data,
const T_data *  data,
int  nb_streams 
)

◆ ~CarmaObj()

CarmaObj::~CarmaObj ( )

Member Function Documentation

◆ add_stream() [1/2]

int CarmaObj::add_stream ( )

Definition at line 176 of file carma_obj.h.

176  {
177  this->streams->add_stream();
178  return this->streams->get_nb_streams();
179  }

◆ add_stream() [2/2]

int CarmaObj::add_stream ( int  nb)

Definition at line 180 of file carma_obj.h.

180  {
181  this->streams->add_stream(nb);
182  return this->streams->get_nb_streams();
183  }

◆ aimax()

int CarmaObj::aimax ( int  incx)

◆ aimin()

int CarmaObj::aimin ( int  incx)

◆ asum()

T_data CarmaObj::asum ( int  incx)

◆ axpy()

void CarmaObj::axpy ( T_data  alpha,
CarmaObj< T_data > *  source,
int  incx,
int  incy,
int  offset = 0 
)

◆ c_str()

char const* CarmaObj::c_str ( )

Definition at line 223 of file carma_obj.h.

223 { return this->to_string().c_str(); }

◆ clip()

void CarmaObj::clip ( T_data  min,
T_data  max 
)

transpose

◆ copy()

void CarmaObj::copy ( CarmaObj< T_data > *  source,
int  incx,
int  incy 
)
Here is the caller graph for this function:

◆ copy_from()

int CarmaObj::copy_from ( const T_data *  data,
int  nb_elem 
)

◆ copy_into()

int CarmaObj::copy_into ( T_data *  data,
int  nb_elem 
)

◆ dealloc()

void CarmaObj::dealloc ( )

General Utilities

Definition at line 209 of file carma_obj.h.

Here is the caller graph for this function:

◆ del_stream() [1/2]

int CarmaObj::del_stream ( )

Definition at line 184 of file carma_obj.h.

184  {
185  this->streams->del_stream();
186  return this->streams->get_nb_streams();
187  }

◆ del_stream() [2/2]

int CarmaObj::del_stream ( int  nb)

Definition at line 188 of file carma_obj.h.

188  {
189  this->streams->del_stream(nb);
190  return this->streams->get_nb_streams();
191  }

◆ destroy_prng()

int CarmaObj::destroy_prng ( )

◆ destroy_prng_host()

int CarmaObj::destroy_prng_host ( )

◆ device2host()

template<typename T_dest >
int CarmaObj::device2host ( T_dest *  data)
Here is the caller graph for this function:

◆ device2host_async()

int CarmaObj::device2host_async ( T_data *  data,
cudaStream_t  stream 
)

◆ device2host_mat()

int CarmaObj::device2host_mat ( T_data *  data,
int  lda,
int  ldb 
)

◆ device2host_opt()

int CarmaObj::device2host_opt ( T_data *  data)

◆ device2host_vect()

int CarmaObj::device2host_vect ( T_data *  data,
int  incx,
int  incy 
)

◆ dgmm()

void CarmaObj::dgmm ( char  side,
CarmaObj< T_data > *  matA,
int  lda,
CarmaObj< T_data > *  vectx,
int  incx,
int  ldc 
)

Curand

◆ dot()

T_data CarmaObj::dot ( CarmaObj< T_data > *  source,
int  incx,
int  incy 
)

◆ geam()

void CarmaObj::geam ( char  transa,
char  transb,
T_data  alpha,
CarmaObj< T_data > *  matA,
int  lda,
T_data  beta,
CarmaObj< T_data > *  matB,
int  ldb,
int  ldc 
)

◆ gemm()

void CarmaObj::gemm ( char  transa,
char  transb,
T_data  alpha,
CarmaObj< T_data > *  matA,
int  lda,
CarmaObj< T_data > *  matB,
int  ldb,
T_data  beta,
int  ldc 
)

◆ gemv()

void CarmaObj::gemv ( char  trans,
T_data  alpha,
CarmaObj< T_data > *  matA,
int  lda,
CarmaObj< T_data > *  vectx,
int  incx,
T_data  beta,
int  incy 
)
Here is the caller graph for this function:

◆ ger()

void CarmaObj::ger ( T_data  alpha,
CarmaObj< T_data > *  vectx,
int  incx,
CarmaObj< T_data > *  vecty,
int  incy,
int  lda 
)

◆ get_context()

CarmaContext* CarmaObj::get_context ( )

Definition at line 242 of file carma_obj.h.

242 { return current_context; }
Here is the caller graph for this function:

◆ get_cuda_stream()

cudaStream_t CarmaObj::get_cuda_stream ( int  stream)

Definition at line 192 of file carma_obj.h.

192  {
193  return this->streams->get_stream(stream);
194  }

◆ get_data()

T_data* CarmaObj::get_data ( )

Definition at line 230 of file carma_obj.h.

230 { return d_data; }

◆ get_data_at()

T_data* CarmaObj::get_data_at ( int  index)

Definition at line 231 of file carma_obj.h.

231 { return &d_data[index]; }
Here is the caller graph for this function:

◆ get_device()

int CarmaObj::get_device ( )

Definition at line 244 of file carma_obj.h.

244 { return device; }
Here is the caller graph for this function:

◆ get_dims() [1/2]

const long* CarmaObj::get_dims ( )

Definition at line 239 of file carma_obj.h.

239 { return dims_data; }
Here is the caller graph for this function:

◆ get_dims() [2/2]

long CarmaObj::get_dims ( int  i)

Definition at line 240 of file carma_obj.h.

240 { return dims_data[i]; }

◆ get_h_data()

T_data* CarmaObj::get_h_data ( )

Definition at line 169 of file carma_obj.h.

169 { return h_data.data(); }

◆ get_nb_elements()

int CarmaObj::get_nb_elements ( )

Definition at line 241 of file carma_obj.h.

241 { return nb_elem; }
Here is the caller graph for this function:

◆ get_nb_streams()

int CarmaObj::get_nb_streams ( ) const

get the number of streams attached to the host object

Definition at line 171 of file carma_obj.h.

171  {
174  return streams->get_nb_streams();
175  }

◆ get_o_data()

T_data* CarmaObj::get_o_data ( )

Definition at line 232 of file carma_obj.h.

232 { return o_data; }

◆ get_o_data_value()

const T_data CarmaObj::get_o_data_value ( ) const

Definition at line 233 of file carma_obj.h.

233  {
234  T_data tmp_float;
236  cudaMemcpy(&tmp_float, o_data, sizeof(T_data), cudaMemcpyDeviceToHost));
237  return tmp_float;
238  }

◆ get_plan()

cufftHandle* CarmaObj::get_plan ( )

FFT plan.

Definition at line 278 of file carma_obj.h.

◆ get_type_plan()

cufftType CarmaObj::get_type_plan ( )

FFT plan type.

Definition at line 280 of file carma_obj.h.

◆ get_values()

unsigned int* CarmaObj::get_values ( )

optional data (used for sort)

sum

Definition at line 283 of file carma_obj.h.

◆ host2device()

template<typename T_dest >
int CarmaObj::host2device ( const T_dest *  data)

◆ host2device_async()

int CarmaObj::host2device_async ( const T_data *  data,
cudaStream_t  stream 
)

◆ host2device_mat()

int CarmaObj::host2device_mat ( const T_data *  data,
int  lda,
int  ldb 
)

◆ host2device_vect()

int CarmaObj::host2device_vect ( const T_data *  data,
int  incx,
int  incy 
)

◆ init()

void CarmaObj::init ( CarmaContext current_context,
const long *  dims_data,
const T_data *  data,
bool  fromHost,
int  nb_streams 
)
protected

◆ init_prng() [1/2]

int CarmaObj::init_prng ( )

◆ init_prng() [2/2]

int CarmaObj::init_prng ( long  seed)

◆ init_prng_host()

int CarmaObj::init_prng_host ( int  seed)

◆ init_reduceCub()

void CarmaObj::init_reduceCub ( )

◆ is_rng_init()

bool CarmaObj::is_rng_init ( )

Memory transfers both ways

Definition at line 246 of file carma_obj.h.

◆ memset()

int CarmaObj::memset ( T_data  value)

Definition at line 273 of file carma_obj.h.

273  {
274  return fill_array_with_value(
275  this->d_data, value, this->nb_elem,
276  this->current_context->get_device(this->device));
277  }

◆ nrm2()

T_data CarmaObj::nrm2 ( int  incx)

◆ operator std::string()

CarmaObj::operator std::string ( )

Definition at line 222 of file carma_obj.h.

222 { return this->to_string(); }

◆ operator T_data *()

CarmaObj::operator T_data * ( )

Definition at line 214 of file carma_obj.h.

214 { return d_data; }

◆ operator[]()

const T_data CarmaObj::operator[] ( int  index) const

Definition at line 224 of file carma_obj.h.

224  {
225  T_data tmp_float;
226  carma_safe_call(cudaMemcpy(&tmp_float, &d_data[index], sizeof(T_data),
227  cudaMemcpyDeviceToHost));
228  return tmp_float;
229  }

◆ prng() [1/5]

int CarmaObj::prng ( char  gtype)

◆ prng() [2/5]

int CarmaObj::prng ( char  gtype,
float  alpha 
)

◆ prng() [3/5]

int CarmaObj::prng ( char  gtype,
float  alpha,
float  beta 
)

◆ prng() [4/5]

int CarmaObj::prng ( T_data *  output,
char  gtype,
float  alpha 
)

◆ prng() [5/5]

int CarmaObj::prng ( T_data *  output,
char  gtype,
float  alpha,
float  beta 
)

◆ prng_host() [1/3]

int CarmaObj::prng_host ( char  gtype)

◆ prng_host() [2/3]

int CarmaObj::prng_host ( char  gtype,
T_data  stddev 
)

◆ prng_host() [3/3]

int CarmaObj::prng_host ( char  gtype,
T_data  stddev,
T_data  alpha 
)

◆ prng_montagn()

int CarmaObj::prng_montagn ( float  init_montagn)

◆ reduceCub()

void CarmaObj::reduceCub ( )

◆ reset()

int CarmaObj::reset ( )

Definition at line 270 of file carma_obj.h.

270  {
271  return cudaMemset(this->d_data, 0, this->nb_elem * sizeof(T_data));
272  }
Here is the caller graph for this function:

◆ rot()

void CarmaObj::rot ( CarmaObj< T_data > *  source,
int  incx,
int  incy,
T_data  sc,
T_data  ss 
)

◆ scale()

void CarmaObj::scale ( T_data  alpha,
int  incx 
)

◆ sum()

T_data CarmaObj::sum ( )

◆ swap()

void CarmaObj::swap ( CarmaObj< T_data > *  source,
int  incx,
int  incy 
)

◆ swap_ptr()

void CarmaObj::swap_ptr ( T_data *  ptr)

Definition at line 203 of file carma_obj.h.

203  {
204  dealloc();
205  d_data = ptr;
206  owner = false;
207  }

◆ symm()

void CarmaObj::symm ( char  side,
char  uplo,
T_data  alpha,
CarmaObj< T_data > *  matA,
int  lda,
CarmaObj< T_data > *  matB,
int  ldb,
T_data  beta,
int  ldc 
)

◆ symv()

void CarmaObj::symv ( char  uplo,
T_data  alpha,
CarmaObj< T_data > *  matA,
int  lda,
CarmaObj< T_data > *  vectx,
int  incx,
T_data  beta,
int  incy 
)

◆ sync_h_data()

void CarmaObj::sync_h_data ( )

Definition at line 164 of file carma_obj.h.

164  {
165  if (h_data.empty()) h_data = std::vector<T_data>(nb_elem);
166  device2host(h_data.data());
167  }

◆ syrk()

void CarmaObj::syrk ( char  uplo,
char  transa,
T_data  alpha,
CarmaObj< T_data > *  matA,
int  lda,
T_data  beta,
int  ldc 
)

◆ syrkx()

void CarmaObj::syrkx ( char  uplo,
char  transa,
T_data  alpha,
CarmaObj< T_data > *  matA,
int  lda,
CarmaObj< T_data > *  matB,
int  ldb,
T_data  beta,
int  ldc 
)

◆ to_string()

std::string CarmaObj::to_string ( )

Definition at line 216 of file carma_obj.h.

216  {
217  std::ostringstream stream;
218  stream << *this;
219  return stream.str();
220  }
Here is the caller graph for this function:

◆ transpose()

int CarmaObj::transpose ( CarmaObj< T_data > *  source)

◆ wait_all_streams()

int CarmaObj::wait_all_streams ( )

Definition at line 199 of file carma_obj.h.

199  {
200  this->streams->wait_all_streams();
201  return EXIT_SUCCESS;
202  }

◆ wait_stream()

int CarmaObj::wait_stream ( int  stream)

Definition at line 195 of file carma_obj.h.

195  {
196  this->streams->wait_stream(stream);
197  return EXIT_SUCCESS;
198  }

Member Data Documentation

◆ cub_data

T_data* CarmaObj::cub_data
protected

optional data (used for scan / reduction)

Definition at line 122 of file carma_obj.h.

◆ cub_data_size

size_t CarmaObj::cub_data_size
protected

Definition at line 123 of file carma_obj.h.

◆ current_context

CarmaContext* CarmaObj::current_context
protected

Definition at line 128 of file carma_obj.h.

◆ d_data

T_data* CarmaObj::d_data
protected

Input data => change to vector.

Definition at line 119 of file carma_obj.h.

◆ d_num_valid

size_t* CarmaObj::d_num_valid
protected

used for compact

Definition at line 140 of file carma_obj.h.

◆ d_states

curandState* CarmaObj::d_states
protected

Definition at line 131 of file carma_obj.h.

◆ device

int CarmaObj::device
protected

device where the CarmaObj is allocate

Definition at line 127 of file carma_obj.h.

◆ dims_data

long* CarmaObj::dims_data
protected

dimensions of the array

Definition at line 125 of file carma_obj.h.

◆ gen

curandGenerator_t CarmaObj::gen
protected

Definition at line 130 of file carma_obj.h.

◆ h_data

std::vector<T_data> CarmaObj::h_data
protected

Definition at line 120 of file carma_obj.h.

◆ keys_only

bool CarmaObj::keys_only
protected

Definition at line 136 of file carma_obj.h.

◆ nb_blocks

int CarmaObj::nb_blocks
protected

Definition at line 134 of file carma_obj.h.

◆ nb_elem

int CarmaObj::nb_elem
protected

number of elements in the array

Definition at line 126 of file carma_obj.h.

◆ nb_threads

int CarmaObj::nb_threads
protected

Definition at line 133 of file carma_obj.h.

◆ ndim

int CarmaObj::ndim
protected

Definition at line 124 of file carma_obj.h.

◆ o_data

T_data* CarmaObj::o_data
protected

optional data (used for scan / reduction)

Definition at line 121 of file carma_obj.h.

◆ owner

bool CarmaObj::owner = true
protected

Definition at line 137 of file carma_obj.h.

◆ plan

cufftHandle CarmaObj::plan
protected

FFT plan.

Definition at line 142 of file carma_obj.h.

◆ streams

CarmaStreams* CarmaObj::streams
protected

Definition at line 145 of file carma_obj.h.

◆ type_plan

cufftType CarmaObj::type_plan
protected

FFT plan type.

Definition at line 143 of file carma_obj.h.

◆ values

unsigned int* CarmaObj::values
protected

optional data (used for sort)

Definition at line 139 of file carma_obj.h.


The documentation for this class was generated from the following files:
CarmaObj::device2host
int device2host(T_dest *data)
carma_safe_call
#define carma_safe_call(err)
Definition: carma_utils.h:145
CarmaObj::dealloc
void dealloc()
Definition: carma_obj.h:209
CarmaObj::nb_elem
int nb_elem
number of elements in the array
Definition: carma_obj.h:126
debug_pyr.index
int index
Definition: debug_pyr.py:78
CarmaStreams::del_stream
int del_stream()
CarmaStreams::add_stream
int add_stream()
CarmaObj::d_data
T_data * d_data
Input data => change to vector.
Definition: carma_obj.h:119
CarmaStreams::wait_all_streams
int wait_all_streams()
CarmaObj::h_data
std::vector< T_data > h_data
Definition: carma_obj.h:120
CarmaStreams::get_nb_streams
int get_nb_streams()
CarmaObj::o_data
T_data * o_data
optional data (used for scan / reduction)
Definition: carma_obj.h:121
CarmaObj::streams
CarmaStreams * streams
Definition: carma_obj.h:145
CarmaObj::to_string
std::string to_string()
Definition: carma_obj.h:216
CarmaObj::device
int device
device where the CarmaObj is allocate
Definition: carma_obj.h:127
CarmaObj::dims_data
long * dims_data
dimensions of the array
Definition: carma_obj.h:125
CarmaStreams::get_stream
cudaStream_t get_stream(int stream)
CarmaObj::owner
bool owner
Definition: carma_obj.h:137
CarmaObj::current_context
CarmaContext * current_context
Definition: carma_obj.h:128
CarmaContext::get_device
CarmaDevice * get_device(int dev)
Definition: carma_context.h:132
layers_test.i
int i
Definition: layers_test.py:43
CarmaStreams::wait_stream
int wait_stream(int stream)
fill_array_with_value
int fill_array_with_value(T_data *d_data, T_data value, int N, CarmaDevice *device)