samples/atiSamples/BinomialOption/BinomialOption.rkt
#lang racket
(require "../../../c.rkt"
         "../atiUtils/utils.rkt"
         ffi/unsafe
         ffi/cvector
         ffi/unsafe/cvector)

(define RISKFREE 0.02)
(define VOLATILITY 0.30)
(define setupTime -1)
(define totalKernelTime -1)
(define devices #f)
(define context #f)
(define commandQueue #f)
(define program #f)
(define kernel #f)
(define randArray #f)
(define output #f)
(define randBuffer #f)
(define outBuffer #f)
(define numSamples 1024)
(define numSteps 0)

(define (setupBinomialOption)
  (set! randArray (malloc (* numSamples (ctype-sizeof _cl_float4)) 'raw))
  (for ([i (in-range (* 4 numSamples))])
    (ptr-set! randArray _cl_float i (random)))
  (set! output (malloc (* numSamples (ctype-sizeof _cl_float4)) 'raw))
  (memset output 0 (* numSamples (ctype-sizeof _cl_float4))))

(define (setupCL)
  (set!-values (devices context commandQueue program) (init-cl "BinomialOption_Kernels.cl" #:queueProperties 'CL_QUEUE_PROFILING_ENABLE))
  (set! randBuffer (clCreateBuffer context '(CL_MEM_READ_ONLY CL_MEM_USE_HOST_PTR) (* (ctype-sizeof _cl_float4) numSamples) randArray))
  (set! outBuffer (clCreateBuffer context '(CL_MEM_WRITE_ONLY CL_MEM_USE_HOST_PTR) (* (ctype-sizeof _cl_float4) numSamples) output))
  (set! kernel (clCreateKernel program #"binomial_options"))
  (set! numSteps (- (optimum-threads kernel (cvector-ref devices 0) 256) 2)))

(define (runCLKernels)
  (clSetKernelArg:_cl_int kernel 0 numSteps)
  (clSetKernelArg:_cl_mem kernel 1 randBuffer)
  (clSetKernelArg:_cl_mem kernel 2 outBuffer)
  (clSetKernelArg:local kernel 3 (* (add1 numSteps) (ctype-sizeof _cl_float4)))
  (clSetKernelArg:local kernel 4 (* numSteps (ctype-sizeof _cl_float4)))
  (define globalThreads (* numSamples (add1 numSteps)))
  (define localThreads (add1 numSteps))
  (clEnqueueNDRangeKernel commandQueue kernel 1 (vector globalThreads) (vector localThreads) (make-vector 0))
  (clFinish commandQueue)
  (clEnqueueReadBuffer commandQueue outBuffer 'CL_TRUE 0 (* numSamples (ctype-sizeof _cl_float4)) output (make-vector 0)))

(define (binomialOptionCPUReference)
  (define refOutput (malloc (* numSamples (ctype-sizeof _cl_float4))))
  (define stepsArray (make-vector (* (add1 numSteps) 4)))
  (for ([bid (in-range numSamples)])
    (define s (make-vector 4))
    (define x (make-vector 4))
    (define vsdt (make-vector 4))
    (define puByr (make-vector 4))
    (define pdByr (make-vector 4))
    (define optionYears (make-vector 4))
    (define inRand (make-vector 4))
    (for ([i (in-range 4)])
      (vector-set! inRand i (ptr-ref randArray _cl_float (+ bid i)))
      (define val (vector-ref inRand i))
      (vector-set! s i (+ (* (- 1.0 val) 5.0) (* val 30.0)))
      (vector-set! x i (+ (* (- 1.0 val) 1.0) (* val 100.0)))
      (vector-set! optionYears i (+ (* (- 1.0 val) 0.25) (* val 10.0)))
      (define dt (* (vector-ref optionYears i) (/ 1.0 numSteps)))
      (vector-set! vsdt i (* VOLATILITY (sqrt dt)))
      (define rdt (* RISKFREE dt))
      (define r (exp rdt))
      (define rInv (/ 1.0 r))
      (define u (exp (vector-ref vsdt i)))
      (define d (/ 1.0 u))
      (define pu (/ (- r d) (- u d)))
      (define pd (- 1.0 pu))
      (vector-set! puByr i (* pu rInv))
      (vector-set! pdByr i (* pd rInv)))
    (for ([j (in-range (add1 numSteps))])
      (for ([i (in-range 4)])
        (define profit (- (* (vector-ref s i) (exp (* (vector-ref vsdt i) (- (* 2.0 j) numSteps)))) (vector-ref x i)))
        (vector-set! stepsArray (+ i (* j 4)) (if (> profit 0.0) profit 0.0))))
    (for ([j (in-range numSteps 0 -1)])
      (for ([k (in-range j)])
        (for ([i (in-range 4)])
          (vector-set! stepsArray (+ i (* k 4)) (+ (* (vector-ref pdByr i) (vector-ref stepsArray (+ i (* 4 (add1 k)))))
                                                    (* (vector-ref puByr i) (vector-ref stepsArray (+ i (* k 4)))))))))
    (ptr-set! refOutput _cl_float bid (vector-ref stepsArray 0)))
  (compare output refOutput numSamples))
  
      
(define (setup)
  (setupBinomialOption)
  (set! setupTime (time-real setupCL)))

(define (run)
  (set! totalKernelTime (time-real runCLKernels)))

(define (verify-results)
  (define verified (binomialOptionCPUReference))
  (printf "~n~a~n" (if verified "Passed" "Failed")))

(define (cleanup)
  (clReleaseKernel kernel)
  (clReleaseProgram program)
  (clReleaseMemObject randBuffer)
  (clReleaseMemObject outBuffer)
  (clReleaseCommandQueue commandQueue)
  (clReleaseContext context)
  (free randArray)
  (free output))

(define (print-stats)
  (printf "~nOption Samples: ~a, Setup Time: ~a, Kernel Time: ~a, Total Time: ~a, Options/sec: ~a~n"
          numSamples 
          (real->decimal-string setupTime 3) 
          (real->decimal-string totalKernelTime 3)
          (real->decimal-string (+ setupTime totalKernelTime) 3)
          (real->decimal-string (/ numSamples (+ setupTime totalKernelTime)))))

(setup)
(run)
(verify-results)
(cleanup)
(print-stats)