Changeset: 33fb5d0be306 for MonetDB URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=33fb5d0be306 Modified Files: monetdb5/extras/bwd/cl_program_utilities.c monetdb5/extras/bwd/cl_program_utilities.h monetdb5/extras/bwd/operations.c Branch: bwd Log Message:
* using kernel to zero out buffers instead of dma transfer of host-allocated
buffer (performance improvement is tremendous)
Unterschiede (125 Zeilen):
diff --git a/monetdb5/extras/bwd/cl_program_utilities.c
b/monetdb5/extras/bwd/cl_program_utilities.c
--- a/monetdb5/extras/bwd/cl_program_utilities.c
+++ b/monetdb5/extras/bwd/cl_program_utilities.c
@@ -54,6 +54,14 @@ cl_program compileProgram(const char* so
return program;
}
+cl_program getZeroOutProgram(){
+ const char* sourceCode = "__kernel void zeroOut (\n"
+ "__global int* buffer\n){\n"
+ " buffer[get_global_id(0)] = 0;"
+ "}";
+ return compileProgram(sourceCode, "");
+}
+
cl_program getMultiplyProgram(unsigned int approximationBits){
diff --git a/monetdb5/extras/bwd/cl_program_utilities.h
b/monetdb5/extras/bwd/cl_program_utilities.h
--- a/monetdb5/extras/bwd/cl_program_utilities.h
+++ b/monetdb5/extras/bwd/cl_program_utilities.h
@@ -13,5 +13,6 @@ cl_program getProjectionLeftjoinProgram(
cl_program getGroupProgram(unsigned int attributeCount, unsigned int*
approximationBits);
cl_program getHistogramCompactionProgram(void);
cl_program compileProgram(const char* sourceCode, char* options);
+cl_program getZeroOutProgram();
/* cl_program getProjectionSemijoinProgram(unsigned int approximationBits); */
#endif /* _CL_PROGRAM_UTILITIES_H_ */
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
@@ -33,6 +33,7 @@ extern const int VALUES_PER_WORK_ITEM;
/* #define MAX_INTERMEDIATE_RESULT_SIZE 16777216 */
/* #define MAX_INTERMEDIATE_RESULT_SIZE 59142609 */
#define MAX_INTERMEDIATE_RESULT_SIZE 67108864
+/* #define MAX_INTERMEDIATE_RESULT_SIZE 8388608 */
#pragma mark Actual MAL Operations Implementation
#ifdef __APPLE__
@@ -102,6 +103,30 @@ clTail* getApproximateValuesColumn(cl_me
return buffer;
}
+
+static inline cl_int bwdEnqueueNDRangeKernel ( cl_command_queue command_queue,
+
cl_kernel kernel,
+
cl_uint work_dim,
+
const size_t
*global_work_offset,
+
const size_t
*global_work_size,
+
const size_t
*local_work_size,
+
cl_uint
num_events_in_wait_list,
+
const cl_event
*event_wait_list,
+
cl_event *event){
+ return clEnqueueNDRangeKernel ( command_queue,
+
kernel,
+
work_dim,
+
global_work_offset,
+
global_work_size,
+
local_work_size,
+
num_events_in_wait_list,
+
event_wait_list,
+
event);
+
+}
+
+
+const unsigned int zeroIntPattern[1] = {};
#ifndef CL_API_SUFFIX__VERSION_1_2
cl_int clEnqueueFillBuffer(cl_command_queue command_queue ,
cl_mem buffer ,
@@ -113,39 +138,27 @@ cl_int clEnqueueFillBuffer(cl_command_qu
const cl_event * event_wait_list ,
cl_event * event ) {
if(pattern_size == 4 && ((int*)pattern)[0] == 0){
- cl_int err = 0;
- int* tmpBuffer = GDKzalloc(size);
- err = clEnqueueWriteBuffer(command_queue, buffer, CL_TRUE,
offset, size, tmpBuffer, num_events_in_wait_list, event_wait_list, event);
- if(err) printf("#%s, clEnqueueWriteBuffer: %s;\n", __func__,
clError(err));
- GDKfree(tmpBuffer);
- return err;
+ cl_int err = 0;
+ if(0){
+ int* tmpBuffer = GDKzalloc(size);
+ err = clEnqueueWriteBuffer(command_queue, buffer,
CL_TRUE, offset, size, tmpBuffer, num_events_in_wait_list, event_wait_list,
event);
+ if(err) printf("#%s, clEnqueueWriteBuffer: %s;\n",
__func__, clError(err));
+ GDKfree(tmpBuffer);
+ return err;
+ } else {
+ cl_kernel kernel = clCreateKernel(getZeroOutProgram(),
"zeroOut", &err);
+ clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);
+ err = bwdEnqueueNDRangeKernel(getCommandQueue(),
kernel, 1, (const size_t[]){0}, (size_t[]){size/4}, (const size_t[]){64},
num_events_in_wait_list, event_wait_list, event);
+ if(err) printf("#%s, bwdEnqueueNDRangeKernel: %s;\n",
__func__, clError(err));
+ err= clFinish(getCommandQueue());
+ if(err) printf("#%s, clFinish: %s;\n", __func__,
clError(err));
+ }
}
return -1;
};
#endif
-static inline cl_int bwdEnqueueNDRangeKernel ( cl_command_queue command_queue,
-
cl_kernel kernel,
-
cl_uint work_dim,
-
const size_t
*global_work_offset,
-
const size_t
*global_work_size,
-
const size_t
*local_work_size,
-
cl_uint
num_events_in_wait_list,
-
const cl_event
*event_wait_list,
-
cl_event *event){
- return clEnqueueNDRangeKernel ( command_queue,
-
kernel,
-
work_dim,
-
global_work_offset,
-
global_work_size,
-
local_work_size,
-
num_events_in_wait_list,
-
event_wait_list,
-
event);
-
- }
-const unsigned int zeroIntPattern[1] = {};
size_t calculatedBufferSize(size_t headCount, size_t approximationBits){
return (((int)ceil(headCount*approximationBits/8.0))/8)*8+16;
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
