From 3f59ec40cd2f98d9cee3ef3e7ba14db95ee27adc Mon Sep 17 00:00:00 2001 From: Vitor Santos Costa Date: Tue, 28 Jun 2016 23:47:09 +0100 Subject: [PATCH] python support --- packages/ProbLog/problog/timer.yap | 1 - packages/ProbLog/problog_learning.yap | 4 +- packages/cuda/CC_CSSTree.cu | 0 packages/cuda/CC_CSSTree.h | 0 packages/cuda/Makefile.in | 0 packages/cuda/bpreds.cu | 37 +- packages/cuda/bpreds.h | 1 + packages/cuda/bpredscpu.cpp | 0 packages/cuda/creator2.c | 4 +- packages/cuda/cuda.c | 0 packages/cuda/cuda.yap | 0 packages/cuda/dbio.cu | 20 +- packages/cuda/joincpu.cpp | 62 +- packages/cuda/lista.cu | 40 +- packages/cuda/lista.h | 0 packages/cuda/memory.cu | 44 +- packages/cuda/memory.h | 0 packages/cuda/pred.h | 0 packages/cuda/selectproyect.cu | 103 +- packages/cuda/selectproyectcpu.cpp | 0 packages/cuda/treeb.cu | 347 +++---- packages/cuda/union2.cu | 80 +- packages/cuda/union2.h | 0 packages/cuda/unioncpu2.cpp | 0 packages/myddas/example | 11 +- packages/myddas/pl/myddas.ypp | 11 +- packages/myddas/sqlite3/CMakeLists.txt | 9 +- packages/python/setup.py.cmake | 10 + packages/python/yap_kernel/prolog.js | 1283 ++++++++++++++++++++++++ packages/yap-lbfgs/yap_lbfgs.c | 3 - 30 files changed, 1684 insertions(+), 386 deletions(-) mode change 100755 => 100644 packages/cuda/CC_CSSTree.cu mode change 100755 => 100644 packages/cuda/CC_CSSTree.h mode change 100755 => 100644 packages/cuda/Makefile.in mode change 100755 => 100644 packages/cuda/bpreds.cu mode change 100755 => 100644 packages/cuda/bpreds.h mode change 100755 => 100644 packages/cuda/bpredscpu.cpp mode change 100755 => 100644 packages/cuda/creator2.c mode change 100755 => 100644 packages/cuda/cuda.c mode change 100755 => 100644 packages/cuda/cuda.yap mode change 100755 => 100644 packages/cuda/joincpu.cpp mode change 100755 => 100644 packages/cuda/lista.cu mode change 100755 => 100644 packages/cuda/lista.h mode change 100755 => 100644 packages/cuda/memory.cu mode change 100755 => 100644 packages/cuda/memory.h mode change 100755 => 100644 packages/cuda/pred.h mode change 100755 => 100644 packages/cuda/selectproyect.cu mode change 100755 => 100644 packages/cuda/selectproyectcpu.cpp mode change 100755 => 100644 packages/cuda/treeb.cu mode change 100755 => 100644 packages/cuda/union2.cu mode change 100755 => 100644 packages/cuda/union2.h mode change 100755 => 100644 packages/cuda/unioncpu2.cpp create mode 100644 packages/python/setup.py.cmake create mode 100644 packages/python/yap_kernel/prolog.js diff --git a/packages/ProbLog/problog/timer.yap b/packages/ProbLog/problog/timer.yap index 12a3927f9..6bf324cac 100644 --- a/packages/ProbLog/problog/timer.yap +++ b/packages/ProbLog/problog/timer.yap @@ -240,7 +240,6 @@ timer_stop(Name,Duration) :- -> statistics(walltime,[StopTime,_]), Duration is StopTime-StartTime; - throw(timer_not_started(timer_stop(Name,Duration))) ). diff --git a/packages/ProbLog/problog_learning.yap b/packages/ProbLog/problog_learning.yap index ea5c42add..b1127bf13 100644 --- a/packages/ProbLog/problog_learning.yap +++ b/packages/ProbLog/problog_learning.yap @@ -1,4 +1,4 @@ -1%%% -*- Mode: Prolog; -*- +%%% -*- Mode: Prolog; -*- %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % @@ -567,6 +567,8 @@ init_learning :- set_problog_flag(alpha,Alpha) ) ) + ; + true ), %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% diff --git a/packages/cuda/CC_CSSTree.cu b/packages/cuda/CC_CSSTree.cu old mode 100755 new mode 100644 diff --git a/packages/cuda/CC_CSSTree.h b/packages/cuda/CC_CSSTree.h old mode 100755 new mode 100644 diff --git a/packages/cuda/Makefile.in b/packages/cuda/Makefile.in old mode 100755 new mode 100644 diff --git a/packages/cuda/bpreds.cu b/packages/cuda/bpreds.cu old mode 100755 new mode 100644 index 50e216fb0..4fb5a8153 --- a/packages/cuda/bpreds.cu +++ b/packages/cuda/bpreds.cu @@ -1,3 +1,4 @@ +#include "hip/hip_runtime.h" #include #include #include @@ -25,10 +26,10 @@ int maximo(int count, ...) __global__ void bpreds(int *dop1, int *dop2, int rows, int of1, int of2, int *cons, int numc, int nx, int *res, int *res2) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int x, rowact, rowact1, op1, op2; - if(threadIdx.x < numc) - shared[threadIdx.x] = cons[threadIdx.x]; + if(hipThreadIdx_x < numc) + shared[hipThreadIdx_x] = cons[hipThreadIdx_x]; __syncthreads(); if(id < rows) { @@ -110,10 +111,10 @@ __global__ void bpreds(int *dop1, int *dop2, int rows, int of1, int of2, int *co __global__ void bpredsnormal2(int *dop1, int rows, int of1, int *cons, int numc, int *res) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int x, rowact, op1, op2; - if(threadIdx.x < numc) - shared[threadIdx.x] = cons[threadIdx.x]; + if(hipThreadIdx_x < numc) + shared[hipThreadIdx_x] = cons[hipThreadIdx_x]; __syncthreads(); if(id < rows) { @@ -159,10 +160,10 @@ __global__ void bpredsnormal2(int *dop1, int rows, int of1, int *cons, int numc, __global__ void bpredsnormal(int *dop1, int rows, int of1, int *cons, int numc, int *res) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int x, rowact, op1, op2; - if(threadIdx.x < numc) - shared[threadIdx.x] = cons[threadIdx.x]; + if(hipThreadIdx_x < numc) + shared[hipThreadIdx_x] = cons[hipThreadIdx_x]; __syncthreads(); if(id < rows) { @@ -226,10 +227,10 @@ __global__ void bpredsnormal(int *dop1, int rows, int of1, int *cons, int numc, __global__ void bpredsOR(int *dop1, int *dop2, int rows, int of1, int of2, int *cons, int numc, int nx, int *res, int *res2) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int x, rowact, rowact1, op1, op2; - if(threadIdx.x < numc) - shared[threadIdx.x] = cons[threadIdx.x]; + if(hipThreadIdx_x < numc) + shared[hipThreadIdx_x] = cons[hipThreadIdx_x]; __syncthreads(); if(id < rows) { @@ -344,10 +345,10 @@ __global__ void bpredsOR(int *dop1, int *dop2, int rows, int of1, int of2, int * __global__ void bpredsorlogic2(int *dop1, int rows, int of1, int *cons, int numc, int *res) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int x, rowact, op1, op2; - if(threadIdx.x < numc) - shared[threadIdx.x] = cons[threadIdx.x]; + if(hipThreadIdx_x < numc) + shared[hipThreadIdx_x] = cons[hipThreadIdx_x]; __syncthreads(); if(id < rows) { @@ -411,10 +412,10 @@ __global__ void bpredsorlogic2(int *dop1, int rows, int of1, int *cons, int numc __global__ void bpredsorlogic(int *dop1, int rows, int of1, int *cons, int numc, int *res) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int x, rowact, op1, op2; - if(threadIdx.x < numc) - shared[threadIdx.x] = cons[threadIdx.x]; + if(hipThreadIdx_x < numc) + shared[hipThreadIdx_x] = cons[hipThreadIdx_x]; __syncthreads(); if(id < rows) { diff --git a/packages/cuda/bpreds.h b/packages/cuda/bpreds.h old mode 100755 new mode 100644 index dc81e08df..2f1fb92e0 --- a/packages/cuda/bpreds.h +++ b/packages/cuda/bpreds.h @@ -1,3 +1,4 @@ +#include "hip/hip_runtime.h" #ifndef _BPREDS_H_ #define _BPREDS_H_ diff --git a/packages/cuda/bpredscpu.cpp b/packages/cuda/bpredscpu.cpp old mode 100755 new mode 100644 diff --git a/packages/cuda/creator2.c b/packages/cuda/creator2.c old mode 100755 new mode 100644 index 54534bd61..1a702c8a3 --- a/packages/cuda/creator2.c +++ b/packages/cuda/creator2.c @@ -66,7 +66,7 @@ int main(int argc, char *argv[]) fprintf(cuda, "\t\t\t{\n"); fprintf(cuda, "\t\t\t\tsize = nrows * tipo * sizeof(int);\n"); fprintf(cuda, "\t\t\t\treservar(&nres, size);\n"); - fprintf(cuda, "\t\t\t\tcudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice);\n"); + fprintf(cuda, "\t\t\t\tcudaMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice);\n"); fprintf(cuda, "\t\t\t\tcudaFree(*ret);\n"); fprintf(cuda, "\t\t\t\t*ret = nres;\n"); fprintf(cuda, "\t\t\t}\n"); @@ -103,7 +103,7 @@ int main(int argc, char *argv[]) fprintf(cuda, "\t\t\t{\n"); fprintf(cuda, "\t\t\t\tsize = nrows * tipo * sizeof(int);\n"); fprintf(cuda, "\t\t\t\treservar(&nres, size);\n"); - fprintf(cuda, "\t\t\t\tcudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice);\n"); + fprintf(cuda, "\t\t\t\tcudaMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice);\n"); fprintf(cuda, "\t\t\t\tcudaFree(*ret);\n"); fprintf(cuda, "\t\t\t\t*ret = nres;\n"); fprintf(cuda, "\t\t\t}\n"); diff --git a/packages/cuda/cuda.c b/packages/cuda/cuda.c old mode 100755 new mode 100644 diff --git a/packages/cuda/cuda.yap b/packages/cuda/cuda.yap old mode 100755 new mode 100644 diff --git a/packages/cuda/dbio.cu b/packages/cuda/dbio.cu index 6b0fe78b4..14e88a747 100644 --- a/packages/cuda/dbio.cu +++ b/packages/cuda/dbio.cu @@ -27,8 +27,8 @@ void datalogWrite(int query, vector::iterator rul_str, vector::iterator rul_str, vector::iterator rul_str, tipo = res_rows * cols1 * sizeof(int); hres = (int *)malloc(tipo); - cudaMemcpy(hres, dop1, tipo, cudaMemcpyDeviceToHost); - cudaFree(dop1); + hipMemcpy(hres, dop1, tipo, hipMemcpyDeviceToHost); + hipFree(dop1); w = z + 1; strtok(qposr->rulename, "_"); @@ -353,8 +353,8 @@ void postgresWrite(int *inpquery, int ninpf, vector::iterator rul_str, res_rows = abs(res_rows); tipo = res_rows * cols1 * sizeof(int); hres = (int *)malloc(tipo); - cudaMemcpy(hres, dop1, tipo, cudaMemcpyDeviceToHost); - cudaFree(dop1); + hipMemcpy(hres, dop1, tipo, hipMemcpyDeviceToHost); + hipFree(dop1); char file[] = "/dev/shm/buffer.csv"; FILE *fp; @@ -554,7 +554,7 @@ void mysqlWrite(vector::iterator rul_str, vector::iterator f sign = tmpfact.predname; tipo = res_rows * cols1 * sizeof(int); hres = (int *)malloc(tipo); - cudaMemcpy(hres, dop1, tipo, cudaMemcpyDeviceToHost); + hipMemcpy(hres, dop1, tipo, hipMemcpyDeviceToHost); if(sign[0] == 'f' && sign[1] >= '0' && sign[1] <= '9') sumar(tmpfact.name, dop1, cols1, res_rows); } diff --git a/packages/cuda/joincpu.cpp b/packages/cuda/joincpu.cpp old mode 100755 new mode 100644 index 391b0d9f8..ead42edb5 --- a/packages/cuda/joincpu.cpp +++ b/packages/cuda/joincpu.cpp @@ -324,11 +324,11 @@ int joincpu(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list 0 || nsj1 > 0) @@ -359,16 +359,16 @@ int joincpu(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list 0 || nsj2 > 0) @@ -381,16 +381,16 @@ int joincpu(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list L; extern "C" int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, int *inpquery, int **result, char *names, int finalDR) { - cudaSetDevice(0); + hipSetDevice(0); vector rules; int x; @@ -1029,11 +1029,11 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, vector::iterator qposr; #if TIMER - cudaEvent_t start, stop; + hipEvent_t start, stop; float time; - cudaEventCreate(&start); - cudaEventCreate(&stop); - cudaEventRecord(start, 0); + hipEventCreate(&start); + hipEventCreate(&stop); + hipEventRecord(start, 0); #endif while(reglas.size()) /*Here's the main loop*/ @@ -1084,7 +1084,7 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, { num_refs = rows1 * cols1 * sizeof(int); reservar(&res, num_refs); - cudaMemcpyAsync(res, dop1, num_refs, cudaMemcpyDeviceToDevice); + hipMemcpyAsync(res, dop1, num_refs, hipMemcpyDeviceToDevice); registrar(rul_act->name, cols1, res, rows1, itr, 1); genflag = 1; rul_act->gen_ant = rul_act->gen_act; @@ -1251,10 +1251,10 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, if(x == num_refs) { #ifdef TIMER - cudaEvent_t start2, stop2; - cudaEventCreate(&start2); - cudaEventCreate(&stop2); - cudaEventRecord(start2, 0); + hipEvent_t start2, stop2; + hipEventCreate(&start2); + hipEventCreate(&stop2); + hipEventRecord(start2, 0); #endif //cout << rul_act->name << " res_rows = " << res_rows << endl; @@ -1263,11 +1263,11 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, res_rows = unir(res, res_rows, rul_act->num_columns, &res, 0); #ifdef TIMER - cudaEventRecord(stop2, 0); - cudaEventSynchronize(stop2); - cudaEventElapsedTime(&time, start2, stop2); - cudaEventDestroy(start2); - cudaEventDestroy(stop2); + hipEventRecord(stop2, 0); + hipEventSynchronize(stop2); + hipEventElapsedTime(&time, start2, stop2); + hipEventDestroy(start2); + hipEventDestroy(stop2); //cout << "Union = " << time << endl; cuda_stats.union_time += time; #endif @@ -1319,16 +1319,16 @@ int Cuda_Eval(predicate **inpfacts, int ninpf, predicate **inprules, int ninpr, #endif #if TIMER - cudaEventRecord(stop, 0); - cudaEventSynchronize(stop); - cudaEventElapsedTime(&time, start, stop); + hipEventRecord(stop, 0); + hipEventSynchronize(stop); + hipEventElapsedTime(&time, start, stop); cuda_stats.total_time += time; if (time > cuda_stats.max_time) cuda_stats.max_time = time; if (time < cuda_stats.min_time || cuda_stats.calls == 1) cuda_stats.min_time = time; - cudaEventDestroy(start); - cudaEventDestroy(stop); + hipEventDestroy(start); + hipEventDestroy(stop); Cuda_Statistics(); #endif diff --git a/packages/cuda/lista.h b/packages/cuda/lista.h old mode 100755 new mode 100644 diff --git a/packages/cuda/memory.cu b/packages/cuda/memory.cu old mode 100755 new mode 100644 index 70036e07b..e57ca39b3 --- a/packages/cuda/memory.cu +++ b/packages/cuda/memory.cu @@ -144,7 +144,7 @@ void limpiar(const char s[], size_t sz) if(GPUmem.size() == 0) { - cudaMemGetInfo(&free,&total); + hipMemGetInfo(&free,&total); cerr << s << ": not enough GPU memory: have " << free << " of " << total << ", need " << sz << " bytes." << endl; exit(1); } @@ -154,11 +154,11 @@ void limpiar(const char s[], size_t sz) { temp = *ini; temp.dev_address = (int *)malloc(ini->size); - cudaMemcpyAsync(temp.dev_address, ini->dev_address, temp.size, cudaMemcpyDeviceToHost); + hipMemcpyAsync(temp.dev_address, ini->dev_address, temp.size, hipMemcpyDeviceToHost); list::iterator pos = lower_bound(CPUmem.begin(), CPUmem.end(), temp, compareiteration); CPUmem.insert(pos, temp); } - cudaFree(ini->dev_address); + hipFree(ini->dev_address); GPUmem.erase(ini); } @@ -173,19 +173,19 @@ void reservar(int **ptr, size_t size) return; } - cudaMemGetInfo(&free, &total); + hipMemGetInfo(&free, &total); while(free < size) { cout << "Se limpio memoria " << free << " " << total << endl; limpiar("not enough memory", size); - cudaMemGetInfo(&free, &total); + hipMemGetInfo(&free, &total); } - while(cudaMalloc(ptr, size) == cudaErrorMemoryAllocation) + while(hipMalloc(ptr, size) == hipErrorMemoryAllocation) limpiar("Error in memory allocation", size); if (! *ptr ) { size_t free, total; - cudaMemGetInfo( &free, &total ); + hipMemGetInfo( &free, &total ); cerr << "Could not allocate " << size << " bytes, only " << free << " avaliable from total of " << total << " !!!" << endl; cerr << "Exiting CUDA...." << endl; exit(1); @@ -277,7 +277,7 @@ int cargar(int name, int num_rows, int num_columns, int is_fact, int *address_ho } size = num_rows * num_columns * sizeof(int); reservar(&temp, size); - cudaMemcpyAsync(temp, address_host_table, size, cudaMemcpyHostToDevice); + hipMemcpyAsync(temp, address_host_table, size, hipMemcpyHostToDevice); registrar(name, num_columns, temp, num_rows, itr, 0); *ptr = temp; return num_rows; @@ -296,13 +296,13 @@ int cargar(int name, int num_rows, int num_columns, int is_fact, int *address_ho reservar(&temp, size); for(x = 0; x < numgpu; x++) { - cudaMemcpyAsync(temp + inc, temp_storage[x].dev_address, temp_storage[x].size, cudaMemcpyDeviceToDevice); + hipMemcpyAsync(temp + inc, temp_storage[x].dev_address, temp_storage[x].size, hipMemcpyDeviceToDevice); inc += temp_storage[x].size / sizeof(int); - cudaFree(temp_storage[x].dev_address); + hipFree(temp_storage[x].dev_address); } for(; x < numcpu; x++) { - cudaMemcpyAsync(temp + inc, temp_storage[x].dev_address, temp_storage[x].size, cudaMemcpyHostToDevice); + hipMemcpyAsync(temp + inc, temp_storage[x].dev_address, temp_storage[x].size, hipMemcpyHostToDevice); inc += temp_storage[x].size / sizeof(int); free(temp_storage[x].dev_address); } @@ -340,9 +340,9 @@ int cargarcpu(int name, int num_rows, int num_columns, int is_fact, int *address temp = (int *)malloc(size); for(x = 0; x < numgpu; x++) { - cudaMemcpyAsync(temp + inc, temp_storage[x].dev_address, temp_storage[x].size, cudaMemcpyDeviceToHost); + hipMemcpyAsync(temp + inc, temp_storage[x].dev_address, temp_storage[x].size, hipMemcpyDeviceToHost); inc += temp_storage[x].size / sizeof(int); - cudaFree(temp_storage[x].dev_address); + hipFree(temp_storage[x].dev_address); } for(; x < numcpu; x++) { @@ -404,7 +404,7 @@ int cargafinal(int name, int cols, int **ptr) cont = pos->rows; #ifdef TUFFY reservar(&temp, pos->size); - cudaMemcpy(temp, pos->dev_address, pos->size, cudaMemcpyHostToDevice); + hipMemcpy(temp, pos->dev_address, pos->size, hipMemcpyHostToDevice); *ptr = temp; #else *ptr = pos->dev_address; @@ -418,14 +418,14 @@ int cargafinal(int name, int cols, int **ptr) pos = gpu; while(pos != endg && pos->name == name) { - cudaMemcpy(temp, pos->dev_address, pos->size, cudaMemcpyDeviceToDevice); + hipMemcpy(temp, pos->dev_address, pos->size, hipMemcpyDeviceToDevice); temp += pos->size / sizeof(int); pos++; } pos = cpu; while(pos != endc && pos->name == name) { - cudaMemcpy(temp, pos->dev_address, pos->size, cudaMemcpyHostToDevice); + hipMemcpy(temp, pos->dev_address, pos->size, hipMemcpyHostToDevice); temp += pos->size / sizeof(int); pos++; } @@ -493,7 +493,7 @@ void clear_memory() { if(ini->isrule) { - cudaFree(ini->dev_address); + hipFree(ini->dev_address); ini = GPUmem.erase(ini); } else @@ -518,7 +518,7 @@ void clear_memory_all() fin = GPUmem.end(); while(ini != fin) { - cudaFree(ini->dev_address); + hipFree(ini->dev_address); ini++; } GPUmem.clear(); @@ -542,7 +542,7 @@ void liberar(int name) { fact = *i; GPUmem.erase(i); - cudaFree(fact.dev_address); + hipFree(fact.dev_address); } i = buscarhecho(CPUmem.begin(), CPUmem.end(), name); if(i != CPUmem.end()) @@ -566,10 +566,10 @@ void sumar(int name, int *dop1, int cols, int rows) newrows = rows + fact.rows; reservar(&res, newrows * cols * sizeof(int)); offset = fact.rows * cols; - cudaMemcpyAsync(res, fact.dev_address, offset * sizeof(int), cudaMemcpyDeviceToDevice); + hipMemcpyAsync(res, fact.dev_address, offset * sizeof(int), hipMemcpyDeviceToDevice); GPUmem.erase(i); registrar(name, cols, res, newrows, 0, 0); - cudaMemcpyAsync(res + offset, dop1, rows * cols * sizeof(int), cudaMemcpyDeviceToDevice); - cudaFree(fact.dev_address); + hipMemcpyAsync(res + offset, dop1, rows * cols * sizeof(int), hipMemcpyDeviceToDevice); + hipFree(fact.dev_address); } } diff --git a/packages/cuda/memory.h b/packages/cuda/memory.h old mode 100755 new mode 100644 diff --git a/packages/cuda/pred.h b/packages/cuda/pred.h old mode 100755 new mode 100644 diff --git a/packages/cuda/selectproyect.cu b/packages/cuda/selectproyect.cu old mode 100755 new mode 100644 index df2d51419..21912d149 --- a/packages/cuda/selectproyect.cu +++ b/packages/cuda/selectproyect.cu @@ -1,3 +1,4 @@ +#include "hip/hip_runtime.h" #include #include #include @@ -8,10 +9,10 @@ __global__ void marcar2(int *dop1, int rows, int cols, int *cons, int numc, int *res) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int x, rowact, posact; - if(threadIdx.x < numc) - shared[threadIdx.x] = cons[threadIdx.x]; + if(hipThreadIdx_x < numc) + shared[hipThreadIdx_x] = cons[hipThreadIdx_x]; __syncthreads(); if(id < rows) { @@ -30,10 +31,10 @@ we unmark any rows that do not comply with the selections*/ __global__ void marcar(int *dop1, int rows, int cols, int *cons, int numc, int *res) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int x, rowact, posact; - if(threadIdx.x < numc) - shared[threadIdx.x] = cons[threadIdx.x]; + if(hipThreadIdx_x < numc) + shared[hipThreadIdx_x] = cons[hipThreadIdx_x]; __syncthreads(); if(id < rows) { @@ -56,10 +57,10 @@ __global__ void marcar(int *dop1, int rows, int cols, int *cons, int numc, int * __global__ void samejoin(int *dop1, int rows, int cols, int *dhead, int cont, int *res) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int temp, temp2, pos, x, y; - if(threadIdx.x < cont) - shared[threadIdx.x] = dhead[threadIdx.x]; + if(hipThreadIdx_x < cont) + shared[hipThreadIdx_x] = dhead[hipThreadIdx_x]; __syncthreads(); if(id < rows) { @@ -90,10 +91,10 @@ __global__ void samejoin(int *dop1, int rows, int cols, int *dhead, int cont, in __global__ void samejoin2(int *dop1, int rows, int cols, int *dhead, int cont, int *res) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int temp, temp2, pos, x, y; - if(threadIdx.x < cont) - shared[threadIdx.x] = dhead[threadIdx.x]; + if(hipThreadIdx_x < cont) + shared[hipThreadIdx_x] = dhead[hipThreadIdx_x]; __syncthreads(); if(id < rows) { @@ -120,10 +121,10 @@ __global__ void samejoin2(int *dop1, int rows, int cols, int *dhead, int cont, i __global__ void proyectar(int *dop1, int rows, int cols, int *dhead, int hsize, int *res) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int pos, posr, x; - if(threadIdx.x < hsize) - shared[threadIdx.x] = dhead[threadIdx.x]; + if(hipThreadIdx_x < hsize) + shared[hipThreadIdx_x] = dhead[hipThreadIdx_x]; __syncthreads(); if(id < rows) { @@ -139,10 +140,10 @@ selections, selfjoins, etc.). The array 'temp' holds the result of the prefix su __global__ void llenarproyectar(int *dop1, int rows, int cols, int *temp, int *dhead, int hsize, int *res) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int pos, posr, x; - if(threadIdx.x < hsize) - shared[threadIdx.x] = dhead[threadIdx.x]; + if(hipThreadIdx_x < hsize) + shared[hipThreadIdx_x] = dhead[hipThreadIdx_x]; __syncthreads(); if(id < rows) { @@ -184,27 +185,27 @@ int selectproyect(int *dop1, int rows, int cols, int head_size, int *select, int tmplen = rows + 1; size2 = tmplen * sizeof(int); reservar(&temp, size2); - cudaMemset(temp, 0, size2); + hipMemset(temp, 0, size2); size = numselect * sizeof(int); - cudaMemcpy(dhead, select, size, cudaMemcpyHostToDevice); + hipMemcpy(dhead, select, size, hipMemcpyHostToDevice); - marcar2<<>>(dop1, rows, cols, dhead, numselect, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(marcar2), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, dhead, numselect, temp + 1); if(numselfj > 0) { size = numselfj * sizeof(int); - cudaMemcpy(dhead, selfjoin, size, cudaMemcpyHostToDevice); - samejoin<<>>(dop1, rows, cols, dhead, numselfj, temp + 1); + hipMemcpy(dhead, selfjoin, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(samejoin), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, dhead, numselfj, temp + 1); } if(numpreds > 0) { size = numpreds * sizeof(int); - cudaMemcpy(dhead, preds, size, cudaMemcpyHostToDevice); + hipMemcpy(dhead, preds, size, hipMemcpyHostToDevice); if(ANDlogic) - bpredsnormal<<>>(dop1, rows, cols, dhead, numpreds, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsnormal), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, dhead, numpreds, temp + 1); else - bpredsorlogic<<>>(dop1, rows, cols, dhead, numpreds, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsorlogic), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, dhead, numpreds, temp + 1); } res = thrust::device_pointer_cast(temp); @@ -215,10 +216,10 @@ int selectproyect(int *dop1, int rows, int cols, int head_size, int *select, int size = head_size * sizeof(int); reservar(&fres, num * size); - cudaMemcpy(dhead, project, size, cudaMemcpyHostToDevice); - llenarproyectar<<>>(dop1, rows, cols, temp, dhead, head_size, fres); - cudaFree(dhead); - cudaFree(temp); + hipMemcpy(dhead, project, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(llenarproyectar), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, temp, dhead, head_size, fres); + hipFree(dhead); + hipFree(temp); *ret = fres; return num; } @@ -229,19 +230,19 @@ int selectproyect(int *dop1, int rows, int cols, int head_size, int *select, int tmplen = rows + 1; size2 = tmplen * sizeof(int); reservar(&temp, size2); - cudaMemset(temp, 0, size2); + hipMemset(temp, 0, size2); size = numselfj * sizeof(int); - cudaMemcpy(dhead, selfjoin, size, cudaMemcpyHostToDevice); - samejoin2<<>>(dop1, rows, cols, dhead, numselfj, temp + 1); + hipMemcpy(dhead, selfjoin, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(samejoin2), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, dhead, numselfj, temp + 1); if(numpreds > 0) { size = numpreds * sizeof(int); - cudaMemcpy(dhead, preds, size, cudaMemcpyHostToDevice); + hipMemcpy(dhead, preds, size, hipMemcpyHostToDevice); if(ANDlogic) - bpredsnormal<<>>(dop1, rows, cols, dhead, numpreds, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsnormal), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, dhead, numpreds, temp + 1); else - bpredsorlogic<<>>(dop1, rows, cols, dhead, numpreds, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsorlogic), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, dhead, numpreds, temp + 1); } @@ -253,10 +254,10 @@ int selectproyect(int *dop1, int rows, int cols, int head_size, int *select, int size = head_size * sizeof(int); reservar(&fres, num * size); - cudaMemcpy(dhead, project, size, cudaMemcpyHostToDevice); - llenarproyectar<<>>(dop1, rows, cols, temp, dhead, head_size, fres); - cudaFree(dhead); - cudaFree(temp); + hipMemcpy(dhead, project, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(llenarproyectar), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, temp, dhead, head_size, fres); + hipFree(dhead); + hipFree(temp); *ret = fres; return num; } @@ -267,14 +268,14 @@ int selectproyect(int *dop1, int rows, int cols, int head_size, int *select, int tmplen = rows + 1; size2 = tmplen * sizeof(int); reservar(&temp, size2); - cudaMemset(temp, 0, size2); + hipMemset(temp, 0, size2); size = numpreds * sizeof(int); - cudaMemcpy(dhead, preds, size, cudaMemcpyHostToDevice); + hipMemcpy(dhead, preds, size, hipMemcpyHostToDevice); if(ANDlogic) - bpredsnormal2<<>>(dop1, rows, cols, dhead, numpreds, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsnormal2), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, dhead, numpreds, temp + 1); else - bpredsorlogic2<<>>(dop1, rows, cols, dhead, numpreds, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsorlogic2), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, dhead, numpreds, temp + 1); res = thrust::device_pointer_cast(temp); thrust::inclusive_scan(res + 1, res + tmplen, res + 1); num = res[rows]; @@ -284,10 +285,10 @@ int selectproyect(int *dop1, int rows, int cols, int head_size, int *select, int size = head_size * sizeof(int); reservar(&fres, num * size); - cudaMemcpy(dhead, project, size, cudaMemcpyHostToDevice); - llenarproyectar<<>>(dop1, rows, cols, temp, dhead, head_size, fres); - cudaFree(dhead); - cudaFree(temp); + hipMemcpy(dhead, project, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(llenarproyectar), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, temp, dhead, head_size, fres); + hipFree(dhead); + hipFree(temp); *ret = fres; return num; } @@ -295,9 +296,9 @@ int selectproyect(int *dop1, int rows, int cols, int head_size, int *select, int { size = head_size * sizeof(int); reservar(&fres, rows * size); - cudaMemcpy(dhead, project, size, cudaMemcpyHostToDevice); - proyectar<<>>(dop1, rows, cols, dhead, head_size, fres); - cudaFree(dhead); + hipMemcpy(dhead, project, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(proyectar), dim3(blockllen), dim3(numthreads), size, 0, dop1, rows, cols, dhead, head_size, fres); + hipFree(dhead); *ret = fres; return rows; } diff --git a/packages/cuda/selectproyectcpu.cpp b/packages/cuda/selectproyectcpu.cpp old mode 100755 new mode 100644 diff --git a/packages/cuda/treeb.cu b/packages/cuda/treeb.cu old mode 100755 new mode 100644 index 68bd55a7b..8272c41b7 --- a/packages/cuda/treeb.cu +++ b/packages/cuda/treeb.cu @@ -1,3 +1,4 @@ +#include "hip/hip_runtime.h" #include #include #include @@ -160,11 +161,11 @@ __device__ int firstMatchingKeyInDataNode2(Record records[], IKeyType key) __global__ void gCreateIndex(IDataNode data[], IDirectoryNode dir[], int dirSize, int tree_size, int bottom_start, int nNodesPerBlock) { - int startIdx = blockIdx.x * nNodesPerBlock; + int startIdx = hipBlockIdx_x * nNodesPerBlock; int endIdx = startIdx + nNodesPerBlock; if(endIdx > dirSize) endIdx = dirSize; - int keyIdx = threadIdx.x; + int keyIdx = hipThreadIdx_x; // Proceed only when in internal nodes for(int nodeIdx = startIdx; nodeIdx < endIdx; nodeIdx++) @@ -191,11 +192,11 @@ __global__ void gSearchTree(IDataNode* data, int nDataNodes, IDirectoryNode* dir { // Bringing the root node (visited by every tuple) to the faster shared memory __shared__ IKeyType RootNodeKeys[TREE_NODE_SIZE]; - RootNodeKeys[threadIdx.x] = dir->keys[threadIdx.x]; + RootNodeKeys[hipThreadIdx_x] = dir->keys[hipThreadIdx_x]; __syncthreads(); - int OverallThreadIdx = blockIdx.x * THRD_PER_BLCK_search + threadIdx.x; + int OverallThreadIdx = hipBlockIdx_x * THRD_PER_BLCK_search + hipThreadIdx_x; for(int keyIdx = OverallThreadIdx; keyIdx < nSearchKeys; keyIdx += THRD_PER_GRID_search) { @@ -219,7 +220,7 @@ __global__ void gSearchTree(IDataNode* data, int nDataNodes, IDirectoryNode* dir /*Counts the number of times a row in 'S' is to be joined to a row in 'R'.*/ __global__ void gIndexJoin(int *R, int *S, int g_locations[], int sLen, int g_ResNums[]) { - int s_cur = blockIdx.x * blockDim.x + threadIdx.x; + int s_cur = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; if(s_cur < sLen) { @@ -246,11 +247,11 @@ in 'g_locations' those rows that have equal values in the checked columns.*/ __global__ void gIndexMultiJoinNegative(int *R, int *S, int g_locations[], int rLen, int *p1, int *p2, int of1, int of2, int *mloc, int *sloc, int *muljoin, int wj) { extern __shared__ int shared[]; - int r_cur = blockIdx.x * blockDim.x + threadIdx.x; + int r_cur = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int posr, poss, x; - if(threadIdx.x < wj) - shared[threadIdx.x] = muljoin[threadIdx.x]; + if(hipThreadIdx_x < wj) + shared[hipThreadIdx_x] = muljoin[hipThreadIdx_x]; __syncthreads(); if(r_cur < rLen) @@ -287,11 +288,11 @@ times a row in 'S' is to be joined to its corresponding row in 'R', storing the __global__ void gIndexMultiJoin(int *R, int *S, int g_locations[], int sLen, int g_ResNums[], int *p1, int *p2, int of1, int of2, int *mloc, int *sloc, int *muljoin, int wj) { extern __shared__ int shared[]; - int s_cur = blockIdx.x * blockDim.x + threadIdx.x; + int s_cur = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int posr, poss, x; - if(threadIdx.x < wj) - shared[threadIdx.x] = muljoin[threadIdx.x]; + if(hipThreadIdx_x < wj) + shared[hipThreadIdx_x] = muljoin[hipThreadIdx_x]; __syncthreads(); if(s_cur < sLen) @@ -330,10 +331,10 @@ __global__ void multiJoinWithWrite(int g_locations[], int sLen, int g_PrefixSums { extern __shared__ int shared[]; int *extjoins = &shared[lenrul]; - int s_cur = blockIdx.x * blockDim.x + threadIdx.x; + int s_cur = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; - if(threadIdx.x < (lenrul + wj)) - shared[threadIdx.x] = rule[threadIdx.x]; + if(hipThreadIdx_x < (lenrul + wj)) + shared[hipThreadIdx_x] = rule[hipThreadIdx_x]; __syncthreads(); if(s_cur < sLen) @@ -382,10 +383,10 @@ __global__ void multiJoinWithWrite2(int g_locations[], int sLen, int g_PrefixSum { extern __shared__ int shared[]; int *extjoins = &shared[cols]; - int s_cur = blockIdx.x * blockDim.x + threadIdx.x; + int s_cur = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; - if(threadIdx.x < (cols + wj)) - shared[threadIdx.x] = rule[threadIdx.x]; + if(hipThreadIdx_x < (cols + wj)) + shared[hipThreadIdx_x] = rule[hipThreadIdx_x]; __syncthreads(); if(s_cur < sLen) @@ -432,11 +433,11 @@ predicate are projected.*/ __global__ void gJoinWithWriteNegative(int g_locations[], int rLen, int g_joinResultBuffers[], int *p1, int of1, int *rule, int halfrul, int *mloc) { extern __shared__ int shared[]; - int r_cur = blockIdx.x * blockDim.x + threadIdx.x; + int r_cur = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int posr; - if(threadIdx.x < halfrul) - shared[threadIdx.x] = rule[threadIdx.x]; + if(hipThreadIdx_x < halfrul) + shared[hipThreadIdx_x] = rule[hipThreadIdx_x]; __syncthreads(); if(r_cur < rLen) @@ -461,11 +462,11 @@ predicate are projected.*/ __global__ void gJoinWithWriteNegative2(int g_locations[], int rLen, int g_joinResultBuffers[], int *p1, int of1, int *rule, int cols, int *mloc) { extern __shared__ int shared[]; - int r_cur = blockIdx.x * blockDim.x + threadIdx.x; + int r_cur = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int posr; - if(threadIdx.x < cols) - shared[threadIdx.x] = rule[threadIdx.x]; + if(hipThreadIdx_x < cols) + shared[hipThreadIdx_x] = rule[hipThreadIdx_x]; __syncthreads(); if(r_cur < rLen) @@ -489,10 +490,10 @@ __global__ void gJoinWithWriteNegative2(int g_locations[], int rLen, int g_joinR __global__ void gJoinWithWrite(int g_locations[], int sLen, int g_PrefixSums[], int g_joinResultBuffers[], int *p1, int *p2, int of1, int of2, int *rule, int halfrul, int lenrul, int *mloc, int *sloc) { extern __shared__ int shared[]; - int s_cur = blockIdx.x * blockDim.x + threadIdx.x; + int s_cur = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; - if(threadIdx.x < lenrul) - shared[threadIdx.x] = rule[threadIdx.x]; + if(hipThreadIdx_x < lenrul) + shared[hipThreadIdx_x] = rule[hipThreadIdx_x]; __syncthreads(); if(s_cur < sLen) @@ -525,10 +526,10 @@ projection, which is performed based on the variables in the head of the rule.*/ __global__ void gJoinWithWrite2(int g_locations[], int sLen, int g_PrefixSums[], int g_joinResultBuffers[], int *p1, int *p2, int of1, int of2, int *rule, int cols, int *mloc, int *sloc) { extern __shared__ int shared[]; - int s_cur = blockIdx.x * blockDim.x + threadIdx.x; + int s_cur = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; - if(threadIdx.x < cols) - shared[threadIdx.x] = rule[threadIdx.x]; + if(hipThreadIdx_x < cols) + shared[hipThreadIdx_x] = rule[hipThreadIdx_x]; __syncthreads(); if(s_cur < sLen) @@ -563,7 +564,7 @@ __global__ void gJoinWithWrite2(int g_locations[], int sLen, int g_PrefixSums[], /*Load part of column 'wj' of 'p' in 'R'. Which values are loaded is defined by the prefix sum results in 'pos'.*/ __global__ void llenar(int *p, int *R, int len, int of, int wj, int *pos, int *ids) { - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int cond; if(id < len) { @@ -579,7 +580,7 @@ __global__ void llenar(int *p, int *R, int len, int of, int wj, int *pos, int *i /*Load an entire column from 'p' into 'R'.*/ __global__ void llenarnosel(int *p, int *R, int len, int of, int wj) { - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; if(id < len) R[id] = p[id * of + wj]; } @@ -587,10 +588,10 @@ __global__ void llenarnosel(int *p, int *R, int len, int of, int wj) __global__ void projectfinal(int *res, int rows, int cols, int *rule, int *out) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; - if(threadIdx.x < cols) - shared[threadIdx.x] = rule[threadIdx.x]; + if(hipThreadIdx_x < cols) + shared[hipThreadIdx_x] = rule[hipThreadIdx_x]; __syncthreads(); if(id < rows) @@ -614,26 +615,26 @@ void project(int *res, int resrows, int numcols1, int numcols2, int *proj, int * int *pt = (int *)malloc(sizepro); for(z = 0; z < numcols2; z++) pt[z] = proj[z] - 1; - cudaMemcpy(dcons, pt, sizepro, cudaMemcpyHostToDevice); - //cudaDeviceSynchronize(); //Small cudaMemcpys are asynchronous, uncomment this line if the pointer is being liberated before it is copied. + hipMemcpy(dcons, pt, sizepro, hipMemcpyHostToDevice); + //hipDeviceSynchronize(); //Small cudaMemcpys are asynchronous, uncomment this line if the pointer is being liberated before it is copied. free(pt); } else - cudaMemcpy(dcons, proj, sizepro, cudaMemcpyHostToDevice); + hipMemcpy(dcons, proj, sizepro, hipMemcpyHostToDevice); reservar(&d_Rout, resrows * sizepro); - projectfinal<<>>(res, resrows, numcols1, dcons, d_Rout); - cudaFree(dcons); - cudaFree(*ret); + hipLaunchKernel(HIP_KERNEL_NAME(projectfinal), dim3(blockllen), dim3(numthreads), sizepro, 0, res, resrows, numcols1, dcons, d_Rout); + hipFree(dcons); + hipFree(*ret); *ret = d_Rout; } __global__ void projectadd(int *dop1, int *dop2, int rows1, int rows2, int cols1, int cols2, int *dhead, int hsize, int *res) { extern __shared__ int shared[]; - int id = blockIdx.x * blockDim.x + threadIdx.x; + int id = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int pos2, posr, x, y, cond; - if(threadIdx.x < hsize) - shared[threadIdx.x] = dhead[threadIdx.x]; + if(hipThreadIdx_x < hsize) + shared[hipThreadIdx_x] = dhead[hipThreadIdx_x]; __syncthreads(); if(id < rows2) { @@ -662,10 +663,10 @@ void juntar(int *dop1, int *dop2, int rows1, int rows2, int cols1, int cols2, in int blockllen = rows2 / numthreads + 1; sizepro = pcols * sizeof(int); reservar(&dcons, sizepro); - cudaMemcpy(dcons, proj, sizepro, cudaMemcpyHostToDevice); + hipMemcpy(dcons, proj, sizepro, hipMemcpyHostToDevice); reservar(&d_Rout, rows1 * rows2 * sizepro); - projectadd<<>>(dop1, dop2, rows1, rows2, cols1, cols2, dcons, pcols, d_Rout); - cudaFree(dcons); + hipLaunchKernel(HIP_KERNEL_NAME(projectadd), dim3(blockllen), dim3(numthreads), sizepro, 0, dop1, dop2, rows1, rows2, cols1, cols2, dcons, pcols, d_Rout); + hipFree(dcons); *ret = d_Rout; } @@ -743,51 +744,51 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: #ifdef TIMER //cout << "INICIO" << endl; - cudaEvent_t start, stop; + hipEvent_t start, stop; float time; - cudaEventCreate(&start); - cudaEventCreate(&stop); - cudaEventRecord(start, 0); + hipEventCreate(&start); + hipEventCreate(&stop); + hipEventRecord(start, 0); #endif if(npred2.x > 0 || npred2.y > 0 || nsel2 > 0 || nsj2 > 0) { newLen = sLen + 1; - cudaMemsetAsync(temp, 0, newLen * sizeof(int)); + hipMemsetAsync(temp, 0, newLen * sizeof(int)); } if(npred2.x > 0 || npred2.y > 0) { size = npred2tot * sizeof(int); - cudaMemcpy(dcons, pred2, size, cudaMemcpyHostToDevice); + hipMemcpy(dcons, pred2, size, hipMemcpyHostToDevice); if(npred2.y > 0) /*Fix case when a(X,Y),b(Y,Z),Z > Y*/ { reservar(&temp2, sizet2); - cudaMemsetAsync(temp2, 0, newLen * sizeof(int)); + hipMemsetAsync(temp2, 0, newLen * sizeof(int)); //res = thrust::device_pointer_cast(temp2); - bpreds<<>>(p1, p2, sLen, of1, of2, dcons, npred2tot, npred2.x, temp + 1, temp2 + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpreds), dim3(blockllen), dim3(numthreads), size, 0, p1, p2, sLen, of1, of2, dcons, npred2tot, npred2.x, temp + 1, temp2 + 1); } else { if(negative) - bpreds<<>>(p1, p2, sLen, of1, of2, dcons, npred2tot, npred2.x, temp + 1, NULL); + hipLaunchKernel(HIP_KERNEL_NAME(bpreds), dim3(blockllen), dim3(numthreads), size, 0, p1, p2, sLen, of1, of2, dcons, npred2tot, npred2.x, temp + 1, NULL); else - bpredsOR<<>>(p1, p2, sLen, of1, of2, dcons, npred2tot, npred2.x, temp + 1, NULL); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsOR), dim3(blockllen), dim3(numthreads), size, 0, p1, p2, sLen, of1, of2, dcons, npred2tot, npred2.x, temp + 1, NULL); } if(nsel2 > 0) { size = nsel2 * sizeof(int); - cudaMemcpy(dcons, sel2, size, cudaMemcpyHostToDevice); - marcar<<>>(p2, sLen, of2, dcons, nsel2, temp + 1); + hipMemcpy(dcons, sel2, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(marcar), dim3(blockllen), dim3(numthreads), size, 0, p2, sLen, of2, dcons, nsel2, temp + 1); } if(nsj2 > 0) { size = nsj2 * sizeof(int); - cudaMemcpy(dcons, sjoin2, size, cudaMemcpyHostToDevice); - samejoin<<>>(p2, sLen, of2, dcons, nsj2, temp + 1); + hipMemcpy(dcons, sjoin2, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(samejoin), dim3(blockllen), dim3(numthreads), size, 0, p2, sLen, of2, dcons, nsj2, temp + 1); } } else @@ -795,14 +796,14 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: if(nsel2 > 0) { size = nsel2 * sizeof(int); - cudaMemcpy(dcons, sel2, size, cudaMemcpyHostToDevice); - marcar2<<>>(p2, sLen, of2, dcons, nsel2, temp + 1); + hipMemcpy(dcons, sel2, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(marcar2), dim3(blockllen), dim3(numthreads), size, 0, p2, sLen, of2, dcons, nsel2, temp + 1); if(nsj2 > 0) { size = nsj2 * sizeof(int); - cudaMemcpy(dcons, sjoin2, size, cudaMemcpyHostToDevice); - samejoin<<>>(p2, sLen, of2, dcons, nsj2, temp + 1); + hipMemcpy(dcons, sjoin2, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(samejoin), dim3(blockllen), dim3(numthreads), size, 0, p2, sLen, of2, dcons, nsj2, temp + 1); } } else @@ -810,15 +811,15 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: if(nsj2 > 0) { size = nsj2 * sizeof(int); - cudaMemcpy(dcons, sjoin2, size, cudaMemcpyHostToDevice); - samejoin2<<>>(p2, sLen, of2, dcons, nsj2, temp + 1); + hipMemcpy(dcons, sjoin2, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(samejoin2), dim3(blockllen), dim3(numthreads), size, 0, p2, sLen, of2, dcons, nsj2, temp + 1); } else { sizem32S = m32sLen * sizeof(int); reservar(&d_S, sizem32S); - cudaMemsetAsync(d_S + sLen, 0x7f, extraspaceS * sizeof(int)); - llenarnosel<<>>(p2, d_S, sLen, of2, wherej[1]); + hipMemsetAsync(d_S + sLen, 0x7f, extraspaceS * sizeof(int)); + hipLaunchKernel(HIP_KERNEL_NAME(llenarnosel), dim3(blockllen), dim3(numthreads), 0, 0, p2, d_S, sLen, of2, wherej[1]); } } } @@ -842,8 +843,8 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: if(newLen == 0) // && !negative) ARREGLAR { - cudaFree(temp); - cudaFree(dcons); + hipFree(temp); + hipFree(dcons); return 0; } @@ -854,24 +855,24 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: reservar(&d_S, sizem32S); reservar(&posS, sizem32S); - cudaMemsetAsync(d_S + newLen, 0x7f, sizextra); - cudaMemsetAsync(posS + newLen, 0x7f, sizextra); - llenar<<>>(p2, d_S, sLen, of2, wherej[1], temp, posS); + hipMemsetAsync(d_S + newLen, 0x7f, sizextra); + hipMemsetAsync(posS + newLen, 0x7f, sizextra); + hipLaunchKernel(HIP_KERNEL_NAME(llenar), dim3(blockllen), dim3(numthreads), 0, 0, p2, d_S, sLen, of2, wherej[1], temp, posS); sLen = newLen; } #ifdef TIMER - cudaEventRecord(stop, 0); - cudaEventSynchronize(stop); - cudaEventElapsedTime(&time, start, stop); + hipEventRecord(stop, 0); + hipEventSynchronize(stop); + hipEventElapsedTime(&time, start, stop); //cout << "Select1 = " << time << endl; cuda_stats.select1_time += time; - cudaEventDestroy(start); - cudaEventDestroy(stop); - cudaEventCreate(&start); - cudaEventCreate(&stop); - cudaEventRecord(start, 0); + hipEventDestroy(start); + hipEventDestroy(stop); + hipEventCreate(&start); + hipEventCreate(&stop); + hipEventRecord(start, 0); #endif blockllen = rLen / numthreads + 1; @@ -880,30 +881,30 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: { if(temp2 != NULL) { - cudaFree(temp); + hipFree(temp); temp = temp2; res = thrust::device_pointer_cast(temp); newLen = rLen + 1; if(nsel1 > 0) { size = nsel1 * sizeof(int); - cudaMemcpy(dcons, sel1, size, cudaMemcpyHostToDevice); - marcar<<>>(p1, rLen, of1, dcons, nsel1, temp + 1); + hipMemcpy(dcons, sel1, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(marcar), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, nsel1, temp + 1); } if(nsj1 > 0) { size = nsj1 * sizeof(int); - cudaMemcpy(dcons, sjoin1, size, cudaMemcpyHostToDevice); - samejoin<<>>(p1, rLen, of1, dcons, nsj1, temp + 1); + hipMemcpy(dcons, sjoin1, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(samejoin), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, nsj1, temp + 1); } if(npred1.x > 0) { size = npred1.x * sizeof(int); - cudaMemcpy(dcons, pred1, size, cudaMemcpyHostToDevice); + hipMemcpy(dcons, pred1, size, hipMemcpyHostToDevice); if(ANDlogic) - bpredsnormal<<>>(p1, rLen, of1, dcons, npred1.x, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsnormal), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, npred1.x, temp + 1); else - bpredsorlogic<<>>(p1, rLen, of1, dcons, npred1.x, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsorlogic), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, npred1.x, temp + 1); } } else @@ -911,30 +912,30 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: if(npred1.x > 0 || nsel1 > 0 || nsj1 > 0) { newLen = rLen + 1; - cudaMemsetAsync(temp, 0, newLen * sizeof(int)); + hipMemsetAsync(temp, 0, newLen * sizeof(int)); } if(nsel1 > 0) { size = nsel1 * sizeof(int); - cudaMemcpy(dcons, sel1, size, cudaMemcpyHostToDevice); - marcar2<<>>(p1, rLen, of1, dcons, nsel1, temp + 1); + hipMemcpy(dcons, sel1, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(marcar2), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, nsel1, temp + 1); if(nsj1 > 0) { size = nsj1 * sizeof(int); - cudaMemcpy(dcons, sjoin1, size, cudaMemcpyHostToDevice); - samejoin<<>>(p1, rLen, of1, dcons, nsj1, temp + 1); + hipMemcpy(dcons, sjoin1, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(samejoin), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, nsj1, temp + 1); } if(npred1.x > 0) { size = npred1.x * sizeof(int); - cudaMemcpy(dcons, pred1, size, cudaMemcpyHostToDevice); + hipMemcpy(dcons, pred1, size, hipMemcpyHostToDevice); if(ANDlogic) - bpredsnormal<<>>(p1, rLen, of1, dcons, npred1.x, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsnormal), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, npred1.x, temp + 1); else - bpredsorlogic<<>>(p1, rLen, of1, dcons, npred1.x, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsorlogic), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, npred1.x, temp + 1); } } else @@ -942,17 +943,17 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: if(nsj1 > 0) { size = nsj1 * sizeof(int); - cudaMemcpy(dcons, sjoin1, size, cudaMemcpyHostToDevice); - samejoin2<<>>(p1, rLen, of1, dcons, nsj1, temp + 1); + hipMemcpy(dcons, sjoin1, size, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(samejoin2), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, nsj1, temp + 1); if(npred1.x > 0) { size = npred1.x * sizeof(int); - cudaMemcpy(dcons, pred1, size, cudaMemcpyHostToDevice); + hipMemcpy(dcons, pred1, size, hipMemcpyHostToDevice); if(ANDlogic) - bpredsnormal<<>>(p1, rLen, of1, dcons, npred1.x, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsnormal), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, npred1.x, temp + 1); else - bpredsorlogic<<>>(p1, rLen, of1, dcons, npred1.x, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsorlogic), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, npred1.x, temp + 1); } } else @@ -960,11 +961,11 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: if(npred1.x > 0) { size = npred1.x * sizeof(int); - cudaMemcpy(dcons, pred1, size, cudaMemcpyHostToDevice); + hipMemcpy(dcons, pred1, size, hipMemcpyHostToDevice); if(ANDlogic) - bpredsnormal2<<>>(p1, rLen, of1, dcons, npred1.x, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsnormal2), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, npred1.x, temp + 1); else - bpredsorlogic2<<>>(p1, rLen, of1, dcons, npred1.x, temp + 1); + hipLaunchKernel(HIP_KERNEL_NAME(bpredsorlogic2), dim3(blockllen), dim3(numthreads), size, 0, p1, rLen, of1, dcons, npred1.x, temp + 1); } } } @@ -976,11 +977,11 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: newLen = res[rLen]; if(newLen == 0) { - cudaFree(temp); - cudaFree(dcons); - cudaFree(d_S); + hipFree(temp); + hipFree(dcons); + hipFree(d_S); if(posS != NULL) - cudaFree(posS); + hipFree(posS); return 0; } @@ -991,41 +992,41 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: reservar(&d_R, sizem32); reservar(&posR, sizem32); - cudaMemsetAsync(d_R + newLen, 0x7f, sizextra); - cudaMemsetAsync(posR + newLen, 0x7f, sizextra); - llenar<<>>(p1, d_R, rLen, of1, wherej[0], temp, posR); + hipMemsetAsync(d_R + newLen, 0x7f, sizextra); + hipMemsetAsync(posR + newLen, 0x7f, sizextra); + hipLaunchKernel(HIP_KERNEL_NAME(llenar), dim3(blockllen), dim3(numthreads), 0, 0, p1, d_R, rLen, of1, wherej[0], temp, posR); rLen = newLen; } else { sizem32 = m32rLen * sizeof(int); reservar(&d_R, sizem32); - cudaMemsetAsync(d_R + rLen, 0x7f, extraspace * sizeof(int)); - llenarnosel<<>>(p1, d_R, rLen, of1, wherej[0]); + hipMemsetAsync(d_R + rLen, 0x7f, extraspace * sizeof(int)); + hipLaunchKernel(HIP_KERNEL_NAME(llenarnosel), dim3(blockllen), dim3(numthreads), 0, 0, p1, d_R, rLen, of1, wherej[0]); } } else { sizem32 = m32rLen * sizeof(int); reservar(&d_R, sizem32); - cudaMemsetAsync(d_R + rLen, 0x7f, extraspace * sizeof(int)); - llenarnosel<<>>(p1, d_R, rLen, of1, wherej[0]); + hipMemsetAsync(d_R + rLen, 0x7f, extraspace * sizeof(int)); + hipLaunchKernel(HIP_KERNEL_NAME(llenarnosel), dim3(blockllen), dim3(numthreads), 0, 0, p1, d_R, rLen, of1, wherej[0]); } #ifdef TIMER - cudaEventRecord(stop, 0); - cudaEventSynchronize(stop); - cudaEventElapsedTime(&time, start, stop); + hipEventRecord(stop, 0); + hipEventSynchronize(stop); + hipEventElapsedTime(&time, start, stop); //cout << "Select2 = " << time << endl; cuda_stats.select2_time += time; #endif #ifdef TIMER - cudaEventDestroy(start); - cudaEventDestroy(stop); - cudaEventCreate(&start); - cudaEventCreate(&stop); - cudaEventRecord(start, 0); + hipEventDestroy(start); + hipEventDestroy(stop); + hipEventCreate(&start); + hipEventCreate(&stop); + hipEventRecord(start, 0); #endif thrust::device_ptr dvp1; @@ -1084,17 +1085,17 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: } #ifdef TIMER - cudaEventRecord(stop, 0); - cudaEventSynchronize(stop); - cudaEventElapsedTime(&time, start, stop); + hipEventRecord(stop, 0); + hipEventSynchronize(stop); + hipEventElapsedTime(&time, start, stop); //cout << "Sort = " << time << endl; cuda_stats.sort_time += time; - cudaEventDestroy(start); - cudaEventDestroy(stop); - cudaEventCreate(&start); - cudaEventCreate(&stop); - cudaEventRecord(start, 0); + hipEventDestroy(start); + hipEventDestroy(stop); + hipEventCreate(&start); + hipEventCreate(&stop); + hipEventRecord(start, 0); #endif IDataNode* d_data; @@ -1123,7 +1124,7 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: dim3 Dbc(THRD_PER_BLCK_create, 1, 1); dim3 Dgc(BLCK_PER_GRID_create, 1, 1); - gCreateIndex <<>> (d_data, d_dir, nDirNodes, tree_size, bottom_start, nNodesPerBlock); + hipLaunchKernel(HIP_KERNEL_NAME(gCreateIndex), dim3(Dgc), dim3(Dbc), 0, 0, d_data, d_dir, nDirNodes, tree_size, bottom_start, nNodesPerBlock); int *d_locations; int memSizeR; @@ -1132,7 +1133,7 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: { memSizeR = (rLen + 1) * sizeof(int); reservar(&d_locations, memSizeR); - cudaMemsetAsync(d_locations, 0, sizeof(int)); + hipMemsetAsync(d_locations, 0, sizeof(int)); nSearchKeys = rLen; } else @@ -1146,13 +1147,13 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: unsigned int nKeysPerThread = uintCeilingDiv(nSearchKeys, THRD_PER_GRID_search); if(negative) { - gSearchTree <<>> (d_data, nDataNodes, d_dir, nDirNodes, lvlDir, d_R, d_locations + 1, nSearchKeys, nKeysPerThread, tree_size, bottom_start); - cudaMemsetAsync(temp, 0, memSizeR); + hipLaunchKernel(HIP_KERNEL_NAME(gSearchTree), dim3(Dgs), dim3(Dbs), 0, 0, d_data, nDataNodes, d_dir, nDirNodes, lvlDir, d_R, d_locations + 1, nSearchKeys, nKeysPerThread, tree_size, bottom_start); + hipMemsetAsync(temp, 0, memSizeR); } else { - gSearchTree <<>> (d_data, nDataNodes, d_dir, nDirNodes, lvlDir, d_S, d_locations, nSearchKeys, nKeysPerThread, tree_size, bottom_start); - cudaMemsetAsync(temp, 0, memSizeS); + hipLaunchKernel(HIP_KERNEL_NAME(gSearchTree), dim3(Dgs), dim3(Dbs), 0, 0, d_data, nDataNodes, d_dir, nDirNodes, lvlDir, d_S, d_locations, nSearchKeys, nKeysPerThread, tree_size, bottom_start); + hipMemsetAsync(temp, 0, memSizeS); } int muljoin = 0, muljoinsize = 0, sum; @@ -1165,8 +1166,8 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: { muljoin = numj - 2; muljoinsize = muljoin * sizeof(int); - cudaMemcpy(dcons, wherej + 2, muljoinsize, cudaMemcpyHostToDevice); - gIndexMultiJoinNegative<<>> (d_R, d_S, d_locations + 1, rLen, p1, p2, of1, of2, posR, posS, dcons, muljoin); + hipMemcpy(dcons, wherej + 2, muljoinsize, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(gIndexMultiJoinNegative), dim3(blockllen), dim3(numthreads), muljoinsize, 0, d_R, d_S, d_locations + 1, rLen, p1, p2, of1, of2, posR, posS, dcons, muljoin); } res = thrust::device_pointer_cast(d_locations); @@ -1177,21 +1178,21 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: if(pos == (rule->num_rows - 3)) { sizepro = rule->num_columns * sizeof(int); - cudaMemcpy(dcons, proj, sizepro, cudaMemcpyHostToDevice); + hipMemcpy(dcons, proj, sizepro, hipMemcpyHostToDevice); resSize = sum * sizepro; reservar(&d_Rout, resSize); - gJoinWithWriteNegative2<<>> (d_locations, rLen, d_Rout, p1, of1, dcons, rule->num_columns, posR); + hipLaunchKernel(HIP_KERNEL_NAME(gJoinWithWriteNegative2), dim3(blockllen), dim3(numthreads), sizepro, 0, d_locations, rLen, d_Rout, p1, of1, dcons, rule->num_columns, posR); } else { sizepro = projp.x * sizeof(int); - cudaMemcpy(dcons, proj, sizepro, cudaMemcpyHostToDevice); + hipMemcpy(dcons, proj, sizepro, hipMemcpyHostToDevice); resSize = sum * sizepro; reservar(&d_Rout, resSize); - gJoinWithWriteNegative<<>> (d_locations, rLen, d_Rout, p1, of1, dcons, projp.x, posR); + hipLaunchKernel(HIP_KERNEL_NAME(gJoinWithWriteNegative), dim3(blockllen), dim3(numthreads), sizepro, 0, d_locations, rLen, d_Rout, p1, of1, dcons, projp.x, posR); } - cudaFree(d_R); - cudaFree(d_S); + hipFree(d_R); + hipFree(d_S); } else { @@ -1200,26 +1201,26 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: { muljoin = numj - 2; muljoinsize = muljoin * sizeof(int); - cudaMemcpy(dcons, wherej + 2, muljoinsize, cudaMemcpyHostToDevice); - gIndexMultiJoin<<>> (d_R, d_S, d_locations, sLen, temp, p1, p2, of1, of2, posR, posS, dcons, muljoin); + hipMemcpy(dcons, wherej + 2, muljoinsize, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(gIndexMultiJoin), dim3(blockllen), dim3(numthreads), muljoinsize, 0, d_R, d_S, d_locations, sLen, temp, p1, p2, of1, of2, posR, posS, dcons, muljoin); } else - gIndexJoin<<>> (d_R, d_S, d_locations, sLen, temp); - cudaFree(d_R); - cudaFree(d_S); + hipLaunchKernel(HIP_KERNEL_NAME(gIndexJoin), dim3(blockllen), dim3(numthreads), 0, 0, d_R, d_S, d_locations, sLen, temp); + hipFree(d_R); + hipFree(d_S); sum = res[sLen-1]; thrust::exclusive_scan(res, res + sLen, res); sum += res[sLen-1]; if(sum == 0) { - cudaFree(dcons); - cudaFree(d_locations); - cudaFree(temp); + hipFree(dcons); + hipFree(d_locations); + hipFree(temp); if(posS != NULL) - cudaFree(posS); + hipFree(posS); if(posR != NULL) - cudaFree(posR); + hipFree(posR); return 0; } res[sLen] = sum; @@ -1227,49 +1228,49 @@ int join(int *p1, int *p2, int rLen, int sLen, int of1, int of2, list: if(pos == (rule->num_rows - 3)) { sizepro = rule->num_columns * sizeof(int); - cudaMemcpy(dcons, proj, sizepro, cudaMemcpyHostToDevice); + hipMemcpy(dcons, proj, sizepro, hipMemcpyHostToDevice); resSize = sum * sizepro; reservar(&d_Rout, resSize); if(numj > 2) { - cudaMemcpy(dcons + rule->num_columns, wherej + 2, muljoinsize, cudaMemcpyHostToDevice); - multiJoinWithWrite2<<>> (d_locations, sLen, temp, d_Rout, p1, p2, of1, of2, dcons, rule->num_columns, posR, posS, muljoin); + hipMemcpy(dcons + rule->num_columns, wherej + 2, muljoinsize, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(multiJoinWithWrite2), dim3(blockllen), dim3(numthreads), sizepro + muljoinsize, 0, d_locations, sLen, temp, d_Rout, p1, p2, of1, of2, dcons, rule->num_columns, posR, posS, muljoin); } else - gJoinWithWrite2<<>> (d_locations, sLen, temp, d_Rout, p1, p2, of1, of2, dcons, rule->num_columns, posR, posS); + hipLaunchKernel(HIP_KERNEL_NAME(gJoinWithWrite2), dim3(blockllen), dim3(numthreads), sizepro, 0, d_locations, sLen, temp, d_Rout, p1, p2, of1, of2, dcons, rule->num_columns, posR, posS); } else { sizepro = projp.y * sizeof(int); - cudaMemcpy(dcons, proj, sizepro, cudaMemcpyHostToDevice); + hipMemcpy(dcons, proj, sizepro, hipMemcpyHostToDevice); resSize = sum * sizepro; reservar(&d_Rout, resSize); if(numj > 2) { - cudaMemcpy(dcons + projp.y, wherej + 2, muljoinsize, cudaMemcpyHostToDevice); - multiJoinWithWrite<<>> (d_locations, sLen, temp, d_Rout, p1, p2, of1, of2, dcons, projp.x, projp.y, posR, posS, muljoin); + hipMemcpy(dcons + projp.y, wherej + 2, muljoinsize, hipMemcpyHostToDevice); + hipLaunchKernel(HIP_KERNEL_NAME(multiJoinWithWrite), dim3(blockllen), dim3(numthreads), sizepro + muljoinsize, 0, d_locations, sLen, temp, d_Rout, p1, p2, of1, of2, dcons, projp.x, projp.y, posR, posS, muljoin); } else - gJoinWithWrite<<>> (d_locations, sLen, temp, d_Rout, p1, p2, of1, of2, dcons, projp.x, projp.y, posR, posS); + hipLaunchKernel(HIP_KERNEL_NAME(gJoinWithWrite), dim3(blockllen), dim3(numthreads), sizepro, 0, d_locations, sLen, temp, d_Rout, p1, p2, of1, of2, dcons, projp.x, projp.y, posR, posS); } } - cudaFree(dcons); - cudaFree(d_locations); - cudaFree(temp); + hipFree(dcons); + hipFree(d_locations); + hipFree(temp); if(posS != NULL) - cudaFree(posS); + hipFree(posS); if(posR != NULL) - cudaFree(posR); + hipFree(posR); if(*ret != NULL) - cudaFree(*ret); + hipFree(*ret); *ret = d_Rout; #ifdef TIMER - cudaEventRecord(stop, 0); - cudaEventSynchronize(stop); - cudaEventElapsedTime(&time, start, stop); + hipEventRecord(stop, 0); + hipEventSynchronize(stop); + hipEventElapsedTime(&time, start, stop); //cout << "Join = " << time << endl; //cout << "FIN" << endl; cuda_stats.join_time += time; diff --git a/packages/cuda/union2.cu b/packages/cuda/union2.cu old mode 100755 new mode 100644 index 1be5baf5e..125936393 --- a/packages/cuda/union2.cu +++ b/packages/cuda/union2.cu @@ -87,8 +87,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -122,8 +122,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -157,8 +157,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -192,8 +192,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -227,8 +227,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -262,8 +262,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -297,8 +297,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -332,8 +332,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -367,8 +367,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -402,8 +402,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -437,8 +437,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -472,8 +472,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -507,8 +507,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -542,8 +542,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -577,8 +577,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -612,8 +612,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -647,8 +647,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -682,8 +682,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -717,8 +717,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; @@ -752,8 +752,8 @@ int unir(int *res, int rows, int tipo, int **ret, int final) { size = nrows * tipo * sizeof(int); reservar(&nres, size); - cudaMemcpyAsync(nres, res, size, cudaMemcpyDeviceToDevice); - cudaFree(*ret); + hipMemcpyAsync(nres, res, size, hipMemcpyDeviceToDevice); + hipFree(*ret); *ret = nres; } return nrows; diff --git a/packages/cuda/union2.h b/packages/cuda/union2.h old mode 100755 new mode 100644 diff --git a/packages/cuda/unioncpu2.cpp b/packages/cuda/unioncpu2.cpp old mode 100755 new mode 100644 diff --git a/packages/myddas/example b/packages/myddas/example index fb3b4f544..53503fc19 100644 --- a/packages/myddas/example +++ b/packages/myddas/example @@ -2,19 +2,21 @@ sqlite3 side: create table test (id integer, x integer, y integer); .separator "," -.import /home/vsc/Yap/ILP/HH/DaysInHospital_Y3.csv test +.import DaysInHospital_Y3.csv hh +.export hh +.save hh myddas side: use_module(library(myddas)). -db_open(sqlite3,con,'../hh',x,x). +db_open(sqlite3,con,'.hh',x,x). db_close(con). test 2: use_module(library(myddas)). assert(( -t2 :- -db_open(sqlite3,con,'../hh',x,x), +t2 :- +db_open(sqlite3,con,'hh',x,x), db_import(con,test,test), test(A,B,C), writeln(test(A,B,C)), @@ -22,4 +24,3 @@ fail )). trace. t2. - diff --git a/packages/myddas/pl/myddas.ypp b/packages/myddas/pl/myddas.ypp index b5cea49bc..917c4d65d 100644 --- a/packages/myddas/pl/myddas.ypp +++ b/packages/myddas/pl/myddas.ypp @@ -15,7 +15,12 @@ * * *************************************************************************/ -#if USE_MYDDAS +#if 1 //USE_MYDDAS + +:- load_foreign_files([myddas], [], init_myddas). + +/* Initialize MYDDAS GLOBAL STRUCTURES */ +:- c_db_initialize_myddas. #ifdef DEBUG :- yap_flag(single_var_warnings,on). @@ -800,8 +805,6 @@ ]). #ifdef MYDDAS_MYSQL -:- load_foreign_files([], [], init_mysql). - :- use_module(myddas_mysql,[ db_my_result_set/1, db_datalog_describe/1, @@ -867,7 +870,7 @@ % db_open/4 % -#if MYDDAS_DECLARATIONS +#if 1 // MYDDAS_DECLARATIONS :- db_open(Protocol) extra_arguments db=Db, port=Port, diff --git a/packages/myddas/sqlite3/CMakeLists.txt b/packages/myddas/sqlite3/CMakeLists.txt index 342b43862..f89a52225 100644 --- a/packages/myddas/sqlite3/CMakeLists.txt +++ b/packages/myddas/sqlite3/CMakeLists.txt @@ -23,12 +23,11 @@ Android/jni/sqlite/nativehelper/jni.h #sqlite3 is now in the system set (SQLITE3_FOUND ON PARENT_SCOPE) - macro_log_feature (SQLITE3_FOUND "Sqlite3" - "Sqlite3 Data-Base " - "http://www.sqlite3ql.org" FALSE) + message ( + " * Sqlite3 Data-Base (http://www.sqlite3ql.org) is distributed with + MYDDAS" ) - - include_directories (${SQLITE3_INCLUDE_DIRECTORIES} .. . Android/jni/sqlite Android/jni/sqlite/nativehelper) + include_directories ( .. . Android/jni/sqlite Android/jni/sqlite/nativehelper) set (MYDDAS_FLAGS ${MYDDAS_FLAGS} -DMYDDAS_SQLITE3=1 PARENT_SCOPE) set_property( DIRECTORY .. APPEND PROPERTY COMPILE_DEFINITIONS MYDDAS_SQLITE3=1 ) diff --git a/packages/python/setup.py.cmake b/packages/python/setup.py.cmake new file mode 100644 index 000000000..3cd9115b4 --- /dev/null +++ b/packages/python/setup.py.cmake @@ -0,0 +1,10 @@ +from setuptools import setup, Extension + + +setup( + name = "yapex", + version = "0.1", +package_dir = {'': '${CMAKE_SOURCE_DIR}/packages/python' }, +py_modules = ['yapex'] + +) diff --git a/packages/python/yap_kernel/prolog.js b/packages/python/yap_kernel/prolog.js new file mode 100644 index 000000000..f8d172812 --- /dev/null +++ b/packages/python/yap_kernel/prolog.js @@ -0,0 +1,1283 @@ +// CodeMirror, copyright (c) by Marijn Haverbeke and others +// Distributed under an MIT license: http://codemirror.net/LICENSE + +(function(mod) { + if (typeof exports == "object" && typeof module == "object") // CommonJS + mod(require("../../lib/codemirror")); + else if (typeof define == "function" && define.amd) // AMD + define(["../../lib/codemirror"], mod); + else // Plain browser env + mod(CodeMirror); +})(function(CodeMirror) { + "use strict"; + + CodeMirror.defineMode("prolog", function(cmConfig, modeConfig) { + + function chain(stream, state, f) { + state.tokenize = f; + return f(stream, state); + } + + /******************************* + * CONFIG DATA * + *******************************/ + + var config = { quasiQuotations: false, /* {|Syntax||Quotation|} */ + dicts: false, /* tag{k:v, ...} */ + unicodeEscape: true, /* \uXXXX and \UXXXXXXXX */ + multiLineQuoted: true, /* "...\n..." */ + groupedIntegers: false /* 10 000 or 10_000 */ + }; + + var quoteType = { '"': "string", + "'": "qatom", + "`": "bqstring" + }; + + var isSingleEscChar = /[abref\\'"nrtsv]/; + var isOctalDigit = /[0-7]/; + var isHexDigit = /[0-9a-fA-F]/; + + var isSymbolChar = /[-#$&*+./:<=>?@\\^~]/; /* Prolog glueing symbols chars */ + var isSoloChar = /[[\]{}(),;|!]/; /* Prolog solo chars */ + var isNeck = /^(:-|-->)$/; + var isControlOp = /^(,|;|->|\*->|\\+|\|)$/; + + + /******************************* + * CHARACTER ESCAPES * + *******************************/ + + function readDigits(stream, re, count) { + if ( count > 0 ) { + while( count-- > 0 ) { + if ( !re.test(stream.next()) ) + return false; + } + } else { + while ( re.test(stream.peek()) ) + stream.next(); + } + return true; + } + + function readEsc(stream) { + var next = stream.next(); + if ( isSingleEscChar.test(next) ) + return true; + switch( next ) + { case "u": + if ( config.unicodeEscape ) + return readDigits(stream, isHexDigit, 4); /* SWI */ + return false; + case "U": + if ( config.unicodeEscape ) + return readDigits(stream, isHexDigit, 8); /* SWI */ + return false; + case null: return true; /* end of line */ + case "c": stream.eatSpace(); return true; + case "x": return readDigits(stream, isHexDigit, 2); + } + if ( isOctalDigit.test(next) ) { + if ( !readDigits(stream, isOctalDigit, -1) ) + return false; + if ( stream.peek() == "\\" ) /* SWI: optional closing \ */ + stream.next(); + return true; + } + return false; + } + + function nextUntilUnescaped(stream, state, end) { + var next; + while ((next = stream.next()) != null) { + if ( next == end && end != stream.peek() ) + { state.nesting.pop(); + return false; + } + if ( next == "\\" ) + { if ( !readEsc(stream) ) + return false; + } + } + return config.multiLineQuoted; + } + + /******************************* + * CONTEXT NESTING * + *******************************/ + + function nesting(state) { + return state.nesting.slice(-1)[0]; + } + + /* Called on every non-comment token */ + function setArg1(state) { + var nest = nesting(state); + if ( nest ) { + if ( nest.arg == 0 ) /* nested in a compound */ + nest.arg = 1; + else if ( nest.type == "control" ) + state.goalStart = false; + } else + state.goalStart = false; + } + + function setArgAlignment(state) { + var nest = nesting(state); + if ( nest && !nest.alignment && nest.arg != undefined ) { + if ( nest.arg == 0 ) + nest.alignment = nest.leftCol ? nest.leftCol+4 : nest.column+4; + else + nest.alignment = nest.column+1; + } + } + + function nextArg(state) { + var nest = nesting(state); + if ( nest ) { + if ( nest.arg ) /* nested in a compound */ + nest.arg++; + else if ( nest.type == "control" ) + state.goalStart = true; /* FIXME: also needed for ; and -> */ + } else + state.goalStart = true; + } + + function isControl(state) { /* our terms are goals */ + var nest = nesting(state); + if ( nest ) { + if ( nest.type == "control" ) { + return true; + } + return false; + } else + return state.inBody; + } + + // Used as scratch variables to communicate multiple values without + // consing up tons of objects. + var type, content; + function ret(tp, style, cont) { + type = tp; content = cont; + return style; + } + + function peekSpace(stream) { /* TBD: handle block comment as space */ + if ( stream.eol() || + /[\s%]/.test(stream.peek()) ) + return true; + return false; + } + + + /******************************* + * SUB TOKENISERS * + *******************************/ + + function plTokenBase(stream, state) { + var ch = stream.next(); + + if ( ch == "(" ) { + if ( state.lastType == "functor" ) { + state.nesting.push({ functor: state.functorName, + column: stream.column(), + leftCol: state.functorColumn, + arg: 0 + }); + delete state.functorName; + delete state.functorColumn; + } else { + state.nesting.push({ type: "control", + closeColumn: stream.column(), + alignment: stream.column()+4 + }); + } + return ret("solo", null, "("); + } + + if ( ch == "{" && state.lastType == "tag" ) { + state.nesting.push({ tag: state.tagName, + column: stream.column(), + leftCol: state.tagColumn, + arg: 0 + }); + delete state.tagName; + delete state.tagColumn; + return ret("dict_open", null); + } + + if ( ch == "/" && stream.eat("*") ) + return chain(stream, state, plTokenComment); + + if ( ch == "%" ) { + stream.skipToEnd(); + return ret("comment", "comment"); + } + + setArg1(state); + + if ( isSoloChar.test(ch) ) { + switch ( ch ) + { case ")": + state.nesting.pop(); + break; + case "]": + state.nesting.pop(); + return ret("list_close", null); + case "}": + { var nest = nesting(state); + var type = (nest && nest.tag) ? "dict_close" : "brace_term_close"; + + state.nesting.pop(); + return ret(type, null); + } + case ",": + if ( stream.eol() ) + state.commaAtEOL = true; + nextArg(state); + /*FALLTHROUGH*/ + case ";": + if ( isControl(state) ) + state.goalStart = true; + break; + case "[": + state.nesting.push({ type: "list", + closeColumn: stream.column(), + alignment: stream.column()+2 + }); + return ret("list_open", null); + break; + case "{": + if ( config.quasiQuotations && stream.eat("|") ) { + state.nesting.push({ type: "quasi-quotation", + alignment: stream.column()+1 + }); + return ret("qq_open", "qq_open"); + } else { + state.nesting.push({ type: "curly", + closeColumn: stream.column(), + alignment: stream.column()+2 + }); + return ret("brace_term_open", null); + } + break; + case "|": + if ( config.quasiQuotations ) { + if ( stream.eat("|") ) { + state.tokenize = plTokenQuasiQuotation; + return ret("qq_sep", "qq_sep"); + } else if ( stream.eat("}") ) { + state.nesting.pop(); + return ret("qq_close", "qq_close"); + } + } + if ( isControl(state) ) + state.goalStart = true; + break; + } + return ret("solo", null, ch); + } + + if (ch == '"' || ch == "'" || ch == "`") + { state.nesting.push({ type: "quoted", + alignment: stream.column()+1 + }); + return chain(stream, state, plTokenString(ch)); + } + + if ( ch == "0" ) { + if ( stream.eat(/x/i)) { + stream.eatWhile(/[\da-f]/i); + return ret("number", "number"); + } + if ( stream.eat(/o/i)) { + stream.eatWhile(/[0-7]/i); + return ret("number", "number"); + } + if ( stream.eat(/'/) ) { /* 0' */ + var next = stream.next(); + if ( next == "\\" ) { + if ( !readEsc(stream) ) + return ret("error", "error"); + } + return ret("code", "code"); + } + } + + if ( /\d/.test(ch) || /[+-]/.test(ch) && stream.eat(/\d/)) { + if ( config.groupedIntegers ) + stream.match(/^\d*((_|\s+)\d+)*(?:\.\d+)?(?:[eE][+\-]?\d+)?/); + else + stream.match(/^\d*(?:\.\d+)?(?:[eE][+\-]?\d+)?/); + return ret(ch == "-" ? "neg-number" : + ch == "+" ? "pos-number" : + "number"); + } + + if ( isSymbolChar.test(ch) ) { + stream.eatWhile(isSymbolChar); + var atom = stream.current(); + if ( atom == "." && peekSpace(stream) ) { + if ( nesting(state) ) { + return ret("fullstop", "error", atom); + } else { + } return ret("fullstop", "fullstop", atom); + } else if ( isNeck.test(atom) ) { + return ret("neck", "neck", atom); + } else if ( isControl(state) && isControlOp.test(atom) ) { + state.goalStart = true; + return ret("symbol", "operator", atom); + } else + return ret("symbol", "operator", atom); + } + + stream.eatWhile(/[\w_]/); + var word = stream.current(); + if ( stream.peek() == "{" && config.dicts ) { + state.tagName = word; /* tmp state extension */ + state.tagColumn = stream.column(); + return ret("tag", "tag", word); + } else if ( ch == "_" ) { + if ( word.length == 1 ) { + return ret("var", "anon", word); + } else { + var sec = word.charAt(1); + if ( sec == sec.toUpperCase() ) + return ret("var", "var-2", word); + } + return ret("var", "var", word); + } else if ( ch == ch.toUpperCase() ) { + return ret("var", "var", word); + } else if ( stream.peek() == "(" ) { + state.functorName = word; /* tmp state extension */ + state.functorColumn = stream.column(); + return ret("functor", "functor", word); + } else + return ret("atom", "atom", word); + } + + function plTokenString(quote) { + return function(stream, state) { + if (!nextUntilUnescaped(stream, state, quote)) { + state.tokenize = plTokenBase; + if ( stream.peek() == "(" ) { /* 'quoted functor'() */ + var word = stream.current(); + state.functorName = word; /* tmp state extension */ + return ret("functor", "functor", word); + } + if ( stream.peek() == "{" && config.dicts ) { /* 'quoted tag'{} */ + var word = stream.current(); + state.tagName = word; /* tmp state extension */ + return ret("tag", "tag", word); + } + } + return ret(quoteType[quote], quoteType[quote]); + }; + } + + function plTokenQuasiQuotation(stream, state) { + var maybeEnd = false, ch; + while (ch = stream.next()) { + if (ch == "}" && maybeEnd) { + state.tokenize = plTokenBase; + stream.backUp(2); + break; + } + maybeEnd = (ch == "|"); + } + return ret("qq_content", "qq_content"); + } + + function plTokenComment(stream, state) { + var maybeEnd = false, ch; + while (ch = stream.next()) { + if (ch == "/" && maybeEnd) { + state.tokenize = plTokenBase; + break; + } + maybeEnd = (ch == "*"); + } + return ret("comment", "comment"); + } + + + // /******************************* + // * ACTIVE KEYS * + // *******************************/ + + // /* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + // Support if-then-else layout like this: + + // goal :- + // ( Condition + // -> IfTrue + // ; IfFalse + // ). + // - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - */ + + + // CodeMirror.commands.prologStartIfThenElse = function(cm) { + // var start = cm.getCursor("start"); + // var token = cm.getTokenAt(start, true); + + // if ( token.state.goalStart == true ) + // { cm.replaceSelection("( ", "end"); + // return; + // } + + // return CodeMirror.Pass; + // } + + // CodeMirror.commands.prologStartThen = function(cm) { + // var start = cm.getCursor("start"); + // var token = cm.getTokenAt(start, true); + + // /* FIXME: These functions are copied from prolog.js. How + // can we reuse these? + // */ + // function nesting(state) { + // var len = state.nesting.length; + // if ( len > 0 ) + // return state.nesting[len-1]; + // return null; + // } + + // function isControl(state) { /* our terms are goals */ + // var nest = nesting(state); + // if ( nest ) { + // if ( nest.type == "control" ) { + // return true; + // } + // return false; + // } else + // return state.inBody; + // } + + // if ( start.ch == token.end && + // token.type == "operator" && + // token.string == "-" && + // isControl(token.state) ) + // { cm.replaceSelection("> ", "end"); + // return; + // } + + // return CodeMirror.Pass; + // } + + // CodeMirror.commands.prologStartElse = function(cm) { + // var start = cm.getCursor("start"); + // var token = cm.getTokenAt(start, true); + + // if ( token.start == 0 && start.ch == token.end && + // !/\S/.test(token.string) ) + // { cm.replaceSelection("; ", "end"); + // return; + // } + + // return CodeMirror.Pass; + // } + + // CodeMirror.defineOption("prologKeys", null, function(cm, val, prev) { + // if (prev && prev != CodeMirror.Init) + // cm.removeKeyMap("prolog"); + // if ( val ) { + // var map = { name: "prolog", + // "'('": "prologStartIfThenElse", + // "'>'": "prologStartThen", + // "';'": "prologStartElse", + // "Ctrl-L": "refreshHighlight" + // }; + // cm.addKeyMap(map); + // } + // }); + + // }); + //Default (SWI-)Prolog operator table. To be used later to enhance the + //offline experience. + + var ops = { "-->": { p:1200, t:"xfx" }, + ":-": [ { p:1200, t:"xfx" }, + { p:1200, t:"fx" } + ], + "?-": { p:1200, t:"fx" }, + + "dynamic": { p:1150, t:"fx" }, + "discontiguous": { p:1150, t:"fx" }, + "initialization": { p:1150, t:"fx" }, + "meta_predicate": { p:1150, t:"fx" }, + "module_transparent": { p:1150, t:"fx" }, + "multifile": { p:1150, t:"fx" }, + "thread_local": { p:1150, t:"fx" }, + "volatile": { p:1150, t:"fx" }, + + ";": { p:1100, t:"xfy" }, + "|": { p:1100, t:"xfy" }, + + "->": { p:1050, t:"xfy" }, + "*->": { p:1050, t:"xfy" }, + + ",": { p:1000, t:"xfy" }, + + "\\+": { p:900, t:"fy" }, + + "~": { p:900, t:"fx" }, + + "<": { p:700, t:"xfx" }, + "=": { p:700, t:"xfx" }, + "=..": { p:700, t:"xfx" }, + "=@=": { p:700, t:"xfx" }, + "=:=": { p:700, t:"xfx" }, + "=<": { p:700, t:"xfx" }, + "==": { p:700, t:"xfx" }, + "=\\=": { p:700, t:"xfx" }, + ">": { p:700, t:"xfx" }, + ">=": { p:700, t:"xfx" }, + "@<": { p:700, t:"xfx" }, + "@=<": { p:700, t:"xfx" }, + "@>": { p:700, t:"xfx" }, + "@>=": { p:700, t:"xfx" }, + "\\=": { p:700, t:"xfx" }, + "\\==": { p:700, t:"xfx" }, + "is": { p:700, t:"xfx" }, + + ":": { p:600, t:"xfy" }, + + "+": [ { p:500, t:"yfx" }, + { p:200, t:"fy" } + ], + "-": [ { p:500, t:"yfx" }, + { p:200, t:"fy" } + ], + "/\\": { p:500, t:"yfx" }, + "\\/": { p:500, t:"yfx" }, + "xor": { p:500, t:"yfx" }, + + "?": { p:500, t:"fx" }, + + "*": { p:400, t:"yfx" }, + "/": { p:400, t:"yfx" }, + "//": { p:400, t:"yfx" }, + "rdiv": { p:400, t:"yfx" }, + "<<": { p:400, t:"yfx" }, + ">>": { p:400, t:"yfx" }, + "mod": { p:400, t:"yfx" }, + "rem": { p:400, t:"yfx" }, + + "**": { p:200, t:"xfx" }, + "^": { p:200, t:"xfy" }, + + "\\": { p:200, t:"fy" } + }; + + var translType = { + "comment": "comment", + "var": "variable-2", /* JavaScript Types */ + "atom": "atom", + "qatom": "atom", + "bqstring": "string", + "symbol": "atom", + "functor": "keyword", + "tag": "tag", + "number": "number", + "string": "string", + "code": "number", + "neg-number": "number", + "pos-number": "number", + "list_open": "bracket", + "list_close": "bracket", + "qq_open": "bracket", + "qq_sep": "operator", + "qq_close": "bracket", + "dict_open": "bracket", + "dict_close": "bracket", + "brace_term_open": "bracket", + "brace_term_close": "bracket", + "neck": "keyword", + "fullstop": "keyword" + }; + + var builtins = { + "asserta": "prolog", + "atomic_list_concat": "prolog", + "char_type": "prolog", + "compile_expressions": "prolog", + "compile": "prolog", + "create_prolog_flag": "prolog", + "current_module": "prolog", + "current_op": "prolog", + "del_attrs": "prolog", + "depth_bound_call": "prolog", + "dule": "prolog", + "exo_files": "prolog", + "export_list": "prolog", + "foreign_directory": "prolog", + "garbage_collect_atoms": "prolog", + "garbage_collect": "prolog", + "get_attrs": "prolog", + "hread_signal": "prolog", + "ignore": "prolog", + "incore": "prolog", + "initialization": "prolog", + "int_message": "prolog", + "message_to_string": "prolog", + "module_property": "prolog", + "msort": "prolog", + "mutex_unlock_all": "prolog", + "no_style_check": "prolog", + "nospy": "prolog", + "notrace": "prolog", + "ortray_clause": "prolog", + "otherwise": "prolog", + "predsort": "prolog", + "prolog_initialization": "prolog", + "qend_program": "prolog", + "qsave_file": "prolog", + "recordaifnot": "prolog", + "set_base_module": "prolog", + "sformat": "prolog", + "source_file": "prolog", + "split_path_file": "prolog", + "stream_position": "prolog", + "system_error": "prolog", + "system_module": "prolog", + "t_head": "prolog", + "table_statistics": "prolog", + "tabling_mode": "prolog", + "tabling_statistics": "prolog", + "thread_defaults": "prolog", + "thread_local": "prolog", + "thread_set_defaults": "prolog", + "thread_statistics": "prolog", + "unix": "prolog", + "use_system_module": "prolog", + "user_defined_directive": "prolog", + "version": "prolog", + "C": "prolog", + "abolish_all_tables": "prolog", + "abolish_frozen_choice_points": "prolog", + "abolish_module": "prolog", + "abolish_table": "prolog", + "abolish": "prolog", + "abort": "prolog", + "absolute_file_name": "prolog", + "absolute_file_system_path": "prolog", + "access_file": "prolog", + "access": "prolog", + "acyclic_term": "prolog", + "add_import_module": "prolog", + "add_to_array_element": "prolog", + "add_to_path": "prolog", + "alarm": "prolog", + "all": "prolog", + "always_prompt_user": "prolog", + "arena_size": "prolog", + "arg": "prolog", + "array_element": "prolog", + "array": "prolog", + "assert_static": "prolog", + "asserta_static": "prolog", + "assertz_static": "prolog", + "assertz": "prolog", + "assert": "prolog", + "at_end_of_line": "prolog", + "at_end_of_stream_0": "prolog", + "at_end_of_stream": "prolog", + "at_halt": "prolog", + "atom_chars": "prolog", + "atom_codes": "prolog", + "atom_concat": "prolog", + "atom_length": "prolog", + "atom_number": "prolog", + "atom_string": "prolog", + "atom_to_term": "prolog", + "atomic_concat": "prolog", + "atomic_length": "prolog", + "atomic_list_concat": "prolog", + "atomics_to_string": "prolog", + "atomic": "prolog", + "atom": "prolog", + "attvar": "prolog", + "b_getval": "prolog", + "b_setval": "prolog", + "bagof": "prolog", + "bb_delete": "prolog", + "bb_get": "prolog", + "bb_put": "prolog", + "bb_update": "prolog", + "between": "prolog", + "bootstrap": "prolog", + "break": "prolog", + "call_cleanup": "prolog", + "call_count_data": "prolog", + "call_count_reset": "prolog", + "call_count": "prolog", + "call_residue_vars": "prolog", + "call_residue": "prolog", + "call_shared_object_function": "prolog", + "call_with_args": "prolog", + "callable": "prolog", + "call": "prolog", + "catch_ball": "prolog", + "catch": "prolog", + "cd": "prolog", + "cfile_search_path": "prolog", + "char_code": "prolog", + "char_conversion": "prolog", + "char_type": "prolog", + "clause_property": "prolog", + "clause": "prolog", + "close_shared_object": "prolog", + "close_static_array": "prolog", + "close": "prolog", + "code_type": "prolog", + "commons_directory": "prolog", + "commons_library": "prolog", + "compare": "prolog", + "compile_expressions": "prolog", + "compile_predicates": "prolog", + "compile": "prolog", + "compound": "prolog", + "consult_depth": "prolog", + "consult": "prolog", + "context_module": "prolog", + "copy_term_nat": "prolog", + "copy_term": "prolog", + "create_mutable": "prolog", + "create_prolog_flag": "prolog", + "creep_allowed": "prolog", + "current_atom": "prolog", + "current_char_conversion": "prolog", + "current_host": "prolog", + "current_input": "prolog", + "current_key": "prolog", + "current_line_number": "prolog", + "current_module": "prolog", + "current_mutex": "prolog", + "current_op": "prolog", + "current_output": "prolog", + "current_predicate": "prolog", + "current_prolog_flag": "prolog", + "current_reference_count": "prolog", + "current_stream": "prolog", + "current_thread": "prolog", + "db_files": "prolog", + "db_reference": "prolog", + "debugging": "prolog", + "debug": "prolog", + "decrease_reference_count": "prolog", + "del_attrs": "prolog", + "del_attr": "prolog", + "delete_import_module": "prolog", + "depth_bound_call": "prolog", + "dif": "prolog", + "discontiguous": "prolog", + "display": "prolog", + "do_c_built_in": "prolog", + "do_c_built_metacall": "prolog", + "do_not_compile_expressions": "prolog", + "dump_active_goals": "prolog", + "dum": "prolog", + "duplicate_term": "prolog", + "dynamic_predicate": "prolog", + "dynamic_update_array": "prolog", + "dynamic": "prolog", + "eamconsult": "prolog", + "eamtrans": "prolog", + "end_of_file": "prolog", + "ensure_loaded": "prolog", + "eraseall": "prolog", + "erased": "prolog", + "erase": "prolog", + "exists_directory": "prolog", + "exists_file": "prolog", + "exists_source": "prolog", + "exists": "prolog", + "exo_files": "prolog", + "expand_exprs": "prolog", + "expand_expr": "prolog", + "expand_file_name": "prolog", + "expand_goal": "prolog", + "expand_term": "prolog", + "expects_dialect": "prolog", + "export_list": "prolog", + "export_resource": "prolog", + "export": "prolog", + "extend": "prolog", + "fail": "prolog", + "false": "prolog", + "file_base_name": "prolog", + "file_directory_name": "prolog", + "file_exists": "prolog", + "file_name_extension": "prolog", + "file_search_path": "prolog", + "file_size": "prolog", + "fileerrors": "prolog", + "findall": "prolog", + "float": "prolog", + "flush_output": "prolog", + "forall": "prolog", + "foreign_directory": "prolog", + "format": "prolog", + "freeze_choice_point": "prolog", + "freeze": "prolog", + "frozen": "prolog", + "functor": "prolog", + "garbage_collect_atoms": "prolog", + "garbage_collect": "prolog", + "gc": "prolog", + "get0": "prolog", + "get_attr": "prolog", + "get_byte": "prolog", + "get_char": "prolog", + "get_code": "prolog", + "get_depth_limit": "prolog", + "get_mutable": "prolog", + "get_string_code": "prolog", + "get_value": "prolog", + "getcwd": "prolog", + "getenv": "prolog", + "get": "prolog", + "global_trie_statistics": "prolog", + "ground": "prolog", + "grow_heap": "prolog", + "grow_stack": "prolog", + "halt": "prolog", + "heap_space_info": "prolog", + "hide_atom": "prolog", + "hide_predicate": "prolog", + "hostname_address": "prolog", + "hread_get_message": "prolog", + "if": "prolog", + "ignore": "prolog", + "import_module": "prolog", + "incore": "prolog", + "increase_reference_count": "prolog", + "init_random_state": "prolog", + "initialization": "prolog", + "instance_property": "prolog", + "instance": "prolog", + "integer": "prolog", + "is_absolute_file_name": "prolog", + "is_list": "prolog", + "is_mutable": "prolog", + "is_tabled": "prolog", + "isinf": "prolog", + "isnan": "prolog", + "is": "prolog", + "key_erased_statistics": "prolog", + "key_statistics": "prolog", + "keysort": "prolog", + "leash": "prolog", + "length": "prolog", + "libraries_directories": "prolog", + "line_count": "prolog", + "listing": "prolog", + "load_absolute_foreign_files": "prolog", + "load_db": "prolog", + "load_files": "prolog", + "load_foreign_files": "prolog", + "log_event": "prolog", + "logsum": "prolog", + "ls_imports": "prolog", + "ls": "prolog", + "make_directory": "prolog", + "make_library_index": "prolog", + "make": "prolog", + "message_queue_create": "prolog", + "message_queue_destroy": "prolog", + "message_queue_property": "prolog", + "message_to_string": "prolog", + "mmapped_array": "prolog", + "module_property": "prolog", + "module_state": "prolog", + "module": "prolog", + "msort": "prolog", + "multifile": "prolog", + "must_be_of_type": "prolog", + "mutex_create": "prolog", + "mutex_property": "prolog", + "mutex_unlock_all": "prolog", + "name": "prolog", + "nb_create": "prolog", + "nb_current": "prolog", + "nb_delete": "prolog", + "nb_getval": "prolog", + "nb_linkarg": "prolog", + "nb_linkval": "prolog", + "nb_set_bit": "prolog", + "nb_set_shared_arg": "prolog", + "nb_set_shared_val": "prolog", + "nb_setarg": "prolog", + "nb_setval": "prolog", + "new_system_module": "prolog", + "nl": "prolog", + "no_source": "prolog", + "no_style_check": "prolog", + "nodebug": "prolog", + "nofileeleerrors": "prolog", + "nogc": "prolog", + "nonvar": "prolog", + "nospyall": "prolog", + "nospy": "prolog", + "notrace": "prolog", + "not": "prolog", + "nth_clause": "prolog", + "nth_instance": "prolog", + "number_atom": "prolog", + "number_chars": "prolog", + "number_codes": "prolog", + "number_string": "prolog", + "numbervars": "prolog", + "number": "prolog", + "on_exception": "prolog", + "on_signal": "prolog", + "once": "prolog", + "opaque": "prolog", + "open_pipe_stream": "prolog", + "open_shared_object": "prolog", + "open": "prolog", + "opt_statistics": "prolog", + "op": "prolog", + "or_statistics": "prolog", + "otherwise": "prolog", + "parallel_findall": "prolog", + "parallel_findfirst": "prolog", + "parallel_once": "prolog", + "parallel": "prolog", + "path": "prolog", + "peek_byte": "prolog", + "peek_char": "prolog", + "peek_code": "prolog", + "peek": "prolog", + "phrase": "prolog", + "plus": "prolog", + "portray_clause": "prolog", + "predicate_erased_statistics": "prolog", + "predicate_property": "prolog", + "predicate_statistics": "prolog", + "predmerge": "prolog", + "predsort": "prolog", + "primitive": "prolog", + "print_message_lines": "prolog", + "print_message": "prolog", + "print": "prolog", + "private": "prolog", + "profalt": "prolog", + "profend": "prolog", + "profile_data": "prolog", + "profile_reset": "prolog", + "profinit": "prolog", + "profoff": "prolog", + "profon": "prolog", + "prolog_current_frame": "prolog", + "prolog_file_name": "prolog", + "prolog_file_type": "prolog", + "prolog_flag_property": "prolog", + "prolog_flag": "prolog", + "prolog_initialization": "prolog", + "prolog_load_context": "prolog", + "prolog_to_os_filename": "prolog", + "prolog": "prolog", + "prompt1": "prolog", + "prompt": "prolog", + "put_attrs": "prolog", + "put_attr": "prolog", + "put_byte": "prolog", + "put_char1": "prolog", + "put_char": "prolog", + "put_code": "prolog", + "putenv": "prolog", + "put": "prolog", + "pwd": "prolog", + "qend_program": "prolog", + "qload_file": "prolog", + "qload_module": "prolog", + "qpack_clean_up_to_disjunction": "prolog", + "qsave_file": "prolog", + "qsave_module": "prolog", + "qsave_program": "prolog", + "raise_exception": "prolog", + "rational_term_to_tree": "prolog", + "rational": "prolog", + "read_clause": "prolog", + "read_sig": "prolog", + "read_term_from_atomic": "prolog", + "read_term_from_atom": "prolog", + "read_term_from_string": "prolog", + "read_term": "prolog", + "read": "prolog", + "real_path": "prolog", + "reconsult": "prolog", + "recorda_at": "prolog", + "recordaifnot": "prolog", + "recorda": "prolog", + "recorded": "prolog", + "recordz_at": "prolog", + "recordzifnot": "prolog", + "recordz": "prolog", + "release_random_state": "prolog", + "remove_from_path": "prolog", + "rename": "prolog", + "repeat": "prolog", + "reset_static_array": "prolog", + "reset_total_choicepoints": "prolog", + "resize_static_array": "prolog", + "restore": "prolog", + "retractall": "prolog", + "retract": "prolog", + "rmdir": "prolog", + "same_file": "prolog", + "save_program": "prolog", + "seeing": "prolog", + "seen": "prolog", + "see": "prolog", + "set_base_module": "prolog", + "set_input": "prolog", + "set_output": "prolog", + "set_prolog_flag": "prolog", + "set_random_state": "prolog", + "set_stream_position": "prolog", + "set_stream": "prolog", + "set_value": "prolog", + "setarg": "prolog", + "setenv": "prolog", + "setof": "prolog", + "setup_call_catcher_cleanup": "prolog", + "setup_call_cleanup": "prolog", + "sformat": "prolog", + "show_all_local_tables": "prolog", + "show_all_tables": "prolog", + "show_global_trieshow_tabled_predicates": "prolog", + "show_global_trie": "prolog", + "show_low_level_trace": "prolog", + "show_tabled_predicates": "prolog", + "show_table": "prolog", + "showprofres": "prolog", + "sh": "prolog", + "simple": "prolog", + "skip1": "prolog", + "skip": "prolog", + "socket_accept": "prolog", + "socket_bind": "prolog", + "socket_close": "prolog", + "socket_connect": "prolog", + "socket_listen": "prolog", + "socket": "prolog", + "sort2": "prolog", + "sort": "prolog", + "source_file_property": "prolog", + "source_file": "prolog", + "source_location": "prolog", + "source_mode": "prolog", + "source_module": "prolog", + "source": "prolog", + "split_path_file": "prolog", + "spy": "prolog", + "srandom": "prolog", + "start_low_level_trace": "prolog", + "stash_predicate": "prolog", + "static_array_location": "prolog", + "static_array_properties": "prolog", + "static_array_to_term": "prolog", + "static_array": "prolog", + "statistics": "prolog", + "stop_low_level_trace": "prolog", + "stream_position_data": "prolog", + "stream_position": "prolog", + "stream_property": "prolog", + "stream_select": "prolog", + "string_chars": "prolog", + "string_codes": "prolog", + "string_code": "prolog", + "string_concat": "prolog", + "string_length": "prolog", + "string_number": "prolog", + "string_to_atomic": "prolog", + "string_to_atom": "prolog", + "string_to_list": "prolog", + "string": "prolog", + "strip_module": "prolog", + "style_check": "prolog", + "sub_atom": "prolog", + "sub_string": "prolog", + "subsumes_term": "prolog", + "succ": "prolog", + "sys_debug": "prolog", + "system_error": "prolog", + "system_library": "prolog", + "system_module": "prolog", + "system_predicate": "prolog", + "system": "prolog", + "t_body": "prolog", + "t_head": "prolog", + "t_hgoal": "prolog", + "t_hlist": "prolog", + "t_tidy": "prolog", + "tab1": "prolog", + "table_statistics": "prolog", + "table": "prolog", + "tabling_mode": "prolog", + "tabling_statistics": "prolog", + "tab": "prolog", + "telling": "prolog", + "tell": "prolog", + "term_attvars": "prolog", + "term_factorized": "prolog", + "term_to_atom": "prolog", + "term_to_string": "prolog", + "term_variables": "prolog", + "thread_at_exit": "prolog", + "thread_cancel": "prolog", + "thread_create": "prolog", + "thread_defaults": "prolog", + "thread_default": "prolog", + "thread_detach": "prolog", + "thread_exit": "prolog", + "thread_get_message": "prolog", + "thread_join": "prolog", + "thread_local": "prolog", + "thread_peek_message": "prolog", + "thread_property": "prolog", + "thread_self": "prolog", + "thread_send_message": "prolog", + "thread_set_defaults": "prolog", + "thread_set_default": "prolog", + "thread_signal": "prolog", + "thread_sleep": "prolog", + "thread_statistics": "prolog", + "threads": "prolog", + "throw": "prolog", + "time_file64": "prolog", + "time_file": "prolog", + "time": "prolog", + "told": "prolog", + "tolower": "prolog", + "total_choicepoints": "prolog", + "total_erased": "prolog", + "toupper": "prolog", + "trace": "prolog", + "true_file_name": "prolog", + "true": "prolog", + "tthread_peek_message": "prolog", + "ttyget0": "prolog", + "ttyget": "prolog", + "ttynl": "prolog", + "ttyput": "prolog", + "ttyskip": "prolog", + "udi": "prolog", + "unhide_atom": "prolog", + "unify_with_occurs_check": "prolog", + "unix": "prolog", + "unknown": "prolog", + "unload_file": "prolog", + "unload_module": "prolog", + "unnumbervars": "prolog", + "update_array": "prolog", + "update_mutable": "prolog", + "use_module": "prolog", + "use_system_module": "prolog", + "user_defined_directive": "prolog", + "var": "prolog", + "version": "prolog", + "volatile": "prolog", + "wake_choice_point": "prolog", + "when": "prolog", + "with_mutex": "prolog", + "with_output_to": "prolog", + "working_directory": "prolog", + "write_canonical": "prolog", + "write_depth": "prolog", + "write_term": "prolog", + "writeln": "prolog", + "writeq": "prolog", + "write": "prolog", + "yap_flag": "prolog" + }; + + /******************************* + * RETURN OBJECT * + *******************************/ + + return { + startState: function() { + return { + tokenize: plTokenBase, + inBody: false, + goalStart: false, + lastType: null, + nesting: new Array(), /* ([{}]) nesting FIXME: copy this */ + curTerm: null, /* term index in metainfo */ + curToken: null /* token in term */ + }; + }, + + + token: function(stream, state) { + var nest; + + if ( state.curTerm == null && mode + Config.metainfo ) { + state.curTerm = 0; + state.curToken = 0; + } + + if ( stream.sol() ) + delete state.commaAtEOL; + + if ( state.tokenize == plTokenBase && stream.eatSpace() ) { + if ( stream.eol() ) + setArgAlignment(state); + return null; + } + + var style = state.tokenize(stream, state); + + if ( stream.eol() ) + setArgAlignment(state); + + if ( type == "neck" ) { + state.inBody = true; + state.goalStart = true; + } else if ( type == "fullstop" ) { + state.inBody = false; + state.goalStart = false; + } + + state.lastType = type; + + + if ( builtins[state.curToken] == "prolog") + return "builtin"; + if ( ops[state.curToken]) + return "operator"; + return translType[type]; + }, + + indent: function(state, textAfter) { + if (state.tokenize == plTokenComment) return CodeMirror.Pass; + + var nest; + if ( (nest=nesting(state)) ) { + if ( nest.closeColumn && !state.commaAtEOL ) + return nest.closeColumn; + return nest.alignment; + } + if ( !state.inBody ) + return 0; + + return 4; + }, + + // theme: "prolog", + + blockCommentStart: "/*", /* continuecomment.js support */ + blockCommentEnd: "*/", + blockCommentContinue: " * ", + lineComment: "%", + }; + + }); + +CodeMirror.defineMIME("text/x-prolog", "prolog"); +}); diff --git a/packages/yap-lbfgs/yap_lbfgs.c b/packages/yap-lbfgs/yap_lbfgs.c index e04baf40f..fbf224101 100644 --- a/packages/yap-lbfgs/yap_lbfgs.c +++ b/packages/yap-lbfgs/yap_lbfgs.c @@ -78,14 +78,11 @@ static lbfgsfloatval_t evaluate( a1 = YAP_ArgOfTerm(1,call); if (YAP_IsFloatTerm(a1)) { - YAP_ShutdownGoal( TRUE ); return (lbfgsfloatval_t) YAP_FloatOfTerm(a1); } else if (YAP_IsIntTerm(a1)) { - YAP_ShutdownGoal( TRUE ); return (lbfgsfloatval_t) YAP_IntOfTerm(a1); } - YAP_ShutdownGoal( TRUE ); fprintf(stderr, "ERROR: The evaluate call back function did not return a number as first argument.\n"); return 0; }