Skip to content

Commit 287fe0a

Browse files
tgrogerstgrogers
authored
A bunch of maintenance fixes, the largest of which is getting the PTX simulation to work with CUDA 12. (gpgpu-sim#95)
* Fixing the formatter to always use a consistent format and running it on the codebase * Update linux-so-version.txt * Update Makefile * A couple of unnecessary files that are lingering around * Support CUDA 12 * Getting the PTX simulations to work with CUDA 12. The issue is that ptxas added more information (number of barriers and compile time). We have to parse these or lexx/yacc fail. * Update ptxinfo.l debug MACRO was ineffective * Update gpgpusim_check.cmake Update to make the CUDA version print a warning, not an error and updating the print to be more reflective of what the actual problem is.
1 parent 7934dfe commit 287fe0a

34 files changed

+120
-169
lines changed

‎Makefile‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -169,6 +169,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
169169
if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.0; fi
170170
if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.1; fi
171171
if [ ! -f $(SIM_LIB_DIR)/libcudart.so.11.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.11.0; fi
172+
if [ ! -f $(SIM_LIB_DIR)/libcudart.so.12 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.12; fi
172173
if [ ! -f $(SIM_LIB_DIR)/libcudart_mod.so ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart_mod.so; fi
173174

174175
$(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib

‎bitbucket-pipelines.yml‎

Lines changed: 0 additions & 15 deletions
This file was deleted.

‎format-code.sh‎

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,13 @@
11
# This bash script formats GPGPU-Sim using clang-format
22
THIS_DIR="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )"
33
echo "Running clang-format on $THIS_DIR"
4-
clang-format -i ${THIS_DIR}/libcuda/*.h
5-
clang-format -i ${THIS_DIR}/libcuda/*.cc
6-
clang-format -i ${THIS_DIR}/src/*.h
7-
clang-format -i ${THIS_DIR}/src/*.cc
8-
clang-format -i ${THIS_DIR}/src/gpgpu-sim/*.h
9-
clang-format -i ${THIS_DIR}/src/gpgpu-sim/*.cc
10-
clang-format -i ${THIS_DIR}/src/cuda-sim/*.h
11-
clang-format -i ${THIS_DIR}/src/cuda-sim/*.cc
12-
clang-format -i ${THIS_DIR}/src/accelwattch/*.h
13-
clang-format -i ${THIS_DIR}/src/accelwattch/*.cc
4+
clang-format -i ${THIS_DIR}/libcuda/*.h --style=file:${THIS_DIR}/.clang-format
5+
clang-format -i ${THIS_DIR}/libcuda/*.cc --style=file:${THIS_DIR}/.clang-format
6+
clang-format -i ${THIS_DIR}/src/*.h --style=file:${THIS_DIR}/.clang-format
7+
clang-format -i ${THIS_DIR}/src/*.cc --style=file:${THIS_DIR}/.clang-format
8+
clang-format -i ${THIS_DIR}/src/gpgpu-sim/*.h --style=file:${THIS_DIR}/.clang-format
9+
clang-format -i ${THIS_DIR}/src/gpgpu-sim/*.cc --style=file:${THIS_DIR}/.clang-format
10+
clang-format -i ${THIS_DIR}/src/cuda-sim/*.h --style=file:${THIS_DIR}/.clang-format
11+
clang-format -i ${THIS_DIR}/src/cuda-sim/*.cc --style=file:${THIS_DIR}/.clang-format
12+
clang-format -i ${THIS_DIR}/src/accelwattch/*.h --style=file:${THIS_DIR}/.clang-format
13+
clang-format -i ${THIS_DIR}/src/accelwattch/*.cc --style=file:${THIS_DIR}/.clang-format

‎gpgpusim_check.cmake‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -63,8 +63,8 @@ else()
6363
message(CHECK_PASS "${CUDAToolkit_NVCC_EXECUTABLE}")
6464
message(CHECK_START "Checking CUDA compiler version")
6565
message(CHECK_PASS "${CUDAToolkit_VERSION}")
66-
if((CUDAToolkit_VERSION VERSION_LESS 2.0.3) OR (CUDAToolkit_VERSION VERSION_GREATER 11.10.0))
67-
message(FATAL_ERROR "GPGPU-Sim ${CMAKE_PROJECT_VERSION} not tested with CUDA version ${CUDAToolkit_VERSION} (please see README)")
66+
if((CUDAToolkit_VERSION VERSION_LESS 2.0.3) OR (CUDAToolkit_VERSION VERSION_GREATER 13.0.0))
67+
message(WARNING "GPGPU-Sim not tested with CUDA version ${CUDAToolkit_VERSION} (please see README)")
6868
endif()
6969
endif()
7070

@@ -132,4 +132,4 @@ list(POP_BACK CMAKE_MESSAGE_INDENT)
132132
message(CHECK_PASS "done")
133133
message(STATUS "Be sure to run 'source setup' "
134134
"before you run CUDA program with GPGPU-Sim or building with external "
135-
"simulator like SST")
135+
"simulator like SST")

‎libcuda/cuda_api.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2607,12 +2607,12 @@ typedef struct CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st {
26072607
/**
26082608
* Device that represents the CPU
26092609
*/
2610-
#define CU_DEVICE_CPU ((CUdevice)-1)
2610+
#define CU_DEVICE_CPU ((CUdevice) - 1)
26112611

26122612
/**
26132613
* Device that represents an invalid device
26142614
*/
2615-
#define CU_DEVICE_INVALID ((CUdevice)-2)
2615+
#define CU_DEVICE_INVALID ((CUdevice) - 2)
26162616

26172617
/** @} */ /* END CUDA_TYPES */
26182618

‎libcuda/cuda_api_object.h‎

Lines changed: 2 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -35,9 +35,7 @@ struct _cuda_device_id {
3535
m_next = NULL;
3636
m_gpgpu = gpu;
3737
}
38-
struct _cuda_device_id *next() {
39-
return m_next;
40-
}
38+
struct _cuda_device_id *next() { return m_next; }
4139
unsigned num_shader() const { return m_gpgpu->get_config().num_shader(); }
4240
int num_devices() const {
4341
if (m_next == NULL)
@@ -158,9 +156,7 @@ class kernel_config {
158156
void set_grid_dim(dim3 *d) { m_GridDim = *d; }
159157
void set_block_dim(dim3 *d) { m_BlockDim = *d; }
160158
gpgpu_ptx_sim_arg_list_t get_args() { return m_args; }
161-
struct CUstream_st *get_stream() {
162-
return m_stream;
163-
}
159+
struct CUstream_st *get_stream() { return m_stream; }
164160

165161
private:
166162
dim3 m_GridDim;

‎libcuda/cuda_runtime_api.cc‎

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -3046,8 +3046,7 @@ __host__ cudaError_t CUDARTAPI cudaEventSynchronize(cudaEvent_t event) {
30463046
printf("GPGPU-Sim API: cudaEventSynchronize ** waiting for event\n");
30473047
fflush(stdout);
30483048
CUevent_st *e = (CUevent_st *)event;
3049-
while (!e->done())
3050-
;
3049+
while (!e->done());
30513050
printf("GPGPU-Sim API: cudaEventSynchronize ** event detected\n");
30523051
fflush(stdout);
30533052
return g_last_cudaError = cudaSuccess;
@@ -3171,7 +3170,7 @@ __host__ cudaError_t CUDARTAPI cudaGetExportTable(
31713170
* *
31723171
*******************************************************************************/
31733172

3174-
//#include "../../cuobjdump_to_ptxplus/cuobjdump_parser.h"
3173+
// #include "../../cuobjdump_to_ptxplus/cuobjdump_parser.h"
31753174

31763175
// extracts all ptx files from binary and dumps into
31773176
// prog_name.unique_no.sm_<>.ptx files
@@ -4068,9 +4067,9 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit,
40684067
return g_last_cudaError = cudaSuccess;
40694068
}
40704069

4071-
//#if CUDART_VERSION >= 9000
4070+
// #if CUDART_VERSION >= 9000
40724071
//__host__ cudaError_t cudaFuncSetAttribute ( const void* func, enum
4073-
// cudaFuncAttribute attr, int value ) {
4072+
// cudaFuncAttribute attr, int value ) {
40744073

40754074
// ignore this Attribute for now, and the default is that carveout =
40764075
// cudaSharedmemCarveoutDefault; // (-1)

‎linux-so-version.txt‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,5 +10,7 @@ libcudart.so.10.1{
1010
};
1111
libcudart.so.11.0{
1212
};
13+
libcudart.so.12{
14+
};
1315
libcuda.so.1{
1416
};

‎new‎

Lines changed: 0 additions & 7 deletions
This file was deleted.

‎src/abstract_hardware_model.cc‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -216,20 +216,20 @@ new_addr_type line_size_based_tag_func(new_addr_type address,
216216
return address & ~(line_size - 1);
217217
}
218218

219-
const char *mem_access_type_str(enum mem_access_type access_type) {
219+
const char *mem_access_type_str(enum mem_access_type access_type){
220220
#define MA_TUP_BEGIN(X) static const char *access_type_str[] = {
221221
#define MA_TUP(X) #X
222222
#define MA_TUP_END(X) \
223223
} \
224224
;
225-
MEM_ACCESS_TYPE_TUP_DEF
225+
MEM_ACCESS_TYPE_TUP_DEF
226226
#undef MA_TUP_BEGIN
227227
#undef MA_TUP
228228
#undef MA_TUP_END
229229

230-
assert(access_type < NUM_MEM_ACCESS_TYPE);
230+
assert(access_type < NUM_MEM_ACCESS_TYPE);
231231

232-
return access_type_str[access_type];
232+
return access_type_str[access_type];
233233
}
234234

235235
void warp_inst_t::clear_active(const active_mask_t &inactive) {

0 commit comments

Comments
 (0)