Changeset: 561de22dcce3 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=561de22dcce3
Modified Files:
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/utilities.c
Branch: bwd
Log Message:
* implemented uselect approximation kernel (for ints :-))
Unterschiede (102 Zeilen):
diff --git a/monetdb5/extras/bwd/operations.c b/monetdb5/extras/bwd/operations.c
--- a/monetdb5/extras/bwd/operations.c
+++ b/monetdb5/extras/bwd/operations.c
@@ -100,29 +100,47 @@ str BWDThetauselectApproximate(bat *res,
- const char* sourceCodeTemplate = "__kernel void select ("
- "__global %1$s* output,"
- "__global const %1$s* approximation,"
- "__global const %1$s operand,"
- "__global unsigned int output_cursor"
+ const char* sourceCodeTemplate = "__kernel void uselect ("
+ "__global void* output,"
+ "__global const void* approximation,"
+ "const %1$s operand,"
+ "__global unsigned int* output_cursor"
") {"
- " const %1$s value = (*(approximation +
get_global_id(0)*%3$d) << %4$d);"
+ " printf(\"%%d\\n\", get_global_id(0));"
+ " const %1$s value = (*(int*)(approximation +
get_global_id(0)*%3$d) << %4$d);"
" if(value %2$s operand)"
- " output[atomic_inc(output_cursor)] = value;"
+ " *(int*)(output + atomic_inc(output_cursor)*%3$d)
|= value;"
"}";
char* sourceCode = malloc(16384);
- snprintf(sourceCode, 16384, sourceCodeTemplate,
typeNames[BATttype(data)], OP, batTailApproximationBits(data)/8,
32-batTailApproximationBits(data));
+ snprintf(sourceCode, 16384, sourceCodeTemplate,
typeNames[BATttype(data)], *OP, batTailApproximationBits(data)/8,
32-batTailApproximationBits(data));
THRprintf(GDKout, "#%s, sourceCode: %s;\n", __func__,
sourceCode);
cl_int err;
cl_program program = clCreateProgramWithSource(getCLContext(),
1, (const char*[]){sourceCode}, (size_t[]){strlen(sourceCode)}, &err);
if(err) THRprintf(GDKout, "#%s, clCreateProgramWithSource:
%s;\n", __func__, clError(err));
- clBuildProgram(program, 1, (const
cl_device_id[]){getDeviceID()}, "", NULL, NULL);
- cl_kernel selectKernel = clCreateKernel(program, "select",
&err);
+ err = clBuildProgram(program, 1, (const
cl_device_id[]){getDeviceID()}, "", NULL, NULL);
+ if(err) {
+ THRprintf(GDKout, "#%s, clBuildProgram: %s;\n",
__func__, clError(err));
+ size_t bufferSize;
+ char* buffer;
+ clGetProgramBuildInfo(program, getDeviceID(),
CL_PROGRAM_BUILD_LOG, 0, NULL, &bufferSize);
+ clGetProgramBuildInfo(program, getDeviceID(),
CL_PROGRAM_BUILD_LOG, bufferSize, (buffer = alloca(bufferSize)), NULL);
+ THRprintf(GDKout, "#%s, clBuildProgram log: %s;\n",
__func__, buffer);
+ /* return "build failure"; */
+ }
+
+ cl_kernel selectKernel = clCreateKernel(program, "uselect",
&err);
if(err) THRprintf(GDKout, "#%s, clCreateKernel: %s;\n",
__func__, clError(err));
- clSetKernelArg(selectKernel, 0,
BATcount(data)*batTailApproximationBits(data)/8, batTailApproximation(result));
- clSetKernelArg(selectKernel, 1,
BATcount(data)*batTailApproximationBits(data)/8, batTailApproximation(data));
- clSetKernelArg(selectKernel, 2, sizeof(int), val); // type
specific
- clSetKernelArg(selectKernel, 3, sizeof(int), val); // type
specific
+ err = clSetKernelArg(selectKernel, 0, sizeof(cl_mem),
(cl_mem[]){batTailApproximation(result)}); if (err) THRprintf(GDKout, "#%s,
clSetKernelArg(%d): %s;\n", __func__, 0, clError(err));
+ err = clSetKernelArg(selectKernel, 1, sizeof(cl_mem),
(cl_mem[]){batTailApproximation(data)}); if (err) THRprintf(GDKout, "#%s,
clSetKernelArg(%d): %s;\n", __func__, 1, clError(err));
+ err = clSetKernelArg(selectKernel, 2, sizeof(int), val); if
(err) THRprintf(GDKout, "#%s, clSetKernelArg(%d): %s;\n", __func__, 2,
clError(err));
+
+ cl_mem outputCursor;
+ {
+ outputCursor = clCreateBuffer(getCLContext(),
CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(int), (int[]){0}, &err);
+ if(err) THRprintf(GDKout, "#%s, clCreateBuffer, cursor:
%s;\n", __func__, clError(err));
+ }
+
+ err = clSetKernelArg(selectKernel, 3, sizeof(cl_mem),
&outputCursor); if (err) THRprintf(GDKout, "#%s, clSetKernelArg(%d): %s;\n",
__func__, 3, clError(err));
clEnqueueNDRangeKernel(getCommandQueue(), selectKernel, 1,
(const size_t[]){0}, (const size_t[]){BATcount(data)}, (const size_t[]){2}, 0,
NULL, NULL);
free(sourceCode);
BBPkeepref(*res = result->batCacheid);
diff --git a/monetdb5/extras/bwd/utilities.c b/monetdb5/extras/bwd/utilities.c
--- a/monetdb5/extras/bwd/utilities.c
+++ b/monetdb5/extras/bwd/utilities.c
@@ -90,11 +90,11 @@ cl_device_id getDeviceID(){
}
-static cl_context clContextSingleton;
+static cl_context clContextSingleton = NULL;
cl_context getCLContext(){
if(clContextSingleton) return clContextSingleton;
int err;
- clContextSingleton = clCreateContext(0,1,
(cl_device_id[1]){getDeviceID()}, NULL, NULL, &err);
+ clContextSingleton = clCreateContext(0,1,
(cl_device_id[]){getDeviceID()}, NULL, NULL, &err);
if (err == CL_SUCCESS) return clContextSingleton;
else THRprintf(GDKout, "failure when creating the context\n");
return NULL;
@@ -105,7 +105,7 @@ static cl_command_queue queue = NULL;
cl_command_queue getCommandQueue(cl_device_id (*device) ()){
if(queue) return queue;
int err;
- queue =
clCreateCommandQueue(getCLContext(),device(),CL_QUEUE_PROFILING_ENABLE,&err);
+ queue =
clCreateCommandQueue(getCLContext(),getDeviceID(),CL_QUEUE_PROFILING_ENABLE,&err);
if (err == CL_SUCCESS) return queue;
else THRprintf(GDKout, "failed to create queue");
return NULL;
@@ -149,8 +149,8 @@ const unsigned char* batTailResiduals(co
}
const unsigned int getNextFreeDecomposedBATSlotIndex(){
- size_t i;
- for (i = 0; i < MAX_DECOMPOSED_BATS && bwdRegistry[i].residuals !=
NULL;)
+ size_t i = 0;
+ while (i < MAX_DECOMPOSED_BATS && bwdRegistry[i].approximationBits > 0)
i++;
return i;
};
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
