From b59cca70937b15c921d350018f947ce23c914b19 Mon Sep 17 00:00:00 2001 From: ksirg Date: Sun, 22 Sep 2013 23:18:25 +0200 Subject: [PATCH] first working version of CuRBFEllILPEvaluator, which does prediction. So far it doesn't use async memory transfers and streams. --- KMLib.CUDA/CUDAmodules/Config.h | 4 +- KMLib.CUDA/CUDAmodules/KernelsEllpack.cu | 1 + KMLib.CUDA/CudaRBFCSREvaluator.cs | 16 ++- KMLib.CUDA/GPUEvaluators/CuEvaluator.cs | 122 ++++++++++++------ .../GPUEvaluators/CuRBFEllILPEvaluator.cs | 74 +++++------ KMLib.CUDA/GPUKernels/CuRBFEllILPKernel.cs | 2 +- KMLib.CUDA/KMLib.GPU.csproj | 2 +- KMLib/Evaluate/DualEvaluator.cs | 16 +-- KMLib/KMLib.csproj.ksirg.nvuser | 6 + KMLibUsageApp/Program.cs | 2 - 10 files changed, 143 insertions(+), 102 deletions(-) create mode 100644 KMLib/KMLib.csproj.ksirg.nvuser diff --git a/KMLib.CUDA/CUDAmodules/Config.h b/KMLib.CUDA/CUDAmodules/Config.h index d1e71fe..bb70c01 100644 --- a/KMLib.CUDA/CUDAmodules/Config.h +++ b/KMLib.CUDA/CUDAmodules/Config.h @@ -35,7 +35,7 @@ texture mainVec2TexRef; #define maxNNZ 100 -//rho for computing +//rho for computingwse //__constant__ float RHO=-2; //sets -1 for negative values and 1 for gather or equal than 0 //params: @@ -79,7 +79,7 @@ extern "C" __global__ void reduce(const float* input, float* output, const int N unsigned int i = blockIdx.x*blockDim.x*2 + threadIdx.x; unsigned int gridSize = blockDim.x*2*gridDim.x; - shVals[tid]=-FLT_MAX; + shVals[tid]=0; float sum=0; while (i < N) diff --git a/KMLib.CUDA/CUDAmodules/KernelsEllpack.cu b/KMLib.CUDA/CUDAmodules/KernelsEllpack.cu index 6e934c4..6f9f1e9 100644 --- a/KMLib.CUDA/CUDAmodules/KernelsEllpack.cu +++ b/KMLib.CUDA/CUDAmodules/KernelsEllpack.cu @@ -1113,6 +1113,7 @@ extern "C" __global__ void rbfEllpackILPEvaluator(const float * vals, //hack for choosing different texture reference when launch in defferent streams float dot = texSel==1 ? SpMV_Ellpack_ILP<1>(vals,colIdx,rowLength,row,num_rows): SpMV_Ellpack_ILP<2>(vals,colIdx,rowLength,row,num_rows) ; results[row]=svY[row]*svAlpha[row]*expf(-shGamma*(svSelfDot[row]+shVecSelfDot-2*dot)); + //results[row]=dot; } } diff --git a/KMLib.CUDA/CudaRBFCSREvaluator.cs b/KMLib.CUDA/CudaRBFCSREvaluator.cs index d8737c5..c871940 100644 --- a/KMLib.CUDA/CudaRBFCSREvaluator.cs +++ b/KMLib.CUDA/CudaRBFCSREvaluator.cs @@ -48,19 +48,23 @@ public override float[] Predict(SparseVec[] elements) throw new ApplicationException("Evaluator is not initialized. Call init method"); - //tranfsorm elements to matrix in CSR format + //transform elements to matrix in CSR format // elements values float[] vecVals; //elements indexes int[] vecIdx; - //elements lenght + //elements length int[] vecLenght; CudaHelpers.TransformToCSRFormat(out vecVals, out vecIdx, out vecLenght, elements); float[] elSelfDot = new float[elements.Length]; + + Stopwatch t = Stopwatch.StartNew(); + + //elSelfDot = elements.Select(e => e.DotProduct()).ToArray(); for (int j = 0; j < elements.Length; j++) { float res = 0; @@ -99,7 +103,7 @@ public override float[] Predict(SparseVec[] elements) outputIntPtr = cuda.HostAllocate(memElementsSize, CUDADriver.CU_MEMHOSTALLOC_DEVICEMAP); outputPtr = cuda.GetHostDevicePointer(outputIntPtr, 0); - // Set the cuda kernel paramerters + // Set the cuda kernel parameters #region set cuda parameters uint Rows = (uint)elements.Length; uint Cols = (uint)TrainedModel.SupportElements.Length; @@ -128,7 +132,7 @@ public override float[] Predict(SparseVec[] elements) offset += IntPtr.Size; cuda.SetParameter(cuFunc, offset, elSelf.Pointer); offset += IntPtr.Size; - //set output (reslut) param + //set output (result) param cuda.SetParameter(cuFunc, offset, outputPtr.Pointer); offset += IntPtr.Size; //set number of elements param @@ -159,13 +163,13 @@ public override float[] Predict(SparseVec[] elements) CudaHelpers.InitBuffer(TrainedModel.SupportElements[k], svVecIntPtrs[k % 2]); cuda.SynchronizeStream(stream); - //copy asynchronously from buffer to devece + //copy asynchronously from buffer to device cuda.CopyHostToDeviceAsync(mainVecPtr, svVecIntPtrs[k % 2], memSvSize, stream); //set the last parameter in kernel (column index) // colIndexParamOffset cuda.SetParameter(cuFunc, lastParameterOffset, (uint)k); - //launch kernl + //launch kernel cuda.LaunchAsync(cuFunc, gridDimX, 1, stream); if (k > 0) diff --git a/KMLib.CUDA/GPUEvaluators/CuEvaluator.cs b/KMLib.CUDA/GPUEvaluators/CuEvaluator.cs index 62494d1..e7d352b 100644 --- a/KMLib.CUDA/GPUEvaluators/CuEvaluator.cs +++ b/KMLib.CUDA/GPUEvaluators/CuEvaluator.cs @@ -22,7 +22,7 @@ namespace KMLib.GPU /// /// base class for all cuda enabled evaluators /// - /// It sores nesessary data for cuda initialization + /// It sores necessary data for cuda initialization public abstract class CuEvaluator : Evaluator { @@ -63,43 +63,45 @@ public abstract class CuEvaluator : Evaluator /// protected int reductionBlocks; /// - /// numer of theread used for reduction + /// number of threads used for reduction /// protected int reductionThreads; /// /// threads per block, def=128, /// - protected int maxReductionThreads = 256; + protected int maxReductionThreads = 128; /// - /// threads per blokc for eval function + /// threads per block for evaluation function /// protected int evalThreads = CUDAConfig.XBlockSize; /// - /// blocks per grid for eval function + /// blocks per grid for evaluation function /// protected int evalBlocks=-1; /// - /// array of 2 buffers for concurent data transfer + /// array of 2 buffers for concurrent data transfer /// - protected IntPtr[] vecIntPtrs = new IntPtr[NUM_STREAMS]; + protected IntPtr[] mainVecIntPtrs = new IntPtr[NUM_STREAMS]; /// /// dense support vector float buffer size /// - protected uint memSize=0; + protected uint vectorsDimMemSize=0; - + protected int vectorSelfDotParamOffset; + protected int texSelParamOffset; + protected int kernelResultParamOffset; #region cuda types /// - /// Cuda .net class for cuda opeation + /// Cuda .net class for cuda operation /// protected CUDA cuda; @@ -112,7 +114,7 @@ public abstract class CuEvaluator : Evaluator /// /// cuda kernel function for computing evaluation values /// - protected CUfunction[] cuFuncEval= new CUfunction[NUM_STREAMS]; + protected CUfunction cuFuncEval; /// /// cuda kernel function for computing evaluation values @@ -130,19 +132,19 @@ public abstract class CuEvaluator : Evaluator /// /// cuda mainVector pointer to device memory (it stores dense vector for prediction) /// - protected CUdeviceptr[] mainVecPtr= new CUdeviceptr[NUM_STREAMS]; + protected CUdeviceptr[] mainVecCuPtr= new CUdeviceptr[NUM_STREAMS]; - protected CUdeviceptr[] outputCuPtr = new CUdeviceptr[NUM_STREAMS]; + protected CUdeviceptr[] evalOutputCuPtr = new CUdeviceptr[NUM_STREAMS]; /// - /// cuda pointer to labels, neded for coping to texture + /// cuda pointer to labels, needed for coping to texture /// protected CUdeviceptr labelsPtr; /// - /// cuda pointer to support vector non zero alpha coeficients + /// cuda pointer to support vector non zero alpha coefficients /// protected CUdeviceptr alphasPtr; @@ -151,14 +153,14 @@ public abstract class CuEvaluator : Evaluator /// protected CUstream[] stream = new CUstream[NUM_STREAMS]; - private CUdeviceptr[] cuReducePtr = new CUdeviceptr[NUM_STREAMS]; + private CUdeviceptr[] reduceCuPtr = new CUdeviceptr[NUM_STREAMS]; - private IntPtr[] redIntPtrs=new IntPtr[NUM_STREAMS]; + private IntPtr[] reduceIntPtrs=new IntPtr[NUM_STREAMS]; /// - /// Offset in cuda setparameter function for pointer to memory to reduce + /// Offset in cuda 'setparameter' function for pointer to memory to reduce /// private int offsetMemToReduce; @@ -184,7 +186,7 @@ public override float[] Predict(SparseVec[] elements) uint reduceSize = (uint)reductionBlocks * sizeof(float); - int loop = (elements.Length +NUM_STREAMS)/ NUM_STREAMS; + int loop = (elements.Length +NUM_STREAMS-1)/ NUM_STREAMS; for (int i = 0; i < loop; i++) { @@ -195,22 +197,46 @@ public override float[] Predict(SparseVec[] elements) { var vec = elements[idx]; + //remove + //float[] svDots = TrainedModel.SupportElements.Select(sv => sv.DotProduct(vec)).ToArray(); + + //set nonzero values to dense vector accessible through vecIntPtr + CudaHelpers.InitBuffer(vec, mainVecIntPtrs[s]); - //set nonzero values to dense vector accesible throught vecIntPtr - CudaHelpers.InitBuffer(vec, vecIntPtrs[s]); + //cuda.CopyHostToDeviceAsync(mainVecPtr[s], vecIntPtrs[s], memSize, stream[s]); + cuda.CopyHostToDevice(mainVecCuPtr[s], mainVecIntPtrs[s], vectorsDimMemSize); + + //remove + //float[] v = new float[vec.Dim + 1]; + //cuda.CopyDeviceToHost(mainVecCuPtr[s],v); - cuda.CopyHostToDeviceAsync(mainVecPtr[s], vecIntPtrs[s], memSize, stream[s]); //cuFunc user different textures - cuda.LaunchAsync(cuFuncEval[s], evalBlocks, 1, stream[s]); + cuda.SetParameter(cuFuncEval, kernelResultParamOffset, evalOutputCuPtr[s]); + cuda.SetParameter(cuFuncEval, vectorSelfDotParamOffset, vec.DotProduct()); + cuda.SetParameter(cuFuncEval, texSelParamOffset, s+1); + + //cuda.LaunchAsync(cuFuncEval[s], evalBlocks, 1, stream[s]); + cuda.Launch(cuFuncEval, evalBlocks, 1); + //remove + //float[] t = new float[sizeSV]; + //cuda.CopyDeviceToHost(evalOutputCuPtr[s], t); - cuda.SetParameter(cuFuncReduce, offsetMemToReduce, outputCuPtr[s]); - cuda.LaunchAsync(cuFuncReduce, reductionBlocks, 1, stream[s]); - cuda.CopyDeviceToHostAsync(cuReducePtr[s], redIntPtrs[s], reduceSize, stream[s]); + cuda.SetParameter(cuFuncReduce, offsetMemToReduce, evalOutputCuPtr[s]); + cuda.SetParameter(cuFuncReduce, offsetOutMemReduce, reduceCuPtr[s]); + //cuda.LaunchAsync(cuFuncReduce, reductionBlocks, 1, stream[s]); + cuda.Launch(cuFuncReduce, reductionBlocks, 1); + + //cuda.CopyDeviceToHostAsync(cuReducePtr[s], redIntPtrs[s], reduceSize, stream[s]); + cuda.CopyDeviceToHost(reduceCuPtr[s], reduceIntPtrs[s], reduceSize); + + //remove + //float[] r = new float[maxReductionBlocks]; + //cuda.CopyDeviceToHost(reduceCuPtr[s], r); } } @@ -225,11 +251,11 @@ public override float[] Predict(SparseVec[] elements) { var vec = elements[idx]; //clear the buffer - //set nonzero values to dense vector accesible throught vecIntPtr - CudaHelpers.SetBufferIdx(vec, vecIntPtrs[s], 0.0f); - float evalValue = ReduceOnHost(redIntPtrs[s], reduceSize); + //set nonzero values to dense vector accessible thought vecIntPtr + CudaHelpers.SetBufferIdx(vec, mainVecIntPtrs[s], 0.0f); + float evalValue = ReduceOnHost(reduceIntPtrs[s], reductionBlocks); - prediction[i] = evalValue; + prediction[idx] = evalValue; } } @@ -244,7 +270,7 @@ public override float[] Predict(SparseVec[] elements) - private float ReduceOnHost(IntPtr reduceIntPtr, uint reduceSize) + private float ReduceOnHost(IntPtr reduceIntPtr, int reduceSize) { double sum = 0; @@ -252,7 +278,7 @@ private float ReduceOnHost(IntPtr reduceIntPtr, uint reduceSize) { float* vecPtr = (float*)reduceIntPtr.ToPointer(); - for (uint j = 0; j < reduceSize; j++) + for (int j = 0; j < reduceSize; j++) { sum+=vecPtr[j]; @@ -260,6 +286,11 @@ private float ReduceOnHost(IntPtr reduceIntPtr, uint reduceSize) } + + //float[] t = new float[maxReductionBlocks]; + //System.Runtime.InteropServices.Marshal.Copy(reduceIntPtr, t, 0, t.Length); + //float s = t.Sum(); + sum -= TrainedModel.Bias; float ret = sum < 0 ? -1 : 1; @@ -281,7 +312,7 @@ public override void Init() SetCudaRedFunctionParams(); - SetCudaEvalFunctionParams(); + IsInitialized = true; @@ -301,11 +332,11 @@ protected void SetCudaRedFunctionParams() int offset = 0; offsetMemToReduce = offset; - cuda.SetParameter(cuFuncReduce, offset, outputCuPtr[0].Pointer); + cuda.SetParameter(cuFuncReduce, offset, evalOutputCuPtr[0].Pointer); offset += IntPtr.Size; offsetOutMemReduce = offset; - cuda.SetParameter(cuFuncReduce, offset, cuReducePtr[0].Pointer); + cuda.SetParameter(cuFuncReduce, offset, reduceCuPtr[0].Pointer); offset += IntPtr.Size; cuda.SetParameter(cuFuncReduce, offset, (uint)sizeSV); @@ -314,6 +345,7 @@ protected void SetCudaRedFunctionParams() cuda.SetParameterSize(cuFuncReduce, (uint)offset); } + protected void SetCudaData() { @@ -334,21 +366,27 @@ protected void SetCudaData() alphasPtr = cuda.CopyHostToDevice(svAlphas); - memSize = (uint)(TrainedModel.SupportElements[0].Dim * sizeof(float)); + vectorsDimMemSize = (uint)((TrainedModel.SupportElements[0].Dim+1) * sizeof(float)); for (int i = 0; i < NUM_STREAMS; i++) { stream[i] = cuda.CreateStream(); //allocates memory for one vector, size = vector dim - vecIntPtrs[i] = cuda.AllocateHost(memSize); - mainVecPtr[i] = cuda.Allocate(memSize); + mainVecIntPtrs[i] = cuda.AllocateHost(vectorsDimMemSize); + mainVecCuPtr[i] = cuda.CopyHostToDevice(mainVecIntPtrs[i], vectorsDimMemSize); //allocate memory for output, size == #SV - outputCuPtr[i] = cuda.Allocate(svAlphas); + evalOutputCuPtr[i] = cuda.Allocate(svAlphas); cuVecTexRef[i] = cuda.GetModuleTexture(cuModule, cudaVecTexRefName[i]); //cuda.SetTextureFlags(cuVecTexRef[i], 0); - cuda.SetTextureAddress(cuVecTexRef[i], mainVecPtr[i], memSize); + cuda.SetTextureAddress(cuVecTexRef[i], mainVecCuPtr[i], vectorsDimMemSize); + + uint reduceMemSize = (uint)maxReductionBlocks * sizeof(float); + reduceIntPtrs[i] = cuda.AllocateHost(reduceMemSize); + reduceCuPtr[i] = cuda.CopyHostToDevice(reduceIntPtrs[i], reduceMemSize); + //reduceCuPtr[i] = cuda.Allocate((uint)maxReductionBlocks * sizeof(float)); + } @@ -366,8 +404,8 @@ private void InitCuda() cuModule = cuda.LoadModule(Path.Combine(Environment.CurrentDirectory, cudaModuleName)); - cuFuncEval[0] = cuda.GetModuleFunction(cudaEvaluatorKernelName ); - cuFuncEval[1] = cuda.GetModuleFunction(cudaEvaluatorKernelName); + cuFuncEval = cuda.GetModuleFunction(cudaEvaluatorKernelName ); + cuFuncReduce = cuda.GetModuleFunction(cudaReduceKernelName); } diff --git a/KMLib.CUDA/GPUEvaluators/CuRBFEllILPEvaluator.cs b/KMLib.CUDA/GPUEvaluators/CuRBFEllILPEvaluator.cs index c0a0688..b44d76f 100644 --- a/KMLib.CUDA/GPUEvaluators/CuRBFEllILPEvaluator.cs +++ b/KMLib.CUDA/GPUEvaluators/CuRBFEllILPEvaluator.cs @@ -18,11 +18,9 @@ public class CuRBFEllILPEvaluator: CuEvaluator private GASS.CUDA.Types.CUdeviceptr vecLengthPtr; private GASS.CUDA.Types.CUdeviceptr selfDotPtr; - private int kernelResultParamOffset; + private float gamma; - private int vectorSelfDotParamOffset; - private int texSelParamOffset; - + public CuRBFEllILPEvaluator(float gamma) @@ -32,59 +30,55 @@ public CuRBFEllILPEvaluator(float gamma) cudaModuleName = "KernelsEllpack.cubin"; } - protected override void SetCudaEvalFunctionParams() { - - for (int i = 0; i < NUM_STREAMS; i++) - { - - cuda.SetFunctionBlockShape(cuFuncEval[i], evalThreads, 1, 1); - int offset = 0; - cuda.SetParameter(cuFuncEval[i], offset, valsPtr.Pointer); - offset += IntPtr.Size; - cuda.SetParameter(cuFuncEval[i], offset, idxPtr.Pointer); - offset += IntPtr.Size; - cuda.SetParameter(cuFuncEval[i], offset, vecLengthPtr.Pointer); - offset += IntPtr.Size; + cuda.SetFunctionBlockShape(cuFuncEval, evalThreads, 1, 1); - cuda.SetParameter(cuFuncEval[i], offset, selfDotPtr.Pointer); - offset += IntPtr.Size; + int offset = 0; + cuda.SetParameter(cuFuncEval, offset, valsPtr.Pointer); + offset += IntPtr.Size; + cuda.SetParameter(cuFuncEval, offset, idxPtr.Pointer); + offset += IntPtr.Size; + cuda.SetParameter(cuFuncEval, offset, vecLengthPtr.Pointer); + offset += IntPtr.Size; - cuda.SetParameter(cuFuncEval[i], offset, alphasPtr.Pointer); - offset += IntPtr.Size; + cuda.SetParameter(cuFuncEval, offset, selfDotPtr.Pointer); + offset += IntPtr.Size; - cuda.SetParameter(cuFuncEval[i], offset, labelsPtr.Pointer); - offset += IntPtr.Size; - kernelResultParamOffset = offset; - cuda.SetParameter(cuFuncEval[i], offset, outputCuPtr[i].Pointer); - offset += IntPtr.Size; + cuda.SetParameter(cuFuncEval, offset, alphasPtr.Pointer); + offset += IntPtr.Size; - cuda.SetParameter(cuFuncEval[i], offset, (uint)sizeSV); - offset += sizeof(int); + cuda.SetParameter(cuFuncEval, offset, labelsPtr.Pointer); + offset += IntPtr.Size; - vectorSelfDotParamOffset = offset; - cuda.SetParameter(cuFuncEval[i], offset, 0); - offset += sizeof(int); + kernelResultParamOffset = offset; + cuda.SetParameter(cuFuncEval, offset, evalOutputCuPtr[0].Pointer); + offset += IntPtr.Size; - cuda.SetParameter(cuFuncEval[i], offset, gamma); - offset += sizeof(float); + cuda.SetParameter(cuFuncEval, offset, (uint)sizeSV); + offset += sizeof(int); - texSelParamOffset = offset; - cuda.SetParameter(cuFuncEval[i], offset, i+1); - offset += sizeof(int); + vectorSelfDotParamOffset = offset; + cuda.SetParameter(cuFuncEval, offset, 0); + offset += sizeof(int); - cuda.SetParameterSize(cuFuncEval[i], (uint)offset); - } + cuda.SetParameter(cuFuncEval, offset, gamma); + offset += sizeof(float); + texSelParamOffset = offset; + cuda.SetParameter(cuFuncEval, offset, 1); + offset += sizeof(int); + cuda.SetParameterSize(cuFuncEval, (uint)offset); } + + public override void Init() { @@ -92,7 +86,7 @@ public override void Init() SetCudaDataForEllpack(); - + SetCudaEvalFunctionParams(); } private void SetCudaDataForEllpack() @@ -107,7 +101,7 @@ private void SetCudaDataForEllpack() float[] selfLinDot = TrainedModel.SupportElements.Select(c => c.DotProduct()).ToArray(); - evalBlocks = (sizeSV+evalThreads) / evalThreads; + evalBlocks = (sizeSV+evalThreads-1) / evalThreads; //copy data to device, set cuda function parameters valsPtr = cuda.CopyHostToDevice(vecVals); diff --git a/KMLib.CUDA/GPUKernels/CuRBFEllILPKernel.cs b/KMLib.CUDA/GPUKernels/CuRBFEllILPKernel.cs index 9308bc7..0fc012a 100644 --- a/KMLib.CUDA/GPUKernels/CuRBFEllILPKernel.cs +++ b/KMLib.CUDA/GPUKernels/CuRBFEllILPKernel.cs @@ -177,7 +177,7 @@ public override void Init() SetCudaFunctionParameters(); - //allocate memory for main vector, size of this vector is the same as dimenson, so many + //allocate memory for main vector, size of this vector is the same as dimension, so many //indexes will be zero, but cuda computation is faster mainVector = new float[problemElements[0].Dim + 1]; CudaHelpers.FillDenseVector(problemElements[0], mainVector); diff --git a/KMLib.CUDA/KMLib.GPU.csproj b/KMLib.CUDA/KMLib.GPU.csproj index b98244b..b9a75b8 100644 --- a/KMLib.CUDA/KMLib.GPU.csproj +++ b/KMLib.CUDA/KMLib.GPU.csproj @@ -208,7 +208,7 @@ del *.cubin, *.cu, *.ptx, *.h xcopy "$(ProjectDir)CUDAmodules\*.cu" "$(TargetDir)" /Y xcopy "$(ProjectDir)CUDAmodules\*.h" "$(TargetDir)" /Y -nvcc -I./ KernelsEllpackCol2.cu KernelsCSR.cu KernelsEllpack.cu KernelsSlicedEllpack.cu gpuFanSmoSolver.cu gpuFOSmoSolver.cu -ccbin "%25VS100COMNTOOLS%25../../VC/bin" -m64 -cubin -gencode=arch=compute_12,code=sm_12 -Xptxas="-v" +nvcc -g -G -I./ KernelsEllpackCol2.cu KernelsCSR.cu KernelsEllpack.cu KernelsSlicedEllpack.cu gpuFanSmoSolver.cu gpuFOSmoSolver.cu -ccbin "%25VS100COMNTOOLS%25../../VC/bin" -m64 -cubin -gencode=arch=compute_12,code=sm_12 -Xptxas="-v" xcopy "$(TargetDir)*.cubin" "$(SolutionDir)Libs\Cuda\" /Y xcopy "$(TargetDir)*.ptx" "$(SolutionDir)Libs\Cuda\" /Y xcopy "$(TargetDir)*.cu" "$(SolutionDir)Libs\Cuda\" /Y diff --git a/KMLib/Evaluate/DualEvaluator.cs b/KMLib/Evaluate/DualEvaluator.cs index 135fbda..4117106 100644 --- a/KMLib/Evaluate/DualEvaluator.cs +++ b/KMLib/Evaluate/DualEvaluator.cs @@ -83,16 +83,16 @@ public override float[] Predict(TProblemElement[] elements) float[] predictions = new float[elements.Length]; - Parallel.For(0, elements.Length, i => - { - - predictions[i] = Predict(elements[i]); - }); - - //for (int i = 0; i < elements.Length; i++) + //Parallel.For(0, elements.Length, i => //{ + // predictions[i] = Predict(elements[i]); - //} + //}); + + for (int i = 0; i < elements.Length; i++) + { + predictions[i] = Predict(elements[i]); + } return predictions; } diff --git a/KMLib/KMLib.csproj.ksirg.nvuser b/KMLib/KMLib.csproj.ksirg.nvuser new file mode 100644 index 0000000..aee6a1d --- /dev/null +++ b/KMLib/KMLib.csproj.ksirg.nvuser @@ -0,0 +1,6 @@ + + + + + + \ No newline at end of file diff --git a/KMLibUsageApp/Program.cs b/KMLibUsageApp/Program.cs index d9b24d6..1cee6f3 100644 --- a/KMLibUsageApp/Program.cs +++ b/KMLibUsageApp/Program.cs @@ -865,9 +865,7 @@ private static void ChooseDataSet(string dataFolder, out string trainningFile, o Model model; //Evaluator evaluator = new RBFDualEvaluator(gamma); - //Evaluator evaluator = new DualEvaluator(); - Evaluator evaluator = new CuRBFEllILPEvaluator(gamma);