... | ... |
@@ -20,21 +20,18 @@ RCcompilerLevel1 <- function(optMeta1) { |
20 | 20 |
|
21 | 21 |
# matrix index offset |
22 | 22 |
gpu_gp_offset = GPUVar$global_private_offset |
23 |
+ gpu_gs_offset = GPUVar$global_shared_offset |
|
24 |
+ gpu_ls_offset = GPUVar$local_shared_offset |
|
23 | 25 |
|
24 | 26 |
# worker shared data, located in global memory |
25 | 27 |
gpu_gs_data = GPUVar$global_shared_data |
26 |
- gpu_gs_offset = GPUVar$global_shared_offset |
|
27 |
- |
|
28 | 28 |
# worker private data, located in private/local memory |
29 | 29 |
gpu_lp_data = GPUVar$local_private_data |
30 |
- gpu_lp_offset = GPUVar$local_private_offset |
|
31 |
- |
|
32 | 30 |
# worker shared data, located in local memory |
33 | 31 |
gpu_ls_data = GPUVar$local_shared_data |
34 |
- gpu_ls_offset = GPUVar$local_shared_offset |
|
35 | 32 |
|
36 |
- gpu_sizeInfo = GPUVar$size_info |
|
37 | 33 |
|
34 |
+ gpu_sizeInfo = GPUVar$size_info |
|
38 | 35 |
gpu_returnSize = GPUVar$return_size |
39 | 36 |
|
40 | 37 |
# Deducted variable |
... | ... |
@@ -51,8 +48,6 @@ RCcompilerLevel1 <- function(optMeta1) { |
51 | 48 |
|
52 | 49 |
|
53 | 50 |
var_def_code = c() |
54 |
- redirect_code=c() |
|
55 |
- redirect_var=c() |
|
56 | 51 |
gpu_gp_num = -1 |
57 | 52 |
gpu_gs_num = -1 |
58 | 53 |
gpu_lp_num = -1 |
... | ... |
@@ -80,13 +75,6 @@ RCcompilerLevel1 <- function(optMeta1) { |
80 | 75 |
if(isNA(curInfo$address)){ |
81 | 76 |
curInfo$address = curVar |
82 | 77 |
} |
83 |
- if(curInfo$redirect!="NA"){ |
|
84 |
- curCode=paste0("#define ", curVar, " ", curInfo$redirect) |
|
85 |
- curSize1=paste0("#define ", getSizeVar(curVar,1)," ",getSizeVar(curInfo$redirect,1)) |
|
86 |
- curSize2=paste0("#define ", getSizeVar(curVar,2)," ",getSizeVar(curInfo$redirect,2)) |
|
87 |
- redirect_code=c(redirect_code,curCode,curSize1,curSize2) |
|
88 |
- redirect_var=c(redirect_var,curVar) |
|
89 |
- } |
|
90 | 78 |
# If the variable is the function argument |
91 | 79 |
if (curInfo$require) { |
92 | 80 |
gpu_gs_num = gpu_gs_num + 1 |
... | ... |
@@ -101,11 +89,9 @@ RCcompilerLevel1 <- function(optMeta1) { |
101 | 89 |
|
102 | 90 |
|
103 | 91 |
if (!curInfo$isPointer) { |
104 |
- CXXtype = getTypeCXXStr(curInfo$precisionType) |
|
105 |
- curCode = paste0(CXXtype, " ", curVar, ";") |
|
106 |
- var_def_code = rbind(var_def_code,c(curVar, curCode)) |
|
107 |
- curInfo$address = curVar |
|
108 |
- varInfo = setVarInfo(varInfo, curInfo) |
|
92 |
+ res=addVariableDeclaration_NonPointer(varInfo,curInfo) |
|
93 |
+ var_def_code = rbind(var_def_code,c(curVar, res$code)) |
|
94 |
+ varInfo = setVarInfo(varInfo, res$Info) |
|
109 | 95 |
next |
110 | 96 |
} |
111 | 97 |
|
... | ... |
@@ -150,11 +136,37 @@ RCcompilerLevel1 <- function(optMeta1) { |
150 | 136 |
} |
151 | 137 |
|
152 | 138 |
|
139 |
+ gpu_gp_num=gpu_gp_num+1 |
|
140 |
+ gpu_gs_num=gpu_gs_num+1 |
|
141 |
+ gpu_lp_num=gpu_lp_num+1 |
|
142 |
+ gpu_ls_num=gpu_ls_num+1 |
|
143 |
+ # size_info: |
|
144 |
+ # return size, gp,gs,ls offset, gp,gs,lp,ls dim(row,col) |
|
145 |
+ data_offset=1 |
|
146 |
+ data_offset_var=c(gpu_gp_offset,gpu_gs_offset,gpu_ls_offset) |
|
147 |
+ data_offset_var_offset=c(gpu_gp_num,gpu_gs_num,gpu_ls_num) |
|
148 |
+ data_offset_macro=offsetMacro(gpu_sizeInfo,data_offset,data_offset_var,data_offset_var_offset) |
|
149 |
+ |
|
150 |
+ #Define matrix size macro |
|
151 |
+ dim_offset=data_offset+gpu_gp_num+gpu_gs_num+gpu_ls_num |
|
152 |
+ size_var=c(paste0(GPUVar$gp_size,1:2), |
|
153 |
+ paste0(GPUVar$gs_size,1:2), |
|
154 |
+ paste0(GPUVar$lp_size,1:2), |
|
155 |
+ paste0(GPUVar$ls_size,1:2)) |
|
156 |
+ size_offset=c(gpu_gp_num,gpu_gp_num, |
|
157 |
+ gpu_gs_num,gpu_gs_num, |
|
158 |
+ gpu_lp_num,gpu_lp_num, |
|
159 |
+ gpu_ls_num,gpu_ls_num) |
|
160 |
+ size_macro=offsetMacro(gpu_sizeInfo,dim_offset,size_var,size_offset) |
|
161 |
+ |
|
153 | 162 |
|
154 | 163 |
gpu_code = c( |
155 | 164 |
"//Define some useful macro", |
156 | 165 |
paste0("#define ", gpu_worker_offset, " ", gpu_global_id), |
157 | 166 |
paste0("#define ", gpu_element_dist, " ", gpu_global_id), |
167 |
+ paste0("#define ", gpu_returnSize, " ", gpu_sizeInfo,"[0]"), |
|
168 |
+ data_offset_macro, |
|
169 |
+ size_macro, |
|
158 | 170 |
|
159 | 171 |
"//Data preparation", |
160 | 172 |
paste0("size_t ", gpu_global_id, "=get_global_id(0);"), |
... | ... |
@@ -165,19 +177,30 @@ RCcompilerLevel1 <- function(optMeta1) { |
165 | 177 |
"//Matrix dimension optimization") |
166 | 178 |
|
167 | 179 |
|
168 |
- |
|
180 |
+ var_def_code=as.data.frame(var_def_code,stringsAsFactors=FALSE) |
|
181 |
+ colnames(var_def_code)=c("var","def") |
|
169 | 182 |
varInfo$var_def_code=var_def_code |
170 |
- varInfo$redirect_var=redirect_var |
|
171 | 183 |
|
172 | 184 |
GPUExp1$Exp = parsedExp |
173 | 185 |
GPUExp1$varInfo = varInfo |
174 | 186 |
GPUExp1$gpu_code = gpu_code |
175 |
- GPUExp1$macro=c(GPUExp1$macro,redirect_code) |
|
176 | 187 |
|
177 | 188 |
|
178 | 189 |
return(GPUExp1) |
179 | 190 |
} |
180 | 191 |
|
192 |
+offsetMacro<-function(targetVar,offset,macroName,macroLength){ |
|
193 |
+ redundant_ind=which(macroLength==0) |
|
194 |
+ macroName=macroName[-redundant_ind] |
|
195 |
+ macroLength=macroLength[-redundant_ind] |
|
196 |
+ macro_offset=c(0) |
|
197 |
+ for(i in seq_len(length(macroLength)-1)+1){ |
|
198 |
+ macro_offset=c(macro_offset,macro_offset[i-1]+macroLength[i-1]) |
|
199 |
+ } |
|
200 |
+ macro_offset=macro_offset+offset |
|
201 |
+ macro=paste0("#define ",macroName,"(x) ",targetVar,"[",macro_offset,"+x]") |
|
202 |
+ macro |
|
203 |
+} |
|
181 | 204 |
|
182 | 205 |
RCcompilerLevel2 <- function(GPUExp1) { |
183 | 206 |
if (DEBUG) { |
... | ... |
@@ -202,7 +225,8 @@ RCcompilerLevel2 <- function(GPUExp1) { |
202 | 225 |
|
203 | 226 |
RCTranslation <- function(varInfo, parsedExp,isTop=FALSE) { |
204 | 227 |
gpu_code = c() |
205 |
- delim=switch(isTop,"\\Main function delimiter",NULL) |
|
228 |
+ delim=switch(isTop,"//Main function delimiter",NULL) |
|
229 |
+ #delim=NULL |
|
206 | 230 |
for (i in seq_along(parsedExp)) { |
207 | 231 |
curExp = parsedExp[[i]] |
208 | 232 |
if (curExp == "{" || is.symbol(curExp)) |
... | ... |
@@ -43,14 +43,14 @@ C_assignment_dispatch <- function(varInfo, Exp) { |
43 | 43 |
|
44 | 44 |
# ==================element functions=========================== |
45 | 45 |
|
46 |
- |
|
46 |
+#Exp=quote({a1 = A[ind2, ]})[[2]] |
|
47 | 47 |
C_element_OP <- function(varInfo, Exp) { |
48 | 48 |
leftExp = Exp[[2]] |
49 | 49 |
rightExp = Exp[[3]] |
50 | 50 |
|
51 | 51 |
leftInfo = getVarInfo(varInfo, leftExp) |
52 | 52 |
if (leftInfo$dataType == T_scale) { |
53 |
- sub = c("1", "1") |
|
53 |
+ sub = c("0", "0") |
|
54 | 54 |
leftElement = C_element_getCExp(varInfo, leftExp, sub = sub, opt = NULL) |
55 | 55 |
rightElement = C_element_getCExp(varInfo, rightExp, sub = sub, extCode = leftElement$extCode, opt = NULL) |
56 | 56 |
} else { |
... | ... |
@@ -71,11 +71,22 @@ C_element_OP <- function(varInfo, Exp) { |
71 | 71 |
code = c("{", extCode, assignmentCode, "}") |
72 | 72 |
} |
73 | 73 |
} else { |
74 |
+ if(is.null(extCode$L0)){ |
|
74 | 75 |
code = C_matrix_assignment( |
75 | 76 |
assignmentCode, |
76 | 77 |
loopInd1 = "gpu_element_j", loopEnd1 = R_ncol(varInfo, leftExp), |
77 | 78 |
loopInd2 = "gpu_element_i", loopEnd2 = R_nrow(varInfo, leftExp), |
78 |
- loopCode0 = extCode$L0, loopCode1 = extCode$L1, loopCode2 = extCode$L2) |
|
79 |
+ loopCode0 = NULL, loopCode1 = extCode$L1, loopCode2 = extCode$L2) |
|
80 |
+ }else{ |
|
81 |
+ code = c( |
|
82 |
+ "{", |
|
83 |
+ C_matrix_assignment( |
|
84 |
+ assignmentCode, |
|
85 |
+ loopInd1 = "gpu_element_j", loopEnd1 = R_ncol(varInfo, leftExp), |
|
86 |
+ loopInd2 = "gpu_element_i", loopEnd2 = R_nrow(varInfo, leftExp), |
|
87 |
+ loopCode0 = extCode$L0, loopCode1 = extCode$L1, loopCode2 = extCode$L2), |
|
88 |
+ "}") |
|
89 |
+ } |
|
79 | 90 |
} |
80 | 91 |
return(code) |
81 | 92 |
} |
... | ... |
@@ -172,7 +183,7 @@ C_element_abs<-function(varInfo, Exp, sub, opt, extCode){ |
172 | 183 |
return(res) |
173 | 184 |
} |
174 | 185 |
processSub<-function(varInfo, Exp, sub, opt, extCode){ |
175 |
- if(deparse(Exp)==""){ |
|
186 |
+ if(Exp==""){ |
|
176 | 187 |
res = list(value = sub, extCode = extCode) |
177 | 188 |
return(res) |
178 | 189 |
} |
... | ... |
@@ -182,16 +193,32 @@ processSub<-function(varInfo, Exp, sub, opt, extCode){ |
182 | 193 |
} |
183 | 194 |
|
184 | 195 |
C_element_sub<-function(varInfo, Exp, sub, opt, extCode){ |
196 |
+ args=matchBracketFunc(Exp) |
|
185 | 197 |
targetExp=Exp[[2]] |
186 |
- sub1=processSub(varInfo, Exp[[3]],sub[1],extCode = extCode, opt = opt) |
|
187 |
- |
|
198 |
+ sub1=processSub(varInfo, args$i,sub[1],extCode = extCode, opt = opt) |
|
188 | 199 |
|
189 |
- if (length(sub) == 1) { |
|
190 |
- res = C_element_getCExp(varInfo, targetExp,sub=sub1$value, extCode = sub1$extCode, opt = opt) |
|
191 |
- } else { |
|
192 |
- sub2=processSub(varInfo, Exp[[4]],sub[2],extCode = sub1$extCode, opt = opt) |
|
193 |
- res = C_element_getCExp(varInfo, targetExp,sub=c(sub1$value,sub2$value), extCode = sub2$extCode, opt = opt) |
|
200 |
+ if(length(sub) == 1){ |
|
201 |
+ #two index bracket -- one index sub |
|
202 |
+ if(!is.null(args$j)){ |
|
203 |
+ rowNum = R_nrow(varInfo,Exp) |
|
204 |
+ res_twoIndex=one_to_two_index(sub1$value,extCode = sub1$extCode,rowNum=rowNum) |
|
205 |
+ sub_new=c(res$i,res$j) |
|
206 |
+ res = C_element_getCExp(varInfo, Exp, sub=sub_new, extCode = res$extCode, opt = opt) |
|
207 |
+ }else{ |
|
208 |
+ #one index bracket -- one index sub |
|
209 |
+ res = C_element_getCExp(varInfo, targetExp,sub=sub1$value, extCode = sub1$extCode, opt = opt) |
|
210 |
+ } |
|
211 |
+ }else{ |
|
212 |
+ if(!is.null(args$j)){ |
|
213 |
+ sub2=processSub(varInfo, args$j,sub[2],extCode = sub1$extCode, opt = opt) |
|
214 |
+ #two index bracket -- two index sub |
|
215 |
+ res = C_element_getCExp(varInfo, targetExp,sub=c(sub1$value,sub2$value), extCode = sub2$extCode, opt = opt) |
|
216 |
+ }else{ |
|
217 |
+ #one index bracket -- two index sub |
|
218 |
+ res = C_element_getCExp(varInfo, targetExp,sub=sub1$value, extCode = sub1$extCode, opt = opt) |
|
219 |
+ } |
|
194 | 220 |
} |
221 |
+ |
|
195 | 222 |
return(res) |
196 | 223 |
} |
197 | 224 |
|
... | ... |
@@ -535,6 +562,20 @@ C_rowSums_right<-function(varInfo, Exp){ |
535 | 562 |
code |
536 | 563 |
} |
537 | 564 |
|
565 |
+#Exp=quote(compiler.define(A,B)) |
|
566 |
+C_compiler_define<-function(varInfo, Exp){ |
|
567 |
+ vars=extractVars(Exp) |
|
568 |
+ code=c() |
|
569 |
+ for(var in vars){ |
|
570 |
+ curInfo=getVarInfo(varInfo,var) |
|
571 |
+ ind=which(varInfo$var_def_code$var==var) |
|
572 |
+ if(length(ind)==0) |
|
573 |
+ stop("Unable to find the variable ",var," in the variable define table.\n Current expression: ",deparse(Exp)) |
|
574 |
+ varDef=varInfo$var_def_code$def[ind] |
|
575 |
+ code=c(code,varDef) |
|
576 |
+ } |
|
577 |
+ code |
|
578 |
+} |
|
538 | 579 |
|
539 | 580 |
|
540 | 581 |
########################################## Super lengthy function############################## |
... | ... |
@@ -1,8 +1,9 @@ |
1 |
-# C_element_getCExp(varInfo,quote(A[ind,]),c(3,6)-1,opt=NULL) |
|
1 |
+# C_element_getCExp(varInfo,quote((At+A)[ind,]),c(3,6)-1,opt=NULL) |
|
2 | 2 |
|
3 | 3 |
R_expression_sub <- function(varInfo, Exp, sub, opt = NULL, extCode = NULL) { |
4 | 4 |
# fill the vector with the first element to make the length of it |
5 | 5 |
# consistant with the length of Exp |
6 |
+ |
|
6 | 7 |
curInfo = getVarInfo(varInfo, Exp) |
7 | 8 |
if (curInfo$isRef) { |
8 | 9 |
refExp = parse(text = curInfo$specialContent)[[1]] |
... | ... |
@@ -47,33 +48,38 @@ R_expression_sub <- function(varInfo, Exp, sub, opt = NULL, extCode = NULL) { |
47 | 48 |
# Expression should be a variable or a matrix subset |
48 | 49 |
# R_oneIndex_exp_sub(varInfo,quote(A[tmp1]),3) |
49 | 50 |
R_oneIndex_exp_sub <- function(varInfo, Exp, k, opt = NULL, extCode = NULL) { |
50 |
- index_type=GPUVar$default_index_type |
|
51 |
+ |
|
51 | 52 |
# If the expression is a variable |
52 | 53 |
curInfo = getVarInfo(varInfo, Exp) |
53 | 54 |
rowNum = R_nrow(varInfo,Exp) |
54 |
- |
|
55 |
- col_ind_var = GPUVar$getTmpVar() |
|
56 |
- col_ind_value = CSimplify(paste0("(", index_type, ")((", |
|
57 |
- k, ")/(", rowNum, "))")) |
|
58 |
- # if the temporary variable is a constant, it will be plug into the |
|
59 |
- # code Otherwise, create a temporary variable for it |
|
60 |
- if (isNumeric(col_ind_value) || length(grep("/", col_ind_value, fixed = TRUE)) == |
|
61 |
- 0) { |
|
62 |
- col_ind_var = paste0("(", col_ind_value, ")") |
|
63 |
- } else { |
|
64 |
- res = getVarFromExtCode(extCode, index_type, col_ind_value) |
|
65 |
- col_ind_var = res$var |
|
66 |
- extCode = res$extCode |
|
67 |
- } |
|
68 |
- |
|
69 |
- i = paste0(k, "-", rowNum, "*", col_ind_var) |
|
70 |
- j = col_ind_var |
|
71 |
- |
|
72 |
- res = R_getVarSub(varInfo, Exp, i = i, j = j, opt = opt, extCode = extCode) |
|
73 |
- |
|
55 |
+ res_index=one_to_two_index(k,extCode = extCode,rowNum=rowNum) |
|
56 |
+ res = R_getVarSub(varInfo, Exp, i = res_index$i, j = res_index$j, opt = opt, extCode = res_index$extCode) |
|
74 | 57 |
return(res) |
75 | 58 |
} |
76 | 59 |
|
60 |
+#Return: i, j, extCode |
|
61 |
+one_to_two_index<-function(k,extCode,rowNum){ |
|
62 |
+ |
|
63 |
+ index_type=GPUVar$default_index_type |
|
64 |
+ col_ind_var = GPUVar$getTmpVar() |
|
65 |
+ col_ind_value = CSimplify(paste0("(", index_type, ")((", |
|
66 |
+ k, ")/(", rowNum, "))")) |
|
67 |
+ # if the temporary variable is a constant, it will be plug into the |
|
68 |
+ # code Otherwise, create a temporary variable for it |
|
69 |
+ if (isNumeric(col_ind_value) || length(grep("/", col_ind_value, fixed = TRUE)) == |
|
70 |
+ 0) { |
|
71 |
+ col_ind_var = col_ind_value |
|
72 |
+ } else { |
|
73 |
+ res = getVarFromExtCode(extCode, index_type, col_ind_value) |
|
74 |
+ col_ind_var = res$var |
|
75 |
+ extCode = res$extCode |
|
76 |
+ } |
|
77 |
+ |
|
78 |
+ i = CSimplify(paste0(k, "-", rowNum, "*(", col_ind_var,")")) |
|
79 |
+ j = CSimplify(col_ind_var) |
|
80 |
+ |
|
81 |
+ return(list(i=i,j=j,extCode=extCode)) |
|
82 |
+} |
|
77 | 83 |
# Get an element from the matrix(eg. A[i,j]), the transpose will be |
78 | 84 |
# taken into account |
79 | 85 |
# i,j are either a number or a variable in C code , they are 0-based index |
... | ... |
@@ -95,19 +101,24 @@ R_getVarSub <- function(varInfo, var, i, j = 1, opt = NULL, extCode = NULL) { |
95 | 101 |
#If the data is not a pointer, it is a scalar, then directly return the address |
96 | 102 |
if(is.na(isPointer)) stop("Unable to determine the type of the variable,", |
97 | 103 |
" this is a bug in the package, please contact the author.") |
98 |
- if(!isPointer) |
|
99 |
- return(list(value =address,extCode=extCode)) |
|
100 |
- #If the data is a pointer, but the size is 1, then return the first element in the address |
|
101 |
- if((curInfo$size1=="1"&&curInfo$size2=="1")||curInfo$dataType==T_scale){ |
|
102 |
- if(!curInfo$isSeq){ |
|
103 |
- return(list(value =paste0(address,"[0]"),extCode=extCode)) |
|
104 |
- }else{ |
|
105 |
- seqInfo = getSeqAddress(varInfo, var) |
|
106 |
- return(list(value =seqInfo$from,extCode=extCode)) |
|
104 |
+ |
|
105 |
+ |
|
106 |
+ |
|
107 |
+ if(!isPointer) { |
|
108 |
+ if(!curInfo$isSeq) |
|
109 |
+ return(list(value =address,extCode=extCode)) |
|
110 |
+ }else{ |
|
111 |
+ #If the data is a pointer, but the size is 1, then return the first element in the address |
|
112 |
+ if((curInfo$size1=="1"&&curInfo$size2=="1")||curInfo$dataType==T_scale){ |
|
113 |
+ if(!curInfo$isSeq){ |
|
114 |
+ return(list(value =paste0(address,"[0]"),extCode=extCode)) |
|
115 |
+ }else{ |
|
116 |
+ seqInfo = getSeqAddress(varInfo, var) |
|
117 |
+ return(list(value =seqInfo$from,extCode=extCode)) |
|
118 |
+ } |
|
107 | 119 |
} |
108 | 120 |
} |
109 | 121 |
|
110 |
- |
|
111 | 122 |
# compute the matrix offset |
112 | 123 |
size1 = R_getVarSize(varInfo, var,C_symbol = FALSE,ind=1,processTranspose = FALSE) |
113 | 124 |
|
... | ... |
@@ -122,7 +133,6 @@ R_getVarSub <- function(varInfo, var, i, j = 1, opt = NULL, extCode = NULL) { |
122 | 133 |
|
123 | 134 |
if (curInfo$isSeq) { |
124 | 135 |
seqInfo = getSeqAddress(varInfo, var) |
125 |
- |
|
126 | 136 |
res = list(extCode=extCode) |
127 | 137 |
res$value = CSimplify(paste0(seqInfo$from, "+", "(", rowOffset, |
128 | 138 |
"-1)*", seqInfo$by)) |
... | ... |
@@ -185,7 +195,7 @@ R_getVarSize <- function(varInfo, var,C_symbol, ind,processTranspose=TRUE) { |
185 | 195 |
return(res) |
186 | 196 |
} |
187 | 197 |
if (curInfo$isSeq) { |
188 |
- res=ifelse(ind==1,getSizeVar(var_char,ind),1) |
|
198 |
+ res=ifelse(ind==1,getSizeVar(var_char,1),1) |
|
189 | 199 |
return(res) |
190 | 200 |
} |
191 | 201 |
if (curInfo$dataType == T_scale) |
... | ... |
@@ -51,6 +51,11 @@ addVariableDeclaration <- function(varInfo,curInfo, data, offset, offsetInd,pref |
51 | 51 |
location = paste0(prefix, " ") |
52 | 52 |
curInfo$address = curInfo$var |
53 | 53 |
|
54 |
+ if(curInfo$redirect!="NA"){ |
|
55 |
+ redirectInfo=getVarInfo(varInfo,curInfo$redirect) |
|
56 |
+ curCode = paste0("#define ",curInfo$var," ",redirectInfo$address) |
|
57 |
+ return(list(code=curCode,Info=curInfo)) |
|
58 |
+ } |
|
54 | 59 |
|
55 | 60 |
# If the variable is a sequence |
56 | 61 |
if (curInfo$isSeq) { |
... | ... |
@@ -59,7 +64,7 @@ addVariableDeclaration <- function(varInfo,curInfo, data, offset, offsetInd,pref |
59 | 64 |
} else { |
60 | 65 |
curCode = paste0(location, CXXtype, 3, "* ", curInfo$var, |
61 | 66 |
"=", "((", location, CXXtype, 3, "*)(", data, "+", offset, |
62 |
- "[", offsetInd, "]))+",worker_offset,";") |
|
67 |
+ "(", offsetInd, ")))+",worker_offset,";") |
|
63 | 68 |
} |
64 | 69 |
return(list(code=curCode,Info=curInfo)) |
65 | 70 |
} |
... | ... |
@@ -73,12 +78,31 @@ addVariableDeclaration <- function(varInfo,curInfo, data, offset, offsetInd,pref |
73 | 78 |
} |
74 | 79 |
}else{ |
75 | 80 |
curCode = paste0(location, CXXtype, "* ", curInfo$var, "=", |
76 |
- "((", location, CXXtype, "*)(", data, "+", offset, "[", offsetInd, |
|
77 |
- "]))+",worker_offset,";") |
|
81 |
+ "((", location, CXXtype, "*)(", data, "+", offset, "(", offsetInd, |
|
82 |
+ ")))+",worker_offset,";") |
|
83 |
+ } |
|
84 |
+ return(list(code=curCode,Info=curInfo)) |
|
85 |
+} |
|
86 |
+ |
|
87 |
+addVariableDeclaration_NonPointer<-function(varInfo,curInfo){ |
|
88 |
+ curVar=curInfo$var |
|
89 |
+ CXXtype = getTypeCXXStr(curInfo$precisionType) |
|
90 |
+ curInfo$address=curVar |
|
91 |
+ if(curInfo$redirect!="NA"){ |
|
92 |
+ redirectInfo=getVarInfo(varInfo,curInfo$redirect) |
|
93 |
+ curCode = paste0("#define ",curVar," ",redirectInfo$address) |
|
94 |
+ return(list(code=curCode,Info=curInfo)) |
|
95 |
+ } |
|
96 |
+ |
|
97 |
+ if (curInfo$isSeq){ |
|
98 |
+ curCode = paste0(CXXtype, 3, " ", curInfo$var, ";") |
|
99 |
+ }else{ |
|
100 |
+ curCode = paste0(CXXtype, " ", curVar, ";") |
|
78 | 101 |
} |
79 | 102 |
return(list(code=curCode,Info=curInfo)) |
80 | 103 |
} |
81 | 104 |
|
105 |
+ |
|
82 | 106 |
getSizeVar<-function(varName,i){ |
83 | 107 |
paste0(GPUVar$matrix_size_prefix, varName,"_dim_",i) |
84 | 108 |
} |
... | ... |
@@ -208,7 +232,7 @@ getSeqAddress <- function(varInfo, var,C_symbol=FALSE) { |
208 | 232 |
from = paste0(ad, ".s0") |
209 | 233 |
to = paste0(ad, ".s1") |
210 | 234 |
by = paste0(ad, ".s2") |
211 |
- length = getSizeVar(deparse(var),ind) |
|
235 |
+ length = getSizeVar(deparse(var),1) |
|
212 | 236 |
data.frame(from = from, to = to, by = by, length = length, stringsAsFactors = FALSE) |
213 | 237 |
} |
214 | 238 |
|
... | ... |
@@ -241,8 +265,7 @@ C_general_scalar_assignment <- function(varInfo, Exp, funcName, func) { |
241 | 265 |
stop("Unexpected function:", deparse(Exp)) |
242 | 266 |
} |
243 | 267 |
} else { |
244 |
- res = R_expression_sub(varInfo, leftExp, sub = 1, sub_C = TRUE, |
|
245 |
- extCode = extCode) |
|
268 |
+ res = C_element_getCExp(varInfo, leftExp, sub = 0, extCode = extCode) |
|
246 | 269 |
value_left = res$value |
247 | 270 |
extCode = res$extCode |
248 | 271 |
} |
... | ... |
@@ -257,8 +280,7 @@ C_general_scalar_assignment <- function(varInfo, Exp, funcName, func) { |
257 | 280 |
stop("Unexpected function:", deparse(Exp)) |
258 | 281 |
} |
259 | 282 |
} else { |
260 |
- res = R_expression_sub(varInfo, rightExp, sub = 1, sub_C = TRUE, |
|
261 |
- extCode = extCode) |
|
283 |
+ res = C_element_getCExp(varInfo, rightExp, sub = 0, extCode = extCode) |
|
262 | 284 |
value_right = res$value |
263 | 285 |
extCode = res$extCode |
264 | 286 |
} |
... | ... |
@@ -1,6 +1,7 @@ |
1 | 1 |
ROptimizer1 <- function(profileMeta2) { |
2 | 2 |
previousExp=profileMeta2$Exp |
3 | 3 |
varInfo=profileMeta2$varInfo |
4 |
+ #Find the start and end line number of the in used variables |
|
4 | 5 |
#a vector of 2: start,end |
5 | 6 |
varUsedInfo=list() |
6 | 7 |
for(i in seq_along(previousExp)){ |
... | ... |
@@ -15,12 +16,12 @@ ROptimizer1 <- function(profileMeta2) { |
15 | 16 |
} |
16 | 17 |
} |
17 | 18 |
} |
18 |
- |
|
19 |
+ #Remove the global variable |
|
19 | 20 |
varUsedInfo[[GPUVar$gpu_global_id]]=NULL |
20 | 21 |
varUsedInfoTbl=do.call("rbind", varUsedInfo) |
21 | 22 |
|
22 | 23 |
vars=row.names(varUsedInfoTbl) |
23 |
- |
|
24 |
+ #Add compiler.define and compiler.release information in the code |
|
24 | 25 |
Exp=c() |
25 | 26 |
for(i in rev(seq_along(previousExp))){ |
26 | 27 |
curReleasedVarInd=which(varUsedInfoTbl[,2]==i) |
... | ... |
@@ -53,36 +54,143 @@ ROptimizer1 <- function(profileMeta2) { |
53 | 54 |
} |
54 | 55 |
|
55 | 56 |
|
56 |
- |
|
57 |
-variableInUsed<-function(varInfo,vars){ |
|
58 |
- res=lapply(vars,findRootVar,varInfo=varInfo) |
|
59 |
- unlist(res) |
|
57 |
+ROptimizer2 <- function(GPUExp2) { |
|
58 |
+ previousCode=GPUExp2$gpu_code |
|
59 |
+ #Find the start and end line number of the var dimension |
|
60 |
+ varRecord=getVarDimRecord(previousCode) |
|
61 |
+ |
|
62 |
+ insertCode=getDimCode(GPUExp2$varInfo,varRecord) |
|
63 |
+ code=previousCode |
|
64 |
+ for(i in rev(as.numeric(names(insertCode)))){ |
|
65 |
+ pre_len=i-1 |
|
66 |
+ post_len=length(code)-i |
|
67 |
+ code=c(code[seq_len(pre_len)],insertCode[[as.character(i)]],code[seq_len(post_len)+i]) |
|
68 |
+ } |
|
69 |
+ code=code[-which(code=="//Main function delimiter")] |
|
70 |
+ #realize the promise assign |
|
71 |
+ code=realizePromiseAssign(code,varRecord) |
|
72 |
+ |
|
73 |
+ GPUExp3=GPUExp2 |
|
74 |
+ GPUExp3$gpu_code=code |
|
75 |
+ GPUExp3 |
|
60 | 76 |
} |
61 | 77 |
|
62 |
-findRootVar<-function(varInfo,Exp){ |
|
63 |
- Exp=toExpression(Exp) |
|
64 |
- vars=extractVars(Exp) |
|
65 |
- root=c() |
|
66 |
- for(i in vars){ |
|
67 |
- if(is.call(i)){ |
|
68 |
- rootc(root,findRootVar(varInfo,i)) |
|
78 |
+ |
|
79 |
+getVarDimRecord<-function(previousCode){ |
|
80 |
+ varRecord=c() |
|
81 |
+ mainLineNum=0 |
|
82 |
+ pattern=paste0(GPUVar$matrix_size_prefix,"(.+?)_dim_([12])") |
|
83 |
+ for(i in seq_along(previousCode)){ |
|
84 |
+ curCode=previousCode[[i]] |
|
85 |
+ if(curCode=="//Main function delimiter"){ |
|
86 |
+ mainLineNum=i |
|
69 | 87 |
next |
70 | 88 |
} |
71 |
- curvar=toCharacter(i) |
|
72 |
- if(hasVar(varInfo,curvar)){ |
|
73 |
- curInfo=getVarInfo(varInfo,curvar) |
|
74 |
- if(curInfo$redirect!="NA"){ |
|
75 |
- root=c(root,findRootVar(varInfo,curInfo$redirect)) |
|
89 |
+ #If it is a promise assign, skip it |
|
90 |
+ if(length(grep(GPUVar$promiseAssgin,curCode,fixed = TRUE))!=0) |
|
91 |
+ next |
|
92 |
+ #find the first main function delimiter |
|
93 |
+ if(mainLineNum==0) |
|
94 |
+ next |
|
95 |
+ #extract the variable dimension variables |
|
96 |
+ dim_vars=unique(str_extract_all(curCode,pattern)[[1]]) |
|
97 |
+ #Find the target variable |
|
98 |
+ vars=gsub(pattern,"\\1",dim_vars) |
|
99 |
+ vars_dim_ind=gsub(pattern,"\\2",dim_vars) |
|
100 |
+ for(j in seq_along(vars)){ |
|
101 |
+ var=vars[j] |
|
102 |
+ var_dim_ind=vars_dim_ind[j] |
|
103 |
+ record_Ind=which(varRecord[,1]==var&varRecord[,2]==var_dim_ind) |
|
104 |
+ if(length(record_Ind)==1){ |
|
105 |
+ varRecord[record_Ind,4]=i |
|
76 | 106 |
}else{ |
77 |
- if(curInfo$specialType=="ref"){ |
|
78 |
- root=c(root,findRootVar(varInfo,curInfo$specialContent)) |
|
79 |
- }else{ |
|
80 |
- if(!curInfo$shared){ |
|
81 |
- root=c(root,curInfo$var) |
|
82 |
- } |
|
83 |
- } |
|
107 |
+ varRecord=rbind(varRecord,c(var,var_dim_ind,i,i,mainLineNum)) |
|
84 | 108 |
} |
85 | 109 |
} |
86 | 110 |
} |
87 |
- root |
|
88 |
-} |
|
89 | 111 |
\ No newline at end of file |
112 |
+ varRecord=as.data.frame(varRecord,stringsAsFactors=FALSE) |
|
113 |
+ colnames(varRecord)=c("var","dim","start","end","mainLineNum") |
|
114 |
+ varRecord |
|
115 |
+} |
|
116 |
+ |
|
117 |
+getDimCode<-function(varInfo,varRecord){ |
|
118 |
+ insertCode=list() |
|
119 |
+ #tmpname,varName |
|
120 |
+ temp_dim_record=c() |
|
121 |
+ ind=unique(sort(c(varRecord$start))) |
|
122 |
+ matrix_temporary_size=GPUVar$matrix_temporary_size |
|
123 |
+ for(i in seq_len(nrow(varRecord))){ |
|
124 |
+ start_line=varRecord$start[i] |
|
125 |
+ #Remove the unused dim |
|
126 |
+ unused_ind=which(varRecord$end<start_line) |
|
127 |
+ unused_var=varRecord$var[unused_ind] |
|
128 |
+ unused_var_dim=varRecord$dim[unused_ind] |
|
129 |
+ var_dim=getSizeVar(unused_var,unused_var_dim) |
|
130 |
+ temp_dim_record[temp_dim_record%in%var_dim]=NA |
|
131 |
+ |
|
132 |
+ #Assign the size variable a temporary dim variable |
|
133 |
+ mainLineInd=varRecord$mainLineNum[i] |
|
134 |
+ curVar=varRecord$var[i] |
|
135 |
+ curVar_dim=varRecord$dim[i] |
|
136 |
+ curVar_dim_char=getSizeVar(curVar,curVar_dim) |
|
137 |
+ #Check if there is enough temprary variables |
|
138 |
+ #If not, add more temporary variables |
|
139 |
+ #If it is enough, assign the temporary variables to the dim variables |
|
140 |
+ temp_Var_def=NULL |
|
141 |
+ tmpVar_available_size=sum(is.na(temp_dim_record)) |
|
142 |
+ if(tmpVar_available_size==0){ |
|
143 |
+ temp_dim_record=c(temp_dim_record,NA) |
|
144 |
+ temp_Var_def=paste0(GPUVar$default_index_type," ",matrix_temporary_size,length(temp_dim_record),";") |
|
145 |
+ } |
|
146 |
+ tmpVar_available_ind=which(is.na(temp_dim_record)) |
|
147 |
+ tmp_ind=tmpVar_available_ind[1] |
|
148 |
+ temp_dim_record[tmp_ind]=curVar_dim_char |
|
149 |
+ #Define the size macro |
|
150 |
+ dim_macro=paste0("#define ",curVar_dim_char," ",matrix_temporary_size,tmp_ind) |
|
151 |
+ #Find the size data |
|
152 |
+ dim_data=NULL |
|
153 |
+ if(has.key(curVar,varInfo$matrixInd)){ |
|
154 |
+ curInfo=getVarInfo(varInfo,curVar) |
|
155 |
+ if(curInfo$location=="global"&&curInfo$shared){ |
|
156 |
+ size_ad=paste0(GPUVar$gs_size,curVar_dim) |
|
157 |
+ } |
|
158 |
+ if(curInfo$location=="global"&&!curInfo$shared){ |
|
159 |
+ size_ad=paste0(GPUVar$gp_size,curVar_dim) |
|
160 |
+ } |
|
161 |
+ if(curInfo$location=="local"&&curInfo$shared){ |
|
162 |
+ size_ad=paste0(GPUVar$ls_size,curVar_dim) |
|
163 |
+ } |
|
164 |
+ if(curInfo$location=="local"&&!curInfo$shared){ |
|
165 |
+ size_ad=paste0(GPUVar$lp_size,curVar_dim) |
|
166 |
+ } |
|
167 |
+ dim_data=paste0(curVar_dim_char,"=",size_ad,"(",varInfo$matrixInd[[curVar]],");") |
|
168 |
+ } |
|
169 |
+ |
|
170 |
+ insertCode[[mainLineInd]]=c(insertCode[[mainLineInd]], |
|
171 |
+ temp_Var_def, |
|
172 |
+ dim_macro, |
|
173 |
+ dim_data |
|
174 |
+ ) |
|
175 |
+ } |
|
176 |
+ insertCode |
|
177 |
+} |
|
178 |
+realizePromiseAssign<-function(code,varRecord){ |
|
179 |
+ vars_dim=getSizeVar(varRecord$var,varRecord$dim) |
|
180 |
+ vars_dim_group=paste0(vars_dim,collapse = "|") |
|
181 |
+ vars_dim_header=GPUVar$promiseAssgin |
|
182 |
+ vars_dim_pattern=paste0(vars_dim_header,"(",vars_dim_group,")") |
|
183 |
+ remove_ind=NULL |
|
184 |
+ for(i in seq_along(code)){ |
|
185 |
+ curCode=code[i] |
|
186 |
+ if(length(grep(vars_dim_header,curCode))>0){ |
|
187 |
+ if(length(grep(vars_dim_pattern,curCode))>0){ |
|
188 |
+ code[i]=substr(curCode,nchar(vars_dim_header)+1,nchar(curCode)) |
|
189 |
+ }else{ |
|
190 |
+ remove_ind=c(remove_ind,i) |
|
191 |
+ } |
|
192 |
+ } |
|
193 |
+ } |
|
194 |
+ if(!is.null(remove_ind)) |
|
195 |
+ code=code[-remove_ind] |
|
196 |
+ return(code) |
|
197 |
+} |
... | ... |
@@ -0,0 +1,36 @@ |
1 |
+#######################Optimizer 1############################# |
|
2 |
+#Find which variables are in used |
|
3 |
+#Can handle redirect and subref |
|
4 |
+variableInUsed<-function(varInfo,vars){ |
|
5 |
+ res=lapply(vars,findRootVar,varInfo=varInfo) |
|
6 |
+ unlist(res) |
|
7 |
+} |
|
8 |
+#Find the in used variables for the current expression |
|
9 |
+findRootVar<-function(varInfo,Exp){ |
|
10 |
+ Exp=toExpression(Exp) |
|
11 |
+ vars=extractVars(Exp) |
|
12 |
+ root=c() |
|
13 |
+ for(i in vars){ |
|
14 |
+ if(is.call(i)){ |
|
15 |
+ rootc(root,findRootVar(varInfo,i)) |
|
16 |
+ next |
|
17 |
+ } |
|
18 |
+ curvar=toCharacter(i) |
|
19 |
+ if(hasVar(varInfo,curvar)){ |
|
20 |
+ curInfo=getVarInfo(varInfo,curvar) |
|
21 |
+ if(curInfo$redirect!="NA"){ |
|
22 |
+ root=c(root,findRootVar(varInfo,curInfo$redirect)) |
|
23 |
+ }else{ |
|
24 |
+ if(curInfo$specialType=="ref"){ |
|
25 |
+ root=c(root,findRootVar(varInfo,curInfo$specialContent)) |
|
26 |
+ }else{ |
|
27 |
+ if(!curInfo$shared){ |
|
28 |
+ root=c(root,curInfo$var) |
|
29 |
+ } |
|
30 |
+ } |
|
31 |
+ } |
|
32 |
+ } |
|
33 |
+ } |
|
34 |
+ root |
|
35 |
+} |
|
36 |
+#######################Optimizer 2############################# |
|
0 | 37 |
\ No newline at end of file |
... | ... |
@@ -175,6 +175,7 @@ profiler_assignment_exitingVar <- function(level, varInfo, curExp) { |
175 | 175 |
# curExp=quote({D=1:gpu_global_id})[[2]] |
176 | 176 |
profiler_assignment_newVar <- function(level, varInfo, curExp) { |
177 | 177 |
result=list(Exp=curExp,varInfo=varInfo) |
178 |
+ |
|
178 | 179 |
rightInfoPack = getExpInfo(varInfo, curExp[[3]]) |
179 | 180 |
rightInfo = rightInfoPack$ExpInfo |
180 | 181 |
|
... | ... |
@@ -695,8 +696,8 @@ profile_transpose <- function(varInfo, Exp) { |
695 | 696 |
# Exp=quote(t_nocpy(A)) |
696 | 697 |
profile_transpose_nocpy <- function(varInfo, Exp) { |
697 | 698 |
curVar = Exp[[2]] |
698 |
- Exp = parse(text = paste0("subRef(", curVar, ",,)"))[[1]] |
|
699 |
- info = getExpInfo(varInfo, Exp) |
|
699 |
+ Exp_new = parse(text = paste0("subRef(", curVar, ",,)"))[[1]] |
|
700 |
+ info = getExpInfo(varInfo, Exp_new) |
|
700 | 701 |
ExpInfo = info$ExpInfo |
701 | 702 |
size1 = ExpInfo$size2 |
702 | 703 |
size2 = ExpInfo$size1 |
... | ... |
@@ -704,6 +705,7 @@ profile_transpose_nocpy <- function(varInfo, Exp) { |
704 | 705 |
ExpInfo$size2 = size2 |
705 | 706 |
ExpInfo$transpose = !ExpInfo$transpose |
706 | 707 |
info$ExpInfo = ExpInfo |
708 |
+ info$Exp=Exp |
|
707 | 709 |
|
708 | 710 |
return(info) |
709 | 711 |
} |
... | ... |
@@ -142,6 +142,7 @@ getVarsNum <- function(extCode) { |
142 | 142 |
hoistOpt <- function(extCode, Exp) { |
143 | 143 |
code = C_to_R(Exp) |
144 | 144 |
code = vapply(expandExp(code), Simplify,character(1)) |
145 |
+ #code=expandExp(code) |
|
145 | 146 |
codeInfo = list() |
146 | 147 |
baseLevel = c() |
147 | 148 |
# Decompose the code and find the base level for each code |
... | ... |
@@ -226,9 +227,10 @@ constructCode <- function(codeInfo, level) { |
226 | 227 |
} |
227 | 228 |
|
228 | 229 |
# Decompose the code into different level The code should not be able |
229 |
-# to separate by +,- operator The current supported decompose function |
|
230 |
-# is * |
|
230 |
+# to separate by +,- operator |
|
231 |
+# The current supported decompose function is * |
|
231 | 232 |
decomposeCode <- function(extCode, code) { |
233 |
+ code = toExpression(code) |
|
232 | 234 |
code = decomposeCode_hidden(extCode, code) |
233 | 235 |
if (nrow(code) > 1) { |
234 | 236 |
for (i in seq_len(getLevelNum(extCode) - 1)) { |
... | ... |
@@ -241,14 +243,14 @@ decomposeCode <- function(extCode, code) { |
241 | 243 |
code |
242 | 244 |
} |
243 | 245 |
decomposeCode_hidden <- function(extCode, code, operator = "") { |
244 |
- code = toExpression(code) |
|
245 |
- if (is.call(code)) { |
|
246 |
+ |
|
247 |
+ if (is.call(code)&&operator!="/") { |
|
246 | 248 |
func = deparse(code[[1]]) |
247 | 249 |
if (func == "*") { |
248 |
- left = decomposeCode_hidden(extCode, code[[2]]) |
|
249 |
- right = rbind(left, decomposeCode_hidden(extCode, code[[3]], |
|
250 |
- operator = func)) |
|
251 |
- return(right) |
|
250 |
+ left = decomposeCode_hidden(extCode, code[[2]],operator=operator) |
|
251 |
+ right = decomposeCode_hidden(extCode, code[[3]],operator = func) |
|
252 |
+ res=rbind(left,right) |
|
253 |
+ return(res) |
|
252 | 254 |
} |
253 | 255 |
if (func == "-") { |
254 | 256 |
res = decomposeCode_hidden(extCode, code[[2]]) |
... | ... |
@@ -256,7 +258,7 @@ decomposeCode_hidden <- function(extCode, code, operator = "") { |
256 | 258 |
return(res) |
257 | 259 |
} |
258 | 260 |
if (func == "(") { |
259 |
- res = decomposeCode_hidden(extCode, code[[2]]) |
|
261 |
+ res = decomposeCode_hidden(extCode, code[[2]],operator=operator) |
|
260 | 262 |
return(res) |
261 | 263 |
} |
262 | 264 |
} |
... | ... |
@@ -278,7 +280,7 @@ findCodeLevel <- function(extCode, code) { |
278 | 280 |
level = findVarLevel(extCode, vars) |
279 | 281 |
return(level) |
280 | 282 |
} |
281 |
- |
|
283 |
+#code="gpu_element_dist * (10 * gpu_element_j + gpu_element_i)" |
|
282 | 284 |
# Expand the parathesis in the expression |
283 | 285 |
expandExp <- function(code) { |
284 | 286 |
code = toExpression(code) |
... | ... |
@@ -99,8 +99,8 @@ fillGPUdata <- function(GPUcode1, .options, .device) { |
99 | 99 |
|
100 | 100 |
kernel_args = list() |
101 | 101 |
|
102 |
- # gp_totalsize, return size |
|
103 |
- kernel_args$sizeInfo = rep(0, 2) |
|
102 |
+ # return size, gp,gs,lp offset, gp,gs,lp,ls number, gp,gs,lp,ls dim(row,col) |
|
103 |
+ kernel_args$sizeInfo = NULL |
|
104 | 104 |
gp_totalsize=0 |
105 | 105 |
returnSize=1 |
106 | 106 |
|
... | ... |
@@ -110,19 +110,19 @@ fillGPUdata <- function(GPUcode1, .options, .device) { |
110 | 110 |
|
111 | 111 |
sizeInfo_gp = getVarSizeInfo_C_level(varInfo$matrix_gp) |
112 | 112 |
# Total size per worker |
113 |
- gp_totalsize = sizeInfo_gp$totalSize |
|
114 | 113 |
matrix_size_info=c(matrix_size_info,sizeInfo_gp$dim) |
115 | 114 |
|
116 | 115 |
|
117 | 116 |
sizeInfo_gs = getVarSizeInfo_C_level(varInfo$matrix_gs) |
118 | 117 |
matrix_size_info=c(matrix_size_info,sizeInfo_gs$dim) |
119 | 118 |
|
119 |
+ sizeInfo_lp = getVarSizeInfo_C_level(varInfo$matrix_lp) |
|
120 |
+ matrix_size_info=c(matrix_size_info,sizeInfo_lp$dim) |
|
121 |
+ |
|
120 | 122 |
sizeInfo_ls = getVarSizeInfo_C_level(varInfo$matrix_ls) |
121 | 123 |
matrix_size_info=c(matrix_size_info,sizeInfo_ls$dim) |
122 | 124 |
|
123 |
- sizeInfo_lp = getVarSizeInfo_C_level(varInfo$matrix_lp) |
|
124 |
- matrix_size_info=c(matrix_size_info,sizeInfo_lp$dim) |
|
125 |
- if(length(matrix_size_info)==0) matrix_size_info=0 |
|
125 |
+ |
|
126 | 126 |
|
127 | 127 |
|
128 | 128 |
if (!is.null(varInfo$returnInfo)) { |
... | ... |
@@ -140,9 +140,10 @@ fillGPUdata <- function(GPUcode1, .options, .device) { |
140 | 140 |
} |
141 | 141 |
} |
142 | 142 |
|
143 |
+ matrix_offset=c(sizeInfo_gp$matrixOffset,sizeInfo_gs$matrixOffset,sizeInfo_ls$matrixOffset) |
|
144 |
+ kernel_args$sizeInfo = c(returnSize,matrix_offset,matrix_size_info) |
|
143 | 145 |
|
144 |
- kernel_args$sizeInfo = c(gp_totalsize,returnSize) |
|
145 |
- |
|
146 |
+ if(length(matrix_size_info)==0) matrix_size_info=0 |
|
146 | 147 |
# Allocate the gpu memory |
147 | 148 |
totalWorkerNum = length(parms[[1]]) |
148 | 149 |
IntType = GPUVar$default_index_type |
... | ... |
@@ -155,16 +156,6 @@ fillGPUdata <- function(GPUcode1, .options, .device) { |
155 | 156 |
device_argument$ls_data = kernel.getSharedMem(sizeInfo_ls$totalSize, |
156 | 157 |
type = "char") |
157 | 158 |
|
158 |
- device_argument$gp_offset = gpuMatrix(sizeInfo_gp$matrixOffset, type = IntType, |
|
159 |
- device = .device) |
|
160 |
- device_argument$gs_offset = gpuMatrix(sizeInfo_gs$matrixOffset, type = IntType, |
|
161 |
- device = .device) |
|
162 |
- device_argument$ls_offset = gpuMatrix(sizeInfo_ls$matrixOffset, type = IntType, |
|
163 |
- device = .device) |
|
164 |
- |
|
165 |
- device_argument$matrix_size_info = gpuMatrix(matrix_size_info, type = IntType, |
|
166 |
- device = .device) |
|
167 |
- |
|
168 | 159 |
|
169 | 160 |
# The return size for each thread |
170 | 161 |
device_argument$return_var = gpuEmptMatrix(returnSize, totalWorkerNum, type = GPUVar$default_float, device = .device) |
... | ... |
@@ -11,9 +11,30 @@ GPUVar <- local({ |
11 | 11 |
|
12 | 12 |
|
13 | 13 |
#gpu size prefix |
14 |
- GPUVar_env$matrix_size_prefix="gpu_size_of_" |
|
14 |
+ GPUVar_env$matrix_size_prefix="gpu_" |
|
15 |
+ GPUVar_env$matrix_temporary_size="gpu_matrix_temporary_size" |
|
16 |
+ GPUVar_env$promiseAssgin="//compiler promise assign:" |
|
15 | 17 |
|
16 | 18 |
|
19 |
+ #Matrix number |
|
20 |
+ GPUVar_env$gp_number="gpu_gp_number" |
|
21 |
+ GPUVar_env$gs_number="gpu_gs_number" |
|
22 |
+ GPUVar_env$lp_number="gpu_lp_number" |
|
23 |
+ GPUVar_env$ls_number="gpu_ls_number" |
|
24 |
+ |
|
25 |
+ |
|
26 |
+ #This is the offset to find the data in the function argument |
|
27 |
+ #It is a macro |
|
28 |
+ GPUVar_env$global_private_offset="gpu_gp_offset" |
|
29 |
+ GPUVar_env$global_shared_offset="gpu_gs_offset" |
|
30 |
+ GPUVar_env$local_shared_offset="gpu_ls_offset" |
|
31 |
+ |
|
32 |
+ #The macro that find the gp,gs,lp,ls size |
|
33 |
+ GPUVar_env$gp_size="gpu_gp_size" |
|
34 |
+ GPUVar_env$gs_size="gpu_gs_size" |
|
35 |
+ GPUVar_env$lp_size="gpu_lp_size" |
|
36 |
+ GPUVar_env$ls_size="gpu_ls_size" |
|
37 |
+ |
|
17 | 38 |
|
18 | 39 |
|
19 | 40 |
# matrix size info |
... | ... |
@@ -21,37 +42,22 @@ GPUVar <- local({ |
21 | 42 |
GPUVar_env$global_share_size = "gpu_gs_size_arg" |
22 | 43 |
GPUVar_env$local_share_size = "gpu_ls_size_arg" |
23 | 44 |
|
24 |
- # matrix dimension info |
|
25 |
- # the order is gp_size1,gp_size2,gs~,lp~,ls~ |
|
26 |
- GPUVar_env$matrix_size_info="gpu_matrix_size_info" |
|
27 | 45 |
|
28 | 46 |
# worker private data, loacted in global memory |
29 | 47 |
GPUVar_env$global_private_data = "gpu_gp_data" |
30 | 48 |
GPUVar_env$global_private_totalSize = "gpu_gp_totalSize" |
31 | 49 |
|
32 |
- # Per worker length |
|
33 |
- GPUVar_env$global_private_size1 = "gpu_gp_size1" |
|
34 |
- GPUVar_env$global_private_size2 = "gpu_gp_size2" |
|
35 | 50 |
# Per worker offset |
36 | 51 |
GPUVar_env$global_private_offset = "gpu_gp_offset" |
37 | 52 |
|
38 | 53 |
# worker shared data, located in global memory |
39 | 54 |
GPUVar_env$global_shared_data = "gpu_gs_data" |
40 |
- GPUVar_env$global_shared_size1 = "gpu_gs_size1" |
|
41 |
- GPUVar_env$global_shared_size2 = "gpu_gs_size2" |
|
42 |
- GPUVar_env$global_shared_offset = "gpu_gs_offset" |
|
43 | 55 |
|
44 | 56 |
# worker private data, located in private/local memory |
45 | 57 |
GPUVar_env$local_private_data = "gpu_lp_data" |
46 |
- GPUVar_env$local_private_size1 = "gpu_lp_size1" |
|
47 |
- GPUVar_env$local_private_size2 = "gpu_lp_size2" |
|
48 |
- GPUVar_env$local_private_offset = "gpu_lp_offset" |
|
49 | 58 |
|
50 | 59 |
# worker shared data, located in local memory |
51 | 60 |
GPUVar_env$local_shared_data = "gpu_ls_data" |
52 |
- GPUVar_env$local_shared_size1 = "gpu_ls_size1" |
|
53 |
- GPUVar_env$local_shared_size2 = "gpu_ls_size2" |
|
54 |
- GPUVar_env$local_shared_offset = "gpu_ls_offset" |
|
55 | 61 |
|
56 | 62 |
|
57 | 63 |
# return value |
... | ... |
@@ -220,6 +226,8 @@ GPUVar <- local({ |
220 | 226 |
.cFuncs[["next"]] = C_next |
221 | 227 |
.cFuncs[["message"]] = C_message |
222 | 228 |
.cFuncs[["setVersion"]] = C_setVersion |
229 |
+.cFuncs[["compiler.define"]] = C_compiler_define |
|
230 |
+ |
|
223 | 231 |
|
224 | 232 |
#' @include RCParserFunc_Rlevel.R |
225 | 233 |
.sizeFuncs=list() |
... | ... |
@@ -385,8 +393,8 @@ compiler.define<-function(varName,...){ |
385 | 393 |
|
386 | 394 |
} |
387 | 395 |
#If the variable is in used, then define it. |
388 |
-compiler.promiseDefine<-function(precision,varName,definition){ |
|
389 |
- #paste0(precision," ",varName,"=",definition,";") |
|
396 |
+compiler.promiseDefine<-function(precision,varName){ |
|
397 |
+ #paste0(precision," ",varName;") |
|
390 | 398 |
} |
391 | 399 |
#If the variable is in used, then do the assignment |
392 | 400 |
compiler.promiseAssign<-function(varName,value){ |
... | ... |
@@ -289,7 +289,6 @@ def_var<-function(varInfo,curInfo){ |
289 | 289 |
if(!is.null(rediectVar)){ |
290 | 290 |
curInfo$redirect=rediectVar |
291 | 291 |
curInfo$designSize=0 |
292 |
- curInfo$initial_ad=FALSE |
|
293 | 292 |
varInfo$memPool=markUsed(varInfo$memPool,rediectVar) |
294 | 293 |
} |
295 | 294 |
} |
... | ... |
@@ -329,7 +328,7 @@ markUsed<-function(memPool,varName){ |
329 | 328 |
memPool |
330 | 329 |
} |
331 | 330 |
addVar<-function(memPool,varName,precision,len,location,shared=FALSE,isPointer=NA,specialType="NA"){ |
332 |
- if(hasVar(memPool,varName,precision,len,location,shared,isPointer)){ |
|
331 |
+ if(hasVar(memPool,varName,precision,len,location,shared,isPointer,specialType)){ |
|
333 | 332 |
return(memPool) |
334 | 333 |
} |
335 | 334 |
curInfo=createMemPoolInfo(varName,precision,len,location,shared,isPointer,specialType) |