diff --git a/PFAC/include/PFAC.h b/PFAC/include/PFAC.h index c5b2a59..043acaa 100644 --- a/PFAC/include/PFAC.h +++ b/PFAC/include/PFAC.h @@ -163,7 +163,7 @@ PFAC_status_t PFAC_dumpTransitionTable( PFAC_handle_t handle, FILE *fp ) ; * PFAC_STATUS_INTERNAL_ERROR please report bugs * */ -PFAC_status_t PFAC_readPatternFromFile( PFAC_handle_t handle, char *filename ) ; +PFAC_status_t PFAC_readPatternFromFile( PFAC_handle_t handle, const char *filename ) ; /* * return diff --git a/PFAC/include/PFAC_P.h b/PFAC/include/PFAC_P.h index 5a43ceb..1adcf6d 100644 --- a/PFAC/include/PFAC_P.h +++ b/PFAC/include/PFAC_P.h @@ -202,7 +202,7 @@ int dump_reorderPattern( char** rowPtr, int *patternID_table, int *patternLen_ta * (6) *max_state_num_ptr = estimation (upper bound) of total states in PFAC DFA * */ -PFAC_status_t parsePatternFile( char *patternFileName, char ***rowPtr, char **patternPool, +PFAC_status_t parsePatternFile( const char *patternFileName, char ***rowPtr, char **patternPool, int **patternID_table_ptr, int **patternLen_table_ptr, int *max_state_num_ptr, int *pattern_num_ptr ) ; diff --git a/PFAC/src/PFAC.cpp b/PFAC/src/PFAC.cpp index 6032141..f4136d8 100644 --- a/PFAC/src/PFAC.cpp +++ b/PFAC/src/PFAC.cpp @@ -650,7 +650,7 @@ PFAC_status_t PFAC_createHashTable( PFAC_handle_t handle ) /* * if return status is not PFAC_STATUS_SUCCESS, then all reousrces are free. */ -PFAC_status_t PFAC_readPatternFromFile( PFAC_handle_t handle, char *filename ) +PFAC_status_t PFAC_readPatternFromFile( PFAC_handle_t handle, const char *filename ) { if ( NULL == handle ){ return PFAC_STATUS_INVALID_HANDLE ; diff --git a/PFAC/src/PFAC_reduce_inplace_kernel.cu b/PFAC/src/PFAC_reduce_inplace_kernel.cu index 49e7781..2908e9f 100644 --- a/PFAC/src/PFAC_reduce_inplace_kernel.cu +++ b/PFAC/src/PFAC_reduce_inplace_kernel.cu @@ -1120,7 +1120,7 @@ __global__ void PFAC_reduce_space_driven_device( } int k = 0 ; unsigned int match_pattern ; - MANUAL_EXPAND_8( match_pattern = __ballot( match[k] > 0 ); \ + MANUAL_EXPAND_8( match_pattern = __ballot_sync(__activemask(), match[k] > 0 ); \ match_pattern <<= (31-lane_id); \ acc_pos[k] = __popc(match_pattern); \ if ( 31 == lane_id ){ \ diff --git a/PFAC/src/PFAC_reduce_kernel.cu b/PFAC/src/PFAC_reduce_kernel.cu index 3f8e759..33c8b2b 100644 --- a/PFAC/src/PFAC_reduce_kernel.cu +++ b/PFAC/src/PFAC_reduce_kernel.cu @@ -764,7 +764,7 @@ __global__ void PFAC_reduce_kernel_device( } int k = 0 ; unsigned int match_pattern ; - MANUAL_EXPAND_8( match_pattern = __ballot( match[k] > 0 ); \ + MANUAL_EXPAND_8( match_pattern = __ballot_sync(__activemask(), match[k] > 0 ); \ match_pattern <<= (31-lane_id); \ acc_pos[k] = __popc(match_pattern); \ if ( 31 == lane_id ){ \ @@ -966,7 +966,7 @@ __global__ void PFAC_reduce_kernel_device( [code] #pragma unroll for(int k = 0 ; k < 4*NUM_INTS_PER_THREAD ; k++ ){ - unsigned int match_pattern = __ballot( match[k] > 0 ) ; + unsigned int match_pattern = __ballot_sync(__activemask(), match[k] > 0 ) ; match_pattern <<= (31 - lane_id); acc_pos[k] = __popc(match_pattern) ; if ( 31 == lane_id ){ @@ -980,7 +980,7 @@ __global__ void PFAC_reduce_kernel_device( [code] int k = 0 ; unsigned int match_pattern ; - MANUAL_EXPAND_8( match_pattern = __ballot( match[k] > 0 ); \ + MANUAL_EXPAND_8( match_pattern = __ballot_sync(__activemask(), match[k] > 0 ); \ match_pattern <<= (31-lane_id); \ acc_pos[k] = __popc(match_pattern); \ if ( 31 == lane_id ){ \ diff --git a/PFAC/src/PFAC_reorder_Table.cpp b/PFAC/src/PFAC_reorder_Table.cpp index d3fd999..a8057f0 100644 --- a/PFAC/src/PFAC_reorder_Table.cpp +++ b/PFAC/src/PFAC_reorder_Table.cpp @@ -118,7 +118,7 @@ void printString( char *s, const int n, FILE* fp ) * (6) *max_state_num_ptr = estimation (upper bound) of total states in PFAC DFA * */ -PFAC_status_t parsePatternFile( char *patternfilename, +PFAC_status_t parsePatternFile( const char *patternfilename, char ***rowPtr, char **valPtr, int **patternID_table_ptr, int **patternLen_table_ptr, int *max_state_num_ptr, int *pattern_num_ptr ) {