public void TransformToCSCFormatTest() { float[] vecVals = null; // TODO: Initialize to an appropriate value int[] vecIdx = null; // TODO: Initialize to an appropriate value int[] vecLenght = null; // TODO: Initialize to an appropriate value int[] vecIdxExpected = new int[] { 0, 2, 4, 1, 3, 0, 2, 4, 1, 3, 0, 2, 4, 1, 3, 0, 2, 4, 1, 3, 0, 2, 4, 1, 3 }; float[] vecValsExpected = new float[] { 0, 2, 4, 1, 3, 0, 2, 4, 1, 3, 0, 2, 4, 1, 3, 0, 2, 4, 1, 3, 0, 2, 4, 1, 3 }; int[] vecLenghtExpected = new int[] { 0, 3, 5, 8, 10, 13, 15, 18, 20, 23, 25 }; CudaHelpers.TransformToCSCFormat(out vecVals, out vecIdx, out vecLenght, problemElements); for (int i = 0; i < vecIdxExpected.Length; i++) { Assert.AreEqual(vecIdxExpected[i], vecIdx[i], "vec idx different at position" + i); } for (int i = 0; i < vecValsExpected.Length; i++) { Assert.AreEqual(vecValsExpected[i], vecVals[i], "vec val different at position" + i); } for (int i = 0; i < vecLenghtExpected.Length; i++) { Assert.AreEqual(vecLenghtExpected[i], vecLenght[i], "vec lenght different at position" + i); } }
private void SetCudaData(Problem <SparseVec> sub_prob) { int vecDim = sub_prob.FeaturesCount;//.Elements[0].Dim; /* * copy vectors to CUDA device */ #region copy trainning examples to GPU float[] vecVals; int[] vecIdx; int[] vecLenght; CudaHelpers.TransformToCSRFormat(out vecVals, out vecIdx, out vecLenght, sub_prob.Elements); valsCSRPtr = cuda.CopyHostToDevice(vecVals); idxCSRPtr = cuda.CopyHostToDevice(vecIdx); vecLenghtCSRPtr = cuda.CopyHostToDevice(vecLenght); CudaHelpers.TransformToCSCFormat(out vecVals, out vecIdx, out vecLenght, sub_prob.Elements); valsCSCPtr = cuda.CopyHostToDevice(vecVals); idxCSCPtr = cuda.CopyHostToDevice(vecIdx); vecLenghtCSCPtr = cuda.CopyHostToDevice(vecLenght); #endregion /* * allocate memory for gradient */ alphaMemSize = (uint)(sub_prob.ElementsCount * sizeof(float)); gradPtr = cuda.Allocate(alphaMemSize); gradOldPtr = cuda.Allocate(alphaMemSize); alphaPtr = cuda.Allocate(alphaMemSize); alphaOldPtr = cuda.Allocate(alphaMemSize); alphaTmpPtr = cuda.Allocate(alphaMemSize); /* * reduction blocks for computing Obj */ GetNumThreadsAndBlocks(vecDim, 64, threadsPerBlock, ref threadsForReduceObjW, ref bpgReduceW); reduceObjW = new float[bpgReduceW]; uint reduceWBytes = (uint)bpgReduceW * sizeof(float); reduceObjWPtr = cuda.Allocate(reduceWBytes); /* * reduction size for kernels which operate on alpha */ int reductionSize = problem.ElementsCount; threadsForReduceObjAlpha = 0; GetNumThreadsAndBlocks(problem.ElementsCount, 64, threadsPerBlock, ref threadsForReduceObjAlpha, ref bpgReduceAlpha); uint alphaReductionBytes = (uint)bpgReduceAlpha * sizeof(float); /* * reduction array for computing objective function value */ reduceObjAlpha = new float[bpgReduceAlpha]; reduceObjAlphaPtr = cuda.Allocate(alphaReductionBytes); /* * reduction arrays for computing BB step */ alphaPartReduce = new float[bpgReduceAlpha]; gradPartReduce = new float[bpgReduceAlpha]; alphaGradPartReduce = new float[bpgReduceAlpha]; reduceBBAlphaGradPtr = cuda.Allocate(alphaReductionBytes); reduceBBAlphaPtr = cuda.Allocate(alphaReductionBytes); reduceBBGradPtr = cuda.Allocate(alphaReductionBytes); /* * reduction arrays for comuting lin part */ reduceLinPart = new float[bpgReduceAlpha]; reduceLinPartPtr = cuda.Allocate(alphaReductionBytes); //float[] wVec = new float[vecDim]; wVecMemSize = (uint)vecDim * sizeof(float); wTempVecPtr = cuda.Allocate(wVecMemSize); //move W wector SetTextureMemory(ref cuWVecTexRef, cudaWVecTexRefName, ref wVecPtr, wVecMemSize); //set texture memory for labels SetTextureMemory(ref cuLabelsTexRef, cudaLabelsTexRefName, sub_prob.Y, ref labelsPtr); SetTextureMemory(ref cuDeltasTexRef, "deltasTexRef", ref deltasPtr, alphaMemSize); diagPtr = cuda.GetModuleGlobal(cuModule, "diag_shift"); stepBBPtr = cuda.GetModuleGlobal(cuModule, "stepBB"); float[] stepData = new float[] { 0.1f }; cuda.CopyHostToDevice(stepBBPtr, stepData); SetCudaParameters(sub_prob); }
private void SetCudaData(Problem <SparseVec> sub_prob) { int vecDim = sub_prob.Elements[0].Dim; /* * copy vectors to CUDA device */ float[] vecVals; int[] vecIdx; int[] vecLenght; CudaHelpers.TransformToCSRFormat(out vecVals, out vecIdx, out vecLenght, sub_prob.Elements); valsCSRPtr = cuda.CopyHostToDevice(vecVals); idxCSRPtr = cuda.CopyHostToDevice(vecIdx); vecLenghtCSRPtr = cuda.CopyHostToDevice(vecLenght); CudaHelpers.TransformToCSCFormat(out vecVals, out vecIdx, out vecLenght, sub_prob.Elements); valsCSCPtr = cuda.CopyHostToDevice(vecVals); idxCSCPtr = cuda.CopyHostToDevice(vecIdx); vecLenghtCSCPtr = cuda.CopyHostToDevice(vecLenght); /* * allocate memory for gradient */ uint memSize = (uint)(sub_prob.ElementsCount * sizeof(float)); //allocate mapped memory for our results (dot product beetween vector W and all elements) gradIntPtr = cuda.HostAllocate(memSize, CUDADriver.CU_MEMHOSTALLOC_DEVICEMAP); gradPtr = cuda.GetHostDevicePointer(gradIntPtr, 0); //allocate memory for main vector, size of this vector is the same as dimenson, so many //indexes will be zero, but cuda computation is faster mainVector = new float[vecDim]; //move W wector //CudaHelpers.FillDenseVector(problemElements[0], mainVector); CudaHelpers.SetTextureMemory(cuda, cuModule, ref cuMainVecTexRef, cudaMainVecTexRefName, mainVector, ref mainVecPtr); //set texture memory for labels CudaHelpers.SetTextureMemory(cuda, cuModule, ref cuLabelsTexRef, cudaLabelsTexRefName, sub_prob.Y, ref labelsPtr); /* * data for cuda solver */ //normaly for L2 solver QDii= xii*xii+Diag_i //where Diag_i = 0.5/Cp if yi=1 // Diag_i = 0.5/Cn if yi=-1 //but we will add this on GPU QD = new float[sub_prob.ElementsCount]; alpha = new float[sub_prob.ElementsCount]; deltas = new float[sub_prob.ElementsCount]; float[] diag = new float[3]; for (int i = 0; i < sub_prob.ElementsCount; i++) { QD[i] = sub_prob.Elements[i].DotProduct(); alpha[i] = 0f; deltas[i] = 0; } qdPtr = cuda.CopyHostToDevice(QD); alphaPtr = cuda.Allocate(alpha); //deltasPtr = cuda.Allocate(deltas); CudaHelpers.SetTextureMemory(cuda, cuModule, ref cuDeltasTexRef, "deltasTexRef", deltas, ref deltasPtr); diagPtr = cuda.GetModuleGlobal(cuModule, "diag_shift"); //set this in fill function //cuda.CopyHostToDevice(diagPtr, diag); //CUdeviceptr dimPtr = cuda.GetModuleGlobal(cuModule, "Dim"); ////todo: check if it ok ////cuda.Memset(dimPtr,(uint) vecDim, 1); //int[] dimArr = new int[] { vecDim }; //cuda.CopyHostToDevice(dimPtr,dimArr); //CUDARuntime.cudaMemcpyToSymbol("Dim", dimPtr, 1, 0, cudaMemcpyKind.cudaMemcpyHostToDevice); //CUDARuntime.cudaMemcpyToSymbol("Dim", ,1,0, cudaMemcpyKind.cudaMemcpyHostToDevice); CUdeviceptr deltaScalingPtr = cuda.GetModuleGlobal(cuModule, "stepScaling"); //two ways of computing scaling param, should be the same, but it depends on rounding. //stepScaling = (float)(1.0 / Math.Sqrt(sub_prob.ElementsCount)); stepScaling = 0.0002f;// (float)(1.0 / sub_prob.ElementsCount); //set scaling constant float[] scArr = new float[] { stepScaling }; cuda.CopyHostToDevice(deltaScalingPtr, scArr); //cuda.Memset(deltaScalingPtr, (uint) scaling,sizeof(float)); //cuda.CopyHostToDevice(dimPtr, problem.Elements[0].Dim); SetCudaParameters(sub_prob); }