public override void Init()
{
linKernel.ProblemElements = problemElements;
linKernel.Y = Y;
linKernel.Init();
base.Init();
float[] vecVals;
int[] vecColIdx;
int[] vecLenght;
int align = preFetch;
CudaHelpers.TransformToEllpackRFormat(out vecVals, out vecColIdx, out vecLenght, problemElements, align);
// CudaHelpers.TransformToEllpackRFormat(out vecVals, out vecColIdx, out vecLenght, problemElements);
selfLinDot = linKernel.DiagonalDotCache;
#region cuda initialization
InitCudaModule();
//copy data to device, set cuda function parameters
valsPtr = cuda.CopyHostToDevice(vecVals);
idxPtr = cuda.CopyHostToDevice(vecColIdx);
vecLengthPtr = cuda.CopyHostToDevice(vecLenght);
labelsPtr = cuda.CopyHostToDevice(Y);
selfLinDotPtr = cuda.CopyHostToDevice(selfLinDot);
uint memSize = (uint)(2 * problemElements.Length * sizeof(float));
//allocate mapped memory for our results
//CUDARuntime.cudaSetDeviceFlags(CUDARuntime.cudaDeviceMapHost);
// var e= CUDADriver.cuMemHostAlloc(ref outputIntPtr, memSize, 8);
//CUDARuntime.cudaHostAlloc(ref outputIntPtr, memSize, CUDARuntime.cudaHostAllocMapped);
//var errMsg=CUDARuntime.cudaGetErrorString(e);
//cuda.HostRegister(outputIntPtr,memSize, Cuda)
outputIntPtr = cuda.HostAllocate(memSize, CUDADriver.CU_MEMHOSTALLOC_DEVICEMAP);
outputPtr = cuda.GetHostDevicePointer(outputIntPtr, 0);
//normal memory allocation
//outputPtr = cuda.Allocate((uint)(sizeof(float) * problemElements.Length));
#endregion
SetCudaFunctionParameters();
//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
VectorI = new float[problemElements[0].Dim + 1];
VectorJ = new float[problemElements[0].Dim + 1];
CudaHelpers.FillDenseVector(problemElements[0], VectorI);
CudaHelpers.FillDenseVector(problemElements[1], VectorJ);
CudaHelpers.SetTextureMemory(cuda, cuModule, ref cuVecI_TexRef, cuVecITexRefName, VectorI, ref VecIPtr);
CudaHelpers.SetTextureMemory(cuda, cuModule, ref cuVecJ_TexRef, cuVecJTexRefName, VectorJ, ref VecJPtr);
}
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);
}
