Skip to content

[BUG FIXES] const qualify PFAC_readPatternFromFile() filename, update to __ballot_sync to support newer GPU architectures #8

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Arleee1 wants to merge 2 commits into
base: master
Choose a base branch
from
Open

[BUG FIXES] const qualify PFAC_readPatternFromFile() filename, update to __ballot_sync to support newer GPU architectures #8

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
1729c64
update all __ballot() calls to __ballot_sync() to run on sm_70 and hi…
Arleee1 Jan 3, 2025
4644033
change filename in PFAC_readPatternFromFile to be const
Arleee1 Feb 28, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion PFAC/include/PFAC.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
2 changes: 1 addition & 1 deletion PFAC/include/PFAC_P.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 ) ;


Expand Down
2 changes: 1 addition & 1 deletion PFAC/src/PFAC.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 ;
Expand Down
2 changes: 1 addition & 1 deletion PFAC/src/PFAC_reduce_inplace_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 ){ \
Expand Down
6 changes: 3 additions & 3 deletions PFAC/src/PFAC_reduce_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 ){ \
Expand Down Expand Up @@ -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 ){
Expand All @@ -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 ){ \
Expand Down
2 changes: 1 addition & 1 deletion PFAC/src/PFAC_reorder_Table.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 )
{
Expand Down