Hi,
I am having 2 different card K600 (Keppler) and K620 (Maxwell)
But I am not being able to achive the performance benefits from K620 GPU :(
previously my project was built with Cuda 4.0 - in this forum with responce to my previous post related one accuracy settings I changed my project settings from CudA 4 to cuda 6.5. That improved the performance for both K600 and K620 - but still K600 is performing much better than K620
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cuda_runtime.h>
#include <stdio.h>
#include <Windows.h>
#include <io.h>
#include <iostream>
#include <fstream>
#include <stdio.h>
using namespace std;
using namespace std;
__global__ void AddPixel(int* pMat1);
void Writeintdata(char* szFileName,int val) {
std::ofstream oFile, oFilexy;
remove(szFileName);
char szbuff[100];
sprintf(szbuff, "%d ", val);
oFile.open(szFileName);
oFile.write(szbuff, strlen(szbuff));
oFile.write("\n", 1);
oFile.close();
}
//int blockSize; // The launch configurator returned block size
//int minGridSize; // The minimum grid size needed to achieve the
//// maximum occupancy for a full device launch
//int gridSize; // The actual grid size needed, based on input size
void MatLaunch(int* pMat1, int DIMX)
{
//int threadsPerBlock = 256;
//int numBlocks = 4096;
int* pDevMat1 = NULL;
int size = DIMX * sizeof(int);
int iStatus = cudaMalloc( &pDevMat1, size);
if(iStatus == cudaSuccess)
{
iStatus = cudaMemcpy(pDevMat1, pMat1, size,cudaMemcpyHostToDevice);
}
if(iStatus == cudaSuccess)
{
AddPixel<<<4096,256>>>(pDevMat1);
cudaDeviceSynchronize();
iStatus = cudaGetLastError();
}
if(iStatus ==cudaSuccess)
{
iStatus = cudaMemcpy(pMat1, pDevMat1, size,cudaMemcpyDeviceToHost);
}
if(iStatus == cudaSuccess)
{
iStatus = cudaFree(pDevMat1);
}
}
__global__ void AddPixel(int* pMat1)
{
int X = blockIdx.x * blockDim.x + threadIdx.x;
pMat1[X] = pMat1[X] + 25;
}
int main()
{
std::ofstream myfile;
myfile.open ("D:\\CUDA\\CUDA6.5_Old_Timings.csv", std::ios::out | std::ios::app );
__int64 ctr1 = 0, ctr2 = 0;
double wpcnt =0;
long size = 1024*1024;
int* pMatLaunch = new int;
for(int i = 0;i<size;i++)
{
pMatLaunch[i] = 50;
}
//cudaOccupancyMaxPotentialBlockSize( &minGridSize, &blockSize,
// AddPixel, 0, size);
//// Round up according to array size
//gridSize = (size + blockSize - 1) / blockSize;
myfile <<endl;
myfile << "1kx1k Start";
myfile <<endl;
for(int i = 0;i<30;i++) // I am taking 30 iterations and taking avrg time (excluding first value)
{
QueryPerformanceCounter((LARGE_INTEGER *)&ctr1);
MatLaunch(pMatLaunch,size);
QueryPerformanceCounter((LARGE_INTEGER *)&ctr2);
wpcnt = ctr2-ctr1;
myfile << "Time taken for 1kx1k image is "<<","<< wpcnt<<endl;
}
myfile << "1kx1k Complete";
myfile << "-------------";
myfile <<endl;
ctr1 = 0, ctr2 = 0;
size = 2048*2048;
pMatLaunch = new int;
for(int i = 0;i<size;i++)
{
pMatLaunch[i] = 50;
}
myfile <<endl;
myfile << "2kx2k Start";
myfile <<endl;
cudaOccupancyMaxPotentialBlockSize( &minGridSize, &blockSize,
AddPixel, 0, size);
// Round up according to array size
gridSize = (size + blockSize - 1) / blockSize;
for(int i = 0;i<30;i++)
{
QueryPerformanceCounter((LARGE_INTEGER *)&ctr1);
MatLaunch(pMatLaunch,size);
QueryPerformanceCounter((LARGE_INTEGER *)&ctr2);
wpcnt = ctr2-ctr1;
myfile << "Time taken for 2kx2k image is "<<","<< wpcnt<<endl;
}
myfile << "2kx2k Complete";
myfile << "-------------";
myfile <<endl;
myfile <<endl;
ctr1 = 0, ctr2 = 0;
size = 4096*4096;
pMatLaunch = new int;
for(int i = 0;i<size;i++)
{
pMatLaunch[i] = 50;
}
myfile << "4kx4k Start";
myfile <<endl;
cudaOccupancyMaxPotentialBlockSize( &minGridSize, &blockSize,
AddPixel, 0, size);
// Round up according to array size
gridSize = (size + blockSize - 1) / blockSize;
for(int i = 0;i<30;i++)
{
QueryPerformanceCounter((LARGE_INTEGER *)&ctr1);
MatLaunch(pMatLaunch,size);
QueryPerformanceCounter((LARGE_INTEGER *)&ctr2);
wpcnt = ctr2-ctr1;
myfile << "Time taken for 4kx4k image is "<<","<< wpcnt<<endl;
}
myfile << "4kx4k Complete";
myfile << "-------------";
myfile <<endl;
myfile <<endl;
ctr1 = 0, ctr2 = 0;
size = 6000*6000;
pMatLaunch = new int;
for(int i = 0;i<size;i++)
{
pMatLaunch[i] = 50;
}
myfile << "6kx6k Start";
myfile <<endl;
cudaOccupancyMaxPotentialBlockSize( &minGridSize, &blockSize,
AddPixel, 0, size);
// Round up according to array size
gridSize = (size + blockSize - 1) / blockSize;
for(int i = 0;i<30;i++)
{
QueryPerformanceCounter((LARGE_INTEGER *)&ctr1);
MatLaunch(pMatLaunch,size);
QueryPerformanceCounter((LARGE_INTEGER *)&ctr2);
wpcnt = ctr2-ctr1;
myfile << "Time taken for 6kx6k image is "<<","<< wpcnt<<endl;
}
myfile << "6kx6k Complete";
myfile << "-------------";
myfile <<endl;
ctr1 = 0, ctr2 = 0;
size = 10000*10000;
pMatLaunch = new int;
for(int i = 0;i<size;i++)
{
pMatLaunch[i] = 50;
}
myfile << "10kx10k Start";
myfile <<endl;
cudaOccupancyMaxPotentialBlockSize( &minGridSize, &blockSize,
AddPixel, 0, size);
// Round up according to array size
gridSize = (size + blockSize - 1) / blockSize;
for(int i = 0;i<30;i++)
{
QueryPerformanceCounter((LARGE_INTEGER *)&ctr1);
MatLaunch(pMatLaunch,size);
QueryPerformanceCounter((LARGE_INTEGER *)&ctr2);
wpcnt = ctr2-ctr1;
myfile << "Time taken for 10kx10k image is "<<","<< wpcnt<<endl;
}
myfile << "10kx10k Complete";
myfile << "-------------";
myfile <<endl;
myfile.close();
return 0;
}
I tried to change the block size from 256 to 512 in case of matrix size 2K X 2K and above. but still no improvemnt in perfomance.
— With all the possible combinations I am trying always K600 is showing better performance.
Project settings I am using same as what is being used by 6.5 sample program → C:\ProgramData\NVIDIA Corporation\CUDA Samples\v6.5\0_Simple\template_runtime
— If some one can check this and let me know if the problem is in coding or project settings and what changes are needed so that my K620 card can outperform the K600 card - it will be very helpful.
Incase there is some changes needed in project settings it will be very helpful if you upload the project setting file some where - so that I can download and try
Thanks