<commit_before>// -*- c++ -*-
// adsyn.cu
// experimental cuda opcodes
//
// <NAME>, 2013

#include <csdl.h>
#include <pstream.h>

typedef struct cudadsyn_ {
  OPDS h;
  MYFLT *asig;
  PVSDAT *fsig;
  MYFLT *kamp, *kfreq;
  MYFLT *inum;
  float *out;
  float *frame;
  int64_t *ndx;
  float *fp, *previous;
  AUXCH out_;
  int bins, blocks, threads;
  int count;
  int vsamps, mblocks, mthreads;
  int framecount;
} CUDADSYN;

static int destroy_cudadsyn(CSOUND *csound, void *pp);

static int init_cudadsyn(CSOUND *csound, CUDADSYN *p){

  int asize, ipsize, fpsize, blockspt;
  if(p->fsig->overlap > 1024)
     return csound->InitError(csound, "overlap is too large\n");
  cudaDeviceProp deviceProp;
  cudaGetDeviceProperties(&deviceProp, 0);
  blockspt = deviceProp.maxThreadsPerBlock;
  if(deviceProp.major < 1)
   csound->InitError(csound,
		     "this opcode requires device capability 1.0 minimum. Device is %d.%d\n", 
        deviceProp.major, deviceProp.minor );

  p->bins = (p->fsig->N)/2;

  if(*p->inum > 0 && *p->inum < p->bins) p->bins = *p->inum;

  p->vsamps = p->fsig->overlap;
  p->threads = p->bins*p->vsamps;
  p->blocks = p->threads > blockspt ? p->threads/blockspt : 1;
  p->mthreads = p->bins;
  p->mblocks = p->mthreads >  blockspt ? p->mthreads/blockspt : 1;

  p->threads /= p->blocks;
  p->mthreads /= p->mblocks;

  asize =  p->vsamps*p->bins*sizeof(float);
  ipsize = p->fsig->N*sizeof(int64_t)/2;
  fpsize = p->fsig->N*sizeof(float);

  cudaMalloc(&p->out, asize);
  cudaMalloc(&p->ndx, ipsize);
  cudaMalloc(&p->frame, fpsize);
  cudaMalloc(&p->previous, fpsize/2);
  cudaMemset(p->previous, 0, fpsize/2);
  cudaMemset(p->ndx, 0, ipsize);

  asize = p->vsamps*sizeof(float);
  if(p->out_.auxp == NULL ||
     p->out_.size < asize)
    csound->AuxAlloc(csound, asize , &p->out_);

  csound->RegisterDeinitCallback(csound, p, destroy_cudadsyn);
  p->count = 0;
  return OK;
}

__global__ void sample(float *out, float *frame, float pitch, int64_t *ph,
                       float *amps, int bins, int vsize, float sr) {

  int t = (threadIdx.x + blockIdx.x*blockDim.x);
  int n =  t%vsize;  /* sample index */
  int h = t/vsize;  /* bin index */
  int k = h<<1;
  int64_t lph;
  float a = amps[h], ascl = ((float)n)/vsize;
  float fscal = pitch*FMAXLEN/sr;
  lph = (ph[h] + (int64_t)(n*round(frame[k+1]*fscal))) & PHMASK;
  a += ascl*(frame[k] - a);
  out[t] = a*sinf((2*PI*lph)/FMAXLEN);
  if(t >= vsize) return;
  syncthreads();
  for(int i=vsize; i < vsize*bins; i+=vsize)
    out[t] += out[t + i];
}

__global__ void update(float *frame, float *amps,
      int64_t *ph,float pitch, int vsize, float sr){

 int h = threadIdx.x + blockIdx.x*blockDim.x;
 int k = h << 1;
 /* update phases and amps */
 ph[h]  = (ph[h] + (int64_t)(vsize*round(pitch*frame[k+1]*FMAXLEN/sr))) & PHMASK;
 amps[h] = frame[k];
}
<commit_msg>added caps 1.1 version of adsyn  opcode<commit_after>// -*- c++ -*-
// adsyn.cu
// experimental cuda opcodes
//
// <NAME>, 2013

#include <csdl.h>
#include <pstream.h>

