Add files via upload

2019-11-25 01:49:17 -05:00
parent 1ff385daaf
commit a38987e7af
9 changed files with 1694 additions and 0 deletions
--- a/benchmarks/opencl/BlackScholes/oclBlackScholes_launcher.cpp
+++ b/benchmarks/opencl/BlackScholes/oclBlackScholes_launcher.cpp
@@ -0,0 +1,125 @@
+/*
+ * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
+ *
+ * Please refer to the NVIDIA end user license agreement (EULA) associated
+ * with this source code for terms and conditions that govern your use of
+ * this software. Any use, reproduction, disclosure, or distribution of
+ * this software and related documentation outside the terms of the EULA
+ * is strictly prohibited.
+ *
+ */
+
+
+
+#include <oclUtils.h>
+#include "oclBlackScholes_common.h"
+
+static cl_program cpBlackScholes;   //OpenCL program
+static cl_kernel  ckBlackScholes;   //OpenCL kernel
+static cl_command_queue cqDefaultCommandQueue;
+
+extern "C" void initBlackScholes(cl_context cxGPUContext, cl_command_queue cqParamCommandQueue, const char **argv){
+    cl_int ciErrNum;
+    size_t kernelLength;
+
+    shrLog("...loading BlackScholes.cl\n");
+        char *cPathAndName = shrFindFilePath("BlackScholes.cl", argv[0]);
+        shrCheckError(cPathAndName != NULL, shrTRUE);
+        char *cBlackScholes = oclLoadProgSource(cPathAndName, "// My comment\n", &kernelLength);
+        shrCheckError(cBlackScholes != NULL, shrTRUE);
+
+    shrLog("...creating BlackScholes program\n");
+        //cpBlackScholes = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&cBlackScholes, &kernelLength, &ciErrNum);
+        cpBlackScholes = clCreateProgramWithBuiltInKernels(context, 1, &device_id, "BlackScholes", NULL);
+        shrCheckError(ciErrNum, CL_SUCCESS);
+
+    shrLog("...building BlackScholes program\n");
+        ciErrNum = clBuildProgram(cpBlackScholes, 0, NULL, "-cl-fast-relaxed-math -Werror", NULL, NULL);
+
+        if(ciErrNum != CL_BUILD_SUCCESS){
+            shrLog("*** Compilation failure ***\n");
+
+            size_t deviceNum;
+            cl_device_id *cdDevices;
+            ciErrNum = clGetContextInfo(cxGPUContext, CL_CONTEXT_DEVICES, 0, NULL, &deviceNum);
+            shrCheckError(ciErrNum, CL_SUCCESS);
+
+            cdDevices = (cl_device_id *)malloc(deviceNum * sizeof(cl_device_id));
+            shrCheckError(cdDevices != NULL, shrTRUE);
+
+            ciErrNum = clGetContextInfo(cxGPUContext, CL_CONTEXT_DEVICES, deviceNum * sizeof(cl_device_id), cdDevices, NULL);
+            shrCheckError(ciErrNum, CL_SUCCESS);
+
+            size_t logSize;
+            char *logTxt;
+
+            ciErrNum = clGetProgramBuildInfo(cpBlackScholes, cdDevices[0], CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
+            shrCheckError(ciErrNum, CL_SUCCESS);
+
+            logTxt = (char *)malloc(logSize);
+            shrCheckError(logTxt != NULL, shrTRUE);
+
+            ciErrNum = clGetProgramBuildInfo(cpBlackScholes, cdDevices[0], CL_PROGRAM_BUILD_LOG, logSize, logTxt, NULL);
+            shrCheckError(ciErrNum, CL_SUCCESS);
+
+            shrLog("%s\n", logTxt);
+            shrLog("*** Exiting ***\n");
+            free(logTxt);
+            free(cdDevices);
+            exit(666);
+        }
+
+    //Save ptx code to separate file
+    oclLogPtx(cpBlackScholes, oclGetFirstDev(cxGPUContext), "BlackScholes.ptx");
+
+    shrLog("...creating BlackScholes kernels\n");
+        ckBlackScholes = clCreateKernel(cpBlackScholes, "BlackScholes", &ciErrNum);
+        shrCheckError(ciErrNum, CL_SUCCESS);
+
+    cqDefaultCommandQueue = cqParamCommandQueue;
+    free(cBlackScholes);
+    free(cPathAndName);
+}
+
+extern "C" void closeBlackScholes(void){
+    cl_int ciErrNum;
+    ciErrNum  = clReleaseKernel(ckBlackScholes);
+    ciErrNum |= clReleaseProgram(cpBlackScholes);
+    shrCheckError(ciErrNum, CL_SUCCESS);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// OpenCL Black-Scholes kernel launcher
+////////////////////////////////////////////////////////////////////////////////
+extern "C" void BlackScholes(
+    cl_command_queue cqCommandQueue,
+    cl_mem d_Call, //Call option price
+    cl_mem d_Put,  //Put option price
+    cl_mem d_S,    //Current stock price
+    cl_mem d_X,    //Option strike price
+    cl_mem d_T,    //Option years
+    cl_float R,    //Riskless rate of return
+    cl_float V,    //Stock volatility
+    cl_uint optionCount
+){
+    cl_int ciErrNum;
+
+    if(!cqCommandQueue)
+        cqCommandQueue = cqDefaultCommandQueue;
+
+    ciErrNum  = clSetKernelArg(ckBlackScholes, 0, sizeof(cl_mem),   (void *)&d_Call);
+    ciErrNum |= clSetKernelArg(ckBlackScholes, 1, sizeof(cl_mem),   (void *)&d_Put);
+    ciErrNum |= clSetKernelArg(ckBlackScholes, 2, sizeof(cl_mem),   (void *)&d_S);
+    ciErrNum |= clSetKernelArg(ckBlackScholes, 3, sizeof(cl_mem),   (void *)&d_X);
+    ciErrNum |= clSetKernelArg(ckBlackScholes, 4, sizeof(cl_mem),   (void *)&d_T);
+    ciErrNum |= clSetKernelArg(ckBlackScholes, 5, sizeof(cl_float), (void *)&R);
+    ciErrNum |= clSetKernelArg(ckBlackScholes, 6, sizeof(cl_float), (void *)&V);
+    ciErrNum |= clSetKernelArg(ckBlackScholes, 7, sizeof(cl_uint),  (void *)&optionCount);
+    shrCheckError(ciErrNum, CL_SUCCESS);
+
+    //Run the kernel
+    size_t globalWorkSize = 60 * 1024;
+	size_t localWorkSize = 128;
+    ciErrNum = clEnqueueNDRangeKernel(cqCommandQueue, ckBlackScholes, 1, NULL, &globalWorkSize, &localWorkSize, 0, NULL, NULL);
+    shrCheckError(ciErrNum, CL_SUCCESS);
+}