////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char **argv)
{
printf("[Matrix Multiply CUBLAS] - Starting...\n");
int devID = 0, sizeMult = 5;
sMatrixSize matrix_size;
initializeCUDA(argc, argv, devID, sizeMult, matrix_size);
int matrix_result = matrixMultiply(argc, argv, devID, matrix_size);
return matrix_result;
}
void PD_flow_mrpt::initializePDFlow()
{
//Initialize Visualization
initializeScene();
//Initialize CUDA
mrpt::system::sleep(500);
initializeCUDA();
//Start video streaming
OpenCamera();
//Fill empty matrices
CaptureFrame();
createImagePyramidGPU();
CaptureFrame();
createImagePyramidGPU();
solveSceneFlowGPU();
}
/*Function: pgm_cublas_sgemm
Multiplica matriz A por matriz B.
Onde o numero de colunas de A deve ser igual ao numero de linhas de B
Em caso de erro, retorno :<b>"Numero de colunas de A diferente de numero de linhas de B"
Essa função so pode ser invocada se houve uma GPU NVidea e estiver instalado o driver do CUDA.
Com o uso dessa função se perde um pouco da precisão.
Otimização:
Essa funçao pode ser otimazada de duas formas:
- Passando true no use_temp_matrix, porem será necessario mais memoria
Parameter:
PGM_Matriz_Double* matrix_a - Matriz A
PGM_Matriz_Double* matrix_b - Matriz B
Return:
PGM_Matriz_Double* - Ponteiro para a matriz resultante
See also:
<blas_multiply>, <pgm_cublas_dgemm>, <pgm_cublas_sgemm>
*/
Datum pgm_cublas_sgemm_v2(PG_FUNCTION_ARGS){
PGM_Matriz_Double *A = (PGM_Matriz_Double*) PG_GETARG_POINTER(0);
PGM_Matriz_Double *B = (PGM_Matriz_Double*) PG_GETARG_POINTER(1);
PGM_Matriz_Double *result;
MemoryContext contextoAnterior;
if(A->n_colunas == B->n_linhas){
int info;
contextoAnterior = MemoryContextSwitchTo( CurTransactionContext );
result = pgm_create_matrix_double(A->n_linhas, B->n_colunas);
MemoryContextSwitchTo( contextoAnterior );
if( initializeCUDA() == -1)
elog(ERROR,"Ao inicializar o CUDA\n");
info = cublas_multiply_matrix_float_v2(A, B, result);
if( info ){
info *= -1;
switch((info)){
case 1: elog(ERROR,"Ao alocar memoria no dispositivo, excedeu o limite de memoria da GPU (%d)\n",info);break;
case 2: elog(ERROR,"Ao transferir dados do Host para a GPU (%d)\n",info);break;
case 3: elog(ERROR,"Ao inicializar o ambiente (%d) \n",info);break;
case 4: elog(ERROR,"Ao executar a multiplicação (%d)\n",info); break;
case 5: elog(ERROR,"Ao transferir dados da GPU para o Host (%d)\n",info);break;
case 6: elog(ERROR,"Ao alocar memoria para vetor temporario (%d)\n",info);break;
}
}
else{
PG_RETURN_POINTER(result);
}
}
else{
elog(ERROR, "Numero de colunas de A(%d) é diferente no numero de linhas de B(%d)\n", A->n_colunas, B->n_linhas);
}
PG_RETURN_VOID();
}
int main(int argc, char *argv[]) {
int i,j,k;
machineInformation currentMachine;
counterSessionInfo session;
initializeCUDA();
// Set machine information from CounterHomeBrew.h
currentMachine.cpu_model = CPU_MODEL;
currentMachine.num_sockets = NUM_SOCKETS;
currentMachine.num_phys_cores_per_socket = NUM_PHYS_CORES_PER_SOCKET;
currentMachine.num_cores_per_socket = NUM_CORES_PER_SOCKET;
currentMachine.num_cores = NUM_CORES;
currentMachine.num_cbos = NUM_PHYS_CORES_PER_SOCKET; // should multiply by NUM_SOCKETS???
currentMachine.core_gen_counter_num_max = CORE_GEN_COUNTER_MAX;
currentMachine.cbo_counter_num_max = CBO_COUNTER_NUM_MAX;
// Set session events, umasks and counters used
// int32 core_event_numbers[] = {FP_COMP_OPS_EXE_EVTNR,SIMD_FP_256_EVTNR,0x51,0xF1,0x80};
// int32 core_umasks[] = {FP_COMP_OPS_EXE_SCALAR_DOUBLE_UMASK,SIMD_FP_256_PACKED_DOUBLE_UMASK,0x01, 0x07,0x01};
session.core_gen_counter_num_used = 5;
int32 core_event_numbers[] = {0x10,0x10,0x11,0x51,0xF1};
int32 core_umasks[] = {0x20,0x40,0x01,0x01, 0x07};
session.cbo_counter_num_used = 1;
int32 cbo_event_numbers[] = {0x37};
int32 cbo_umasks[] = {0xf};
session.cbo_filter = 0x1f;
for (i = 0; i < session.core_gen_counter_num_used; i++) {
session.core_event_numbers[i] = core_event_numbers[i];
session.core_umasks[i] = core_umasks[i];
}
for (i = 0; i < session.cbo_counter_num_used; i++) {
session.cbo_event_numbers[i] = cbo_event_numbers[i];
session.cbo_umasks[i] = cbo_umasks[i];
}
int fd[NUM_CORES];
// Arrays to hold counter data...
counterData before;
counterData after;
// some data for doing a naive matmul to test flop counting...
// initloop(N);
// M,N,K are multiples of the block size....
int gpuOuter = atoi(argv[1]);
int gpuInner = atoi(argv[2]);
int cpuInner = atoi(argv[3]);
double minRuntime = atoi(argv[4]);
int Md = atoi(argv[5])*block_size;
int Nd = atoi(argv[6])*block_size;
int Kd = atoi(argv[7])*block_size;
int Mh = atoi(argv[8]);
int Nh = atoi(argv[9]);
int Kh = atoi(argv[10]);
char *ts1,*ts2,*ts3,*ts4;
char *ts5,*ts6,*ts7,*ts8;
double fineTimeStamps[8];
double gTime = 0.0;
double cTime = 0.0;
double seconds = 0.0;
int num_iters;
uint64 *coreSums;
coreSums = (uint64*)calloc(currentMachine.num_sockets*session.core_gen_counter_num_used,sizeof(uint64));
uint64 *sums;
sums = (uint64*)calloc(currentMachine.num_sockets*session.cbo_counter_num_used,sizeof(uint64));
float *Atmp = NULL;
float *Btmp = NULL;
float *Ctmp = NULL;
Atmp = (float*) malloc( Mh * Nh * sizeof(float) );
Btmp = (float*) malloc( Nh * sizeof(float) );
Ctmp = (float*) malloc( Mh * sizeof(float) );
randomInit(Atmp,Mh*Nh);
randomInit(Btmp,Nh);
for (num_iters = cpuInner; seconds < minRuntime; num_iters *=2) {
seconds = 0.0;
for (i =0; i < num_iters; i++)
BLASFUNC( CblasColMajor,CblasNoTrans,Mh,Nh, 1, Atmp,Mh, Btmp,1, 1, Ctmp,1 );
seconds = read_timer()-seconds;
}
// num_iters /= 2;
free(Atmp);
free(Btmp);
free(Ctmp);
int readyThreads = 0;
#pragma omp parallel
{
int threadNum = omp_get_thread_num();
int numThreads = omp_get_num_threads();
assert(numThreads==2);
if (threadNum == 0) {
cudaError_t error;
int memSizeA = sizeof(float)*Md*Nd;
int memSizeB = sizeof(float)*Nd;
int memSizeC = sizeof(float)*Md;
float *Ahost,*Bhost,*Chost;
// use pinned memory on the host for BW and asynch memory transfers..
int flags = cudaHostAllocDefault;
ts5 = getTimeStamp();
fineTimeStamps[0] = read_timer();
error = cudaHostAlloc((void**)&Ahost,memSizeA,flags);if (error != cudaSuccess){printf("cudaHostMalloc Ahost returned error code %d, line(%d)\n", error, __LINE__);exit(EXIT_FAILURE);}
error = cudaHostAlloc((void**)&Bhost,memSizeB,flags);if (error != cudaSuccess){printf("cudaHostMalloc Bhost returned error code %d, line(%d)\n", error, __LINE__);exit(EXIT_FAILURE);}
error = cudaHostAlloc((void**)&Chost,memSizeC,flags);if (error != cudaSuccess){printf("cudaHostMalloc Chost returned error code %d, line(%d)\n", error, __LINE__);exit(EXIT_FAILURE);}
// set local arrays
randomInit(Ahost,Md*Nd);
randomInit(Bhost,Nd);
// allocate device memory
float *Adevice,*Bdevice,*Cdevice;
error = cudaMalloc((void**)&Adevice,memSizeA); if (error != cudaSuccess){printf("cudaMalloc Adevice returned error code %d, line(%d)\n", error, __LINE__);exit(EXIT_FAILURE);}
error = cudaMalloc((void**)&Bdevice,memSizeB); if (error != cudaSuccess){printf("cudaMalloc Bdevice returned error code %d, line(%d)\n", error, __LINE__);exit(EXIT_FAILURE);}
error = cudaMalloc((void**)&Cdevice,memSizeC); if (error != cudaSuccess){printf("cudaMalloc Cdevice returned error code %d, line(%d)\n", error, __LINE__);exit(EXIT_FAILURE);}
fineTimeStamps[1] = read_timer();
ts6 = getTimeStamp();
#pragma omp critical
{
readyThreads += 1;
}
// fprintf(stderr,"Incremented ready GPU\n");
while (readyThreads < 2){sleep(1);fprintf(stderr,"Thread 0: %d\n",readyThreads);};
//#pragma omp single
//{
cudaStream_t stream1;
cudaStreamCreate ( &stream1) ;
ts3 = getTimeStamp();
fineTimeStamps[2] = read_timer();
gTime = read_timer();
for (int i = 0; i < gpuOuter; i++)
GPUsgemv(gpuInner,Md,Nd,Kd,Adevice,Bdevice,Cdevice,Ahost,Bhost,Chost,&stream1);
cudaStreamSynchronize(stream1);
gTime = read_timer() - gTime;
fineTimeStamps[3] = read_timer();
ts4 = getTimeStamp();
cudaFreeHost(Ahost);
cudaFreeHost(Bhost);
cudaFreeHost(Chost);
} else {
// uint64 min_iters = strtoull(argv[4],NULL,0);
float *A = NULL;
float *B = NULL;
float *C = NULL;
ts7 = getTimeStamp();
fineTimeStamps[4] = read_timer();
A = (float*) malloc( Mh * Nh * sizeof(float) );
B = (float*) malloc( Nh * sizeof(float) );
C = (float*) malloc( Mh * sizeof(float) );
randomInit(A,Mh*Nh);
randomInit(B,Nh);
fineTimeStamps[5] = read_timer();
ts8 = getTimeStamp();
#pragma omp critical
{
readyThreads += 1;
}
// fprintf(stderr,"Incremented ready CPU\n");
while (readyThreads < 2){sleep(1);fprintf(stderr,"Thread 1: %d\n",readyThreads);};
// open the msr files for each core on the machine
for (i = 0; i < currentMachine.num_cores; i++)
open_msr_file(i,&fd[i]);
int socketsProgrammed = 0;
for (i = 0; i < currentMachine.num_cores; i++) {
int currentCoreFD = fd[i];
stopCounters(i, currentCoreFD, ¤tMachine, &session);
programCoreFixedCounters(currentCoreFD);
programGeneralPurposeRegisters(currentCoreFD, ¤tMachine, &session);
/* Program the Uncore as desired...*/
// Only program the first physical core on each socket.
// NOTE: Some assumptions about topology here...check /proc/cpuinfo to confirm.
if (i % currentMachine.num_phys_cores_per_socket == 0 && socketsProgrammed < currentMachine.num_sockets) {
programUncoreCounters( currentCoreFD, ¤tMachine, &session);
socketsProgrammed++;
}
}
seconds = 0.0;
// start the programmed counters...
for (i = 0; i < currentMachine.num_cores; i++)
startCounters( i, fd[i], ¤tMachine, &session);
/* READ COUNTERS BEFORE STUFF */
readCounters(fd,¤tMachine,&session, &before);
ts1 = getTimeStamp();
fineTimeStamps[6] = read_timer();
seconds = read_timer();
/* DO STUFF */
for (i =0; i < num_iters; i++)
BLASFUNC( CblasColMajor,CblasNoTrans,Mh,Nh, 1, A,Mh, B,1, 1, C,1 );
/* END DOING STUFF */
seconds = read_timer()-seconds;
fineTimeStamps[7] = read_timer();
ts2 = getTimeStamp();
/* READ COUNTERS AFTER STUFF */
for (i = 0; i < currentMachine.num_cores; i++)
stopCounters(i,fd[i],¤tMachine, &session);
// printf("num_iters = %"PRIu64", runtime is %g\n",num_iters,seconds);
readCounters(fd,¤tMachine,&session,&after);
diffCounterData(¤tMachine, &session, &after, &before, &after);
for (i = 0; i < currentMachine.num_sockets; i++) {
// printf("Socket %d\n",i);
for (j = 0; j < currentMachine.num_cores_per_socket; j++) {
// printf("%d,",j);
for (k = 0; k < session.core_gen_counter_num_used; k++){
// printf("%"PRIu64",",after.generalCore[i*currentMachine.num_cores_per_socket + j][k]);
// bug in the indexing of the core sums???
// coreSums[i*session.core_gen_counter_num_used + k] += after.generalCore[i*currentMachine.num_cores_per_socket + j][k];
coreSums[k] += after.generalCore[i*currentMachine.num_cores_per_socket + j][k];
}
// printf("\n");
}
}
for (i = 0; i < currentMachine.num_sockets; i++) {
// printf("%d,",i);
for (j = 0; j < currentMachine.num_cbos; j++) {
// printf("%d,",j);
for (k = 0; k < session.cbo_counter_num_used; k++) {
// printf("%llu,",after.cboUncore[i*currentMachine.num_phys_cores_per_socket + j][k]);
// bug in the indexing of the core sums???
// sums[i*session.cbo_counter_num_used + k] += after.cboUncore[i*currentMachine.num_phys_cores_per_socket + j][k];
sums[k] += after.cboUncore[i*currentMachine.num_phys_cores_per_socket + j][k];
}
}
}
// printf("\n");
// Stop counters, reset PMU, close msr files
cleanup(fd,¤tMachine,&session);
free(A);
free(B);
free(C);
}
} // end parallel region
printf("%s,%s,%s,%s,%s,%s,%s,%s,%d,%d,%d,%d,%d,%d,%d,%d,%d,%f,%f,%f,",ts7,ts8,ts1,ts2,ts5,ts6,ts3,ts4,Mh,Nh,Kh,Md/block_size,Nd/block_size,Kd/block_size,num_iters,gpuOuter,gpuInner,seconds,gTime,(float)(gpuOuter*(Md*Kd+Nd+Md))/16.0);
for (int i = 0; i < 8; i++)
printf("%f,",fineTimeStamps[i]);
for (j = 0; j < session.core_gen_counter_num_used; j++)
printf("%llu,",coreSums[j]);
for (j = 0; j < session.cbo_counter_num_used; j++)
if (j == session.cbo_counter_num_used-1)
printf("%llu",sums[j]);
else
printf("%llu,",sums[j]);
printf("\n");
free(sums);
free(coreSums);
return 0;
}