typedef struct cudadsyn_ {
  OPDS h;
  MYFLT *asig;
  PVSDAT *fsig;
  MYFLT *kamp, *kfreq;
  MYFLT *inum;
  float *out;
  float *frame;
  int64_t *ndx;
  float *fp, *previous;
  AUXCH out_;
  int bins, blocks, threads;
  int count;
  int vsamps, mblocks, mthreads;
  int framecount;
} CUDADSYN;

static int destroy_cudadsyn(CSOUND *csound, void *pp);

static int init_cudadsyn(CSOUND *csound, CUDADSYN *p){

  int asize, ipsize, fpsize, blockspt;
  if(p->fsig->overlap > 1024)
     return csound->InitError(csound, "overlap is too large\n");
  cudaDeviceProp deviceProp;
  cudaGetDeviceProperties(&deviceProp, 0);
  blockspt = deviceProp.maxThreadsPerBlock;
  if(deviceProp.major < 1)
   csound->InitError(csound,
		     "this opcode requires device capability 1.0 minimum. Device is %d.%d\n", 
        deviceProp.major, deviceProp.minor );

  p->bins = (p->fsig->N)/2;

  if(*p->inum > 0 && *p->inum < p->bins) p->bins = *p->inum;

  p->vsamps = p->fsig->overlap;
  p->threads = p->bins*p->vsamps;
  p->blocks = p->threads > blockspt ? p->threads/blockspt : 1;
  p->mthreads = p->bins;
  p->mblocks = p->mthreads >  blockspt ? p->mthreads/blockspt : 1;

  p->threads /= p->blocks;
  p->mthreads /= p->mblocks;

  asize =  p->vsamps*p->bins*sizeof(float);
  ipsize = p->fsig->N*sizeof(int64_t)/2;
  fpsize = p->fsig->N*sizeof(float);

  cudaMalloc(&p->out, asize);
  cudaMalloc(&p->ndx, ipsize);
  cudaMalloc(&p->frame, fpsize);
  cudaMalloc(&p->previous, fpsize/2);
  cudaMemset(p->previous, 0, fpsize/2);
  cudaMemset(p->ndx, 0, ipsize);

  asize = p->vsamps*sizeof(float);
  if(p->out_.auxp == NULL ||
     p->out_.size < asize)
    csound->AuxAlloc(csound, asize , &p->out_);

  csound->RegisterDeinitCallback(csound, p, destroy_cudadsyn);
  p->count = 0;
  return OK;
}

__global__ void sample(float *out, float *frame, float pitch, int64_t *ph,
                       float *amps, int bins, int vsize, float sr) {

  int t = (threadIdx.x + blockIdx.x*blockDim.x);
  int n =  t%vsize;  /* sample index */
  int h = t/vsize;  /* bin index */
  int k = h<<1;
  int64_t lph; 
  float a = amps[h], ascl = ((float)n)/vsize;
  float fscal = pitch*FMAXLEN/sr;
  lph = (ph[h] + (int64_t)(n*round(frame[k+1]*fscal))) & PHMASK;
  a += ascl*(frame[k] - a);
  out[t] = a*sinf((2*PI*lph)/FMAXLEN);
  if(t >= vsize) return;
  syncthreads();
  for(int i=vsize; i < vsize*bins; i+=vsize)
    out[t] += out[t + i];
}

__global__ void update(float *frame, float *amps,
      int64_t *ph,float pitch, int vsize, float sr){

 int h = threadIdx.x + blockIdx.x*blockDim.x;
 int k = h << 1;
 /* update phases and amps */
 ph[h]  = (ph[h] + (int64_t)(vsize*round(pitch*frame[k+1]*FMAXLEN/sr))) & PHMASK;
 amps[h] = frame[k];
}

{'fraction_non_alphanumeric': 0.18072289156626506, 'fraction_numerical': 0.014667365112624411, 'mean_word_length': 3.43343653250774, 'pattern_counts': {'":': 0, '<': 25, '<?xml version=': 0, '>': 117, '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': '6503092', 'n_tokens_mistral': 2504, 'n_tokens_neox': 2381, 'n_words': 520}