From 1729c64d1181b629843928c9f7301ecccdb157be Mon Sep 17 00:00:00 2001 From: Ethan Ermovick Date: Fri, 3 Jan 2025 17:21:07 -0500 Subject: [PATCH 1/2] update all __ballot() calls to __ballot_sync() to run on sm_70 and higher --- PFAC/src/PFAC_reduce_inplace_kernel.cu | 2 +- PFAC/src/PFAC_reduce_kernel.cu | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) 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 ){ \ From 464403316ceb7773e515b52a2818675745fd6232 Mon Sep 17 00:00:00 2001 From: Ethan Ermovick Date: Thu, 27 Feb 2025 23:27:07 -0500 Subject: [PATCH 2/2] change filename in PFAC_readPatternFromFile to be const --- PFAC/include/PFAC.h | 2 +- PFAC/include/PFAC_P.h | 2 +- PFAC/src/PFAC.cpp | 2 +- PFAC/src/PFAC_reorder_Table.cpp | 2 +- 4 files changed, 4 insertions(+), 4 deletions(-) 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_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 ) {