diff options
29 files changed, 274 insertions, 249 deletions
diff --git a/.github/workflows/accelsim.yml b/.github/workflows/accelsim.yml deleted file mode 100644 index f152c0d..0000000 --- a/.github/workflows/accelsim.yml +++ /dev/null @@ -1,39 +0,0 @@ -# Test backend changes with Accel-Sim - -name: Short-Tests-AccelSim - -# Controls when the workflow will run -on: - push: - branches-ignore: - - "gh-readonly-queue**" - merge_group: - pull_request: - - # Allows you to run this workflow manually from the Actions tab - workflow_dispatch: - -# By default regress against accel-sim's dev branch -env: - ACCELSIM_REPO: https://github.com/purdue-aalp/accel-sim-framework-public.git - ACCELSIM_BRANCH: dev - -# A workflow run is made up of one or more jobs that can run sequentially or in parallel -jobs: - build-QV100: - strategy: - matrix: - image_tags: - - ubuntu-24.04-cuda-12.8-minimal - - ubuntu-22.04-cuda-11.7.1-minimal - runs-on: ubuntu-latest - container: - image: ghcr.io/accel-sim/accel-sim-framework:${{ matrix.image_tags }} - - # Steps represent a sequence of tasks that will be executed as part of the job - steps: - # Checks-out your repository under $GITHUB_WORKSPACE, so your job can access it - - uses: actions/checkout@v4 - - name: Run Simulation - run: /bin/bash $GITHUB_WORKSPACE/short-tests-accelsim.sh - diff --git a/.github/workflows/cmake.yml b/.github/workflows/cmake.yml deleted file mode 100644 index c61f5b3..0000000 --- a/.github/workflows/cmake.yml +++ /dev/null @@ -1,45 +0,0 @@ -# Workflow with cmake build system -name: Short-Tests-CMake - -# Controls when the workflow will run -on: - # Triggers the workflow on push or pull request events but only for the mydev branch - push: - branches-ignore: - - "gh-readonly-queue**" - merge_group: - pull_request: - - # Allows you to run this workflow manually from the Actions tab - workflow_dispatch: - -# A workflow run is made up of one or more jobs that can run sequentially or in parallel -jobs: - build: - strategy: - matrix: - app: - - rodinia_2.0-ft - config: - - TITANV - - TITANV-LOCALXBAR - - QV100 - - RTX2060 - - RTX3070 - image_tags: - - ubuntu-24.04-cuda-12.8-minimal - - ubuntu-22.04-cuda-11.7.1-minimal - runs-on: ubuntu-latest - name: Build ${{ matrix.app }} on ${{ matrix.image_tags }} with ${{ matrix.config }} - container: - image: ghcr.io/accel-sim/accel-sim-framework:${{ matrix.image_tags }} - env: - CONFIG: ${{ matrix.config }} - APP: ${{ matrix.app }} - - # Steps represent a sequence of tasks that will be executed as part of the job - steps: - # Checks-out your repository under $GITHUB_WORKSPACE, so your job can access it - - uses: actions/checkout@v4 - - name: Run Simulation - run: /bin/bash $GITHUB_WORKSPACE/short-tests-cmake.sh diff --git a/.github/workflows/format_check.yml b/.github/workflows/format_check.yml new file mode 100644 index 0000000..db94ee0 --- /dev/null +++ b/.github/workflows/format_check.yml @@ -0,0 +1,32 @@ +name: Format Check + +on: + pull_request: + +jobs: + check-format: + runs-on: tgrogers-raid + steps: + - uses: actions/checkout@v4 + - name: Check code formatting + run: | + # Run the formatter to see if it would make any changes + /bin/bash ./format-code.sh + + # Check if there are any uncommitted changes + if git diff --quiet; then + echo "✅ Code is properly formatted" + else + echo "❌ Code formatting check failed. The following files need formatting:" + git diff --name-only + echo "" + echo "Please run the formatter locally and commit the changes:" + echo "/bin/bash ./format-code.sh" + exit 1 + fi + + # skip the runs in PRs. Will run in the merge queue + CI-Success: + runs-on: ubuntu-latest + steps: + - run: echo "CI skipped. Will run in the merge queue"
\ No newline at end of file diff --git a/.github/workflows/main.yml b/.github/workflows/main.yml index 4c12f58..0cb61e4 100644 --- a/.github/workflows/main.yml +++ b/.github/workflows/main.yml @@ -1,6 +1,6 @@ # This is a basic workflow to help you get started with Actions -name: Short-Tests +name: CI Tests # Controls when the workflow will run on: @@ -8,33 +8,48 @@ on: branches-ignore: - "gh-readonly-queue**" merge_group: - pull_request: - # Allows you to run this workflow manually from the Actions tab workflow_dispatch: # A workflow run is made up of one or more jobs that can run sequentially or in parallel jobs: - build-TITANV: - runs-on: ubuntu-latest - container: - image: ghcr.io/accel-sim/accel-sim-framework:ubuntu-24.04-cuda-12.8-minimal - env: - CONFIG: TITANV - - # Steps represent a sequence of tasks that will be executed as part of the job + check-format: + runs-on: tgrogers-raid steps: - # Checks-out your repository under $GITHUB_WORKSPACE, so your job can access it - uses: actions/checkout@v4 - - name: Run Simulation - run: /bin/bash $GITHUB_WORKSPACE/short-tests.sh + + - name: Check code formatting + run: | + # Run the formatter to see if it would make any changes + /bin/bash ./format-code.sh + + # Check if there are any uncommitted changes + if git diff --quiet; then + echo "✅ Code is properly formatted" + else + echo "❌ Code formatting check failed. The following files need formatting:" + git diff --name-only + echo "" + echo "Please run the formatter locally and commit the changes:" + echo "/bin/bash ./format-code.sh" + exit 1 + fi - build-TITANV-LOCALXBAR: + build-Makefile: + # minimal tests - Makefile is deprecated + strategy: + matrix: + config: + - QV100 + image_tags: + - ubuntu-24.04-cuda-12.8-minimal runs-on: ubuntu-latest + needs: check-format + name: build-Makefile container: - image: ghcr.io/accel-sim/accel-sim-framework:ubuntu-24.04-cuda-12.8-minimal + image: ghcr.io/accel-sim/accel-sim-framework:${{ matrix.image_tags }} env: - CONFIG: TITANV-LOCALXBAR + CONFIG: ${{ matrix.config }} # Steps represent a sequence of tasks that will be executed as part of the job steps: @@ -43,63 +58,65 @@ jobs: - name: Run Simulation run: /bin/bash $GITHUB_WORKSPACE/short-tests.sh - build-QV100: + build-CMake: + strategy: + matrix: + app: + - rodinia_2.0-ft + config: + - TITANV + - TITANV-LOCALXBAR + - QV100 + - RTX2060 + - RTX3070 + image_tags: + - ubuntu-24.04-cuda-12.8-minimal + - ubuntu-22.04-cuda-11.7.1-minimal runs-on: ubuntu-latest + needs: check-format + name: build-CMake container: - image: ghcr.io/accel-sim/accel-sim-framework:ubuntu-24.04-cuda-12.8-minimal + image: ghcr.io/accel-sim/accel-sim-framework:${{ matrix.image_tags }} env: - CONFIG: QV100 + CONFIG: ${{ matrix.config }} + APP: ${{ matrix.app }} # Steps represent a sequence of tasks that will be executed as part of the job steps: # Checks-out your repository under $GITHUB_WORKSPACE, so your job can access it - uses: actions/checkout@v4 - name: Run Simulation - run: /bin/bash $GITHUB_WORKSPACE/short-tests.sh + run: /bin/bash $GITHUB_WORKSPACE/short-tests-cmake.sh - build-2060: + build-AccelSim: + strategy: + matrix: + image_tags: + - ubuntu-24.04-cuda-12.8-minimal + - ubuntu-22.04-cuda-11.7.1-minimal runs-on: ubuntu-latest + needs: check-format + name: build-AccelSim container: - image: ghcr.io/accel-sim/accel-sim-framework:ubuntu-24.04-cuda-12.8-minimal + image: ghcr.io/accel-sim/accel-sim-framework:${{ matrix.image_tags }} env: - CONFIG: RTX2060 + ACCELSIM_REPO: https://github.com/accel-sim/accel-sim-framework.git + ACCELSIM_BRANCH: dev # Steps represent a sequence of tasks that will be executed as part of the job steps: # Checks-out your repository under $GITHUB_WORKSPACE, so your job can access it - uses: actions/checkout@v4 - name: Run Simulation - run: /bin/bash $GITHUB_WORKSPACE/short-tests.sh + run: /bin/bash $GITHUB_WORKSPACE/short-tests-accelsim.sh - build-3070: + CI-Success: runs-on: ubuntu-latest - container: - image: ghcr.io/accel-sim/accel-sim-framework:ubuntu-24.04-cuda-12.8-minimal - env: - CONFIG: RTX3070 - - # Steps represent a sequence of tasks that will be executed as part of the job + needs: [check-format, build-Makefile, build-CMake, build-AccelSim] steps: - # Checks-out your repository under $GITHUB_WORKSPACE, so your job can access it - - uses: actions/checkout@v4 - - name: Run Simulation - run: /bin/bash $GITHUB_WORKSPACE/short-tests.sh - #format-code: - # if: github.event_name == 'pull_request' - # runs-on: tgrogers-raid - # needs: [build-TITANV, build-TITANV-LOCALXBAR, build-QV100, build-2060, build-3070] - # steps: - # - uses: actions/checkout@v4 - # with: - # ref: ${{github.event.pull_request.head.ref}} - # repository: ${{github.event.pull_request.head.repo.full_name}} - # ssh-key: '' - - # - name: Run clang-format - # run: | - # git config user.name "purdue-jenkins" - # git config user.email "[email protected]" - # git remote set-url origin [email protected]:${{github.event.pull_request.head.repo.full_name}} - # git remote -v - # /bin/bash ./format-code.sh - # if git status --untracked-files=no | grep -q "nothing to commit"; then echo "No changes to commit."; else git commit -a -m "Automated Format"; git push; fi + - run: | + echo "✅ All CI tests completed successfully!" + echo "Format check: ✅" + echo "Makefile build: ✅" + echo "CMake build: ✅" + echo "AccelSim build: ✅" diff --git a/.github/workflows/sst_integration.yml b/.github/workflows/sst_integration.yml index cb94de4..f491f59 100644 --- a/.github/workflows/sst_integration.yml +++ b/.github/workflows/sst_integration.yml @@ -7,14 +7,12 @@ on: push: branches-ignore: - "gh-readonly-queue**" - pull_request: - - # Allows you to run this workflow manually from the Actions tab workflow_dispatch: # A workflow run is made up of one or more jobs that can run sequentially or in parallel jobs: build-QV100: + name: SST-QV100 runs-on: ubuntu-latest defaults: run: diff --git a/libcuda/cuda_api.h b/libcuda/cuda_api.h index 52f36eb..5a970ba 100644 --- a/libcuda/cuda_api.h +++ b/libcuda/cuda_api.h @@ -2607,12 +2607,12 @@ typedef struct CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS_st { /** * Device that represents the CPU */ -#define CU_DEVICE_CPU ((CUdevice) - 1) +#define CU_DEVICE_CPU ((CUdevice)-1) /** * Device that represents an invalid device */ -#define CU_DEVICE_INVALID ((CUdevice) - 2) +#define CU_DEVICE_INVALID ((CUdevice)-2) /** @} */ /* END CUDA_TYPES */ diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index 3e34641..e620e57 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -35,7 +35,9 @@ struct _cuda_device_id { m_next = NULL; m_gpgpu = gpu; } - struct _cuda_device_id *next() { return m_next; } + struct _cuda_device_id *next() { + return m_next; + } unsigned num_shader() const { return m_gpgpu->get_config().num_shader(); } int num_devices() const { if (m_next == NULL) @@ -156,7 +158,9 @@ class kernel_config { void set_grid_dim(dim3 *d) { m_GridDim = *d; } void set_block_dim(dim3 *d) { m_BlockDim = *d; } gpgpu_ptx_sim_arg_list_t get_args() { return m_args; } - struct CUstream_st *get_stream() { return m_stream; } + struct CUstream_st *get_stream() { + return m_stream; + } private: dim3 m_GridDim; diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 8743fe1..5317c86 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1809,8 +1809,9 @@ cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, case 19: *value = 0; break; - case 20: // cudaDevAttrComputeMode for controlling cudaSetDevice for threads - *value = 0; // Dummy value, should not affect simulation + case 20: // cudaDevAttrComputeMode for controlling cudaSetDevice for + // threads + *value = 0; // Dummy value, should not affect simulation case 21: case 22: case 23: @@ -2433,14 +2434,13 @@ void SST_gpgpusim_numcores_equal_check(unsigned sst_numcores) { /** * @brief For SST to check if kernel launch is blocking - * Future: we will need a better interface to the + * Future: we will need a better interface to the * GPGPU-Sim config for integration with outside * simulators. - * + * */ bool SST_gpgpusim_launch_blocking() { return GPGPU_Context()->the_gpgpusim->g_stream_manager->is_blocking(); - } uint64_t cudaMallocSST(void **devPtr, size_t size) { @@ -3016,7 +3016,8 @@ __host__ cudaError_t CUDARTAPI cudaStreamSynchronizeSST(cudaStream_t stream) { return cudaSuccess; } else { // Otherwise we mark we should wait for default strem to sync - ctx->the_gpgpusim->g_stream_manager->get_stream_zero()->set_request_synchronize(); + ctx->the_gpgpusim->g_stream_manager->get_stream_zero() + ->set_request_synchronize(); return cudaErrorNotReady; } } else { @@ -3101,7 +3102,8 @@ __host__ cudaError_t CUDARTAPI cudaEventSynchronize(cudaEvent_t event) { printf("GPGPU-Sim API: cudaEventSynchronize ** waiting for event\n"); fflush(stdout); CUevent_st *e = (CUevent_st *)event; - while (!e->done()); + while (!e->done()) + ; printf("GPGPU-Sim API: cudaEventSynchronize ** event detected\n"); fflush(stdout); return g_last_cudaError = cudaSuccess; @@ -4105,8 +4107,8 @@ cudaError_t CUDARTAPI cudaSetDeviceFlagsSST(int flags) { // SST's simple stream example relies on this // currently just set it to no-op printf( - "GPGPU-Sim PTX: Execution warning: ignoring call to \"%s ( flag=%p)\"\n", - __my_func__, flags); + "GPGPU-Sim PTX: Execution warning: ignoring call to \"%s ( flag=%p)\"\n", + __my_func__, flags); return cudaSuccess; } diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 8743cc7..e8ddf95 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -216,20 +216,20 @@ new_addr_type line_size_based_tag_func(new_addr_type address, return address & ~(line_size - 1); } -const char *mem_access_type_str(enum mem_access_type access_type){ +const char *mem_access_type_str(enum mem_access_type access_type) { #define MA_TUP_BEGIN(X) static const char *access_type_str[] = { #define MA_TUP(X) #X #define MA_TUP_END(X) \ } \ ; - MEM_ACCESS_TYPE_TUP_DEF + MEM_ACCESS_TYPE_TUP_DEF #undef MA_TUP_BEGIN #undef MA_TUP #undef MA_TUP_END - assert(access_type < NUM_MEM_ACCESS_TYPE); + assert(access_type < NUM_MEM_ACCESS_TYPE); -return access_type_str[access_type]; + return access_type_str[access_type]; } void warp_inst_t::clear_active(const active_mask_t &inactive) { diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index cddf523..06c3082 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -248,7 +248,9 @@ class kernel_info_t { } bool running() const { return m_num_cores_running > 0; } bool done() const { return no_more_ctas_to_run() && !running(); } - class function_info *entry() { return m_kernel_entry; } + class function_info *entry() { + return m_kernel_entry; + } const class function_info *entry() const { return m_kernel_entry; } size_t num_blocks() const { @@ -298,7 +300,9 @@ class kernel_info_t { std::list<class ptx_thread_info *> &active_threads() { return m_active_threads; } - class memory_space *get_param_memory() { return m_param_mem; } + class memory_space *get_param_memory() { + return m_param_mem; + } // The following functions access texture bindings present at the kernel's // launch @@ -605,9 +609,15 @@ class gpgpu_t { void memcpy_from_gpu(void *dst, size_t src_start_addr, size_t count); void memcpy_gpu_to_gpu(size_t dst, size_t src, size_t count); - class memory_space *get_global_memory() { return m_global_mem; } - class memory_space *get_tex_memory() { return m_tex_mem; } - class memory_space *get_surf_memory() { return m_surf_mem; } + class memory_space *get_global_memory() { + return m_global_mem; + } + class memory_space *get_tex_memory() { + return m_tex_mem; + } + class memory_space *get_surf_memory() { + return m_surf_mem; + } void gpgpu_ptx_sim_bindTextureToArray(const struct textureReference *texref, const struct cudaArray *array); @@ -1327,7 +1337,9 @@ class core_t { virtual bool warp_waiting_at_barrier(unsigned warp_id) const = 0; virtual void checkExecutionStatusAndUpdate(warp_inst_t &inst, unsigned t, unsigned tid) = 0; - class gpgpu_sim *get_gpu() { return m_gpu; } + class gpgpu_sim *get_gpu() { + return m_gpu; + } void execute_warp_inst_t(warp_inst_t &inst, unsigned warpId = (unsigned)-1); bool ptx_thread_done(unsigned hw_thread_id) const; virtual void updateSIMTStack(unsigned warpId, warp_inst_t *inst); @@ -1337,7 +1349,9 @@ class core_t { void get_pdom_stack_top_info(unsigned warpId, unsigned *pc, unsigned *rpc) const; kernel_info_t *get_kernel_info() { return m_kernel; } - class ptx_thread_info **get_thread_info() { return m_thread; } + class ptx_thread_info **get_thread_info() { + return m_thread; + } unsigned get_warp_size() const { return m_warp_size; } void and_reduction(unsigned ctaid, unsigned barid, bool value) { reduction_storage[ctaid][barid] &= value; diff --git a/src/accelwattch/core.cc b/src/accelwattch/core.cc index fe11e05..556ed6c 100644 --- a/src/accelwattch/core.cc +++ b/src/accelwattch/core.cc @@ -2195,7 +2195,7 @@ EXECU::EXECU(ParseXML* XML_interface, int ithCore_, bypass.area.set_area(bypass.area.get_area() +fpTagBypass->area.get_area()); }*/ - } /* if (coredynp.core_ty==Inorder) */ + } /* if (coredynp.core_ty==Inorder) */ else { // OOO if (coredynp.scheu_ty == PhysicalRegFile) { /* For physical register based OOO, diff --git a/src/accelwattch/logic.cc b/src/accelwattch/logic.cc index 3d3a1a4..7f40189 100644 --- a/src/accelwattch/logic.cc +++ b/src/accelwattch/logic.cc @@ -656,7 +656,7 @@ FunctionalUnit::FunctionalUnit(ParseXML *XML_interface, int ithCore_, } per_access_energy *= 0.5; // According to ARM data embedded processor has // much lower per acc energy - } /* if (XML->sys.Embedded) */ + } /* if (XML->sys.Embedded) */ else { if (fu_type == FPU) { num_fu = coredynp.num_fpus; diff --git a/src/accelwattch/processor.cc b/src/accelwattch/processor.cc index a86b96d..d5c7cdd 100644 --- a/src/accelwattch/processor.cc +++ b/src/accelwattch/processor.cc @@ -664,22 +664,24 @@ void Processor::displayDeviceType(int device_type_, uint32_t indent) { switch (device_type_) { case 0: - cout << indent_str - << "Device Type= " << "ITRS high performance device type" << endl; + cout << indent_str << "Device Type= " + << "ITRS high performance device type" << endl; break; case 1: - cout << indent_str - << "Device Type= " << "ITRS low standby power device type" << endl; + cout << indent_str << "Device Type= " + << "ITRS low standby power device type" << endl; break; case 2: - cout << indent_str - << "Device Type= " << "ITRS low operating power device type" << endl; + cout << indent_str << "Device Type= " + << "ITRS low operating power device type" << endl; break; case 3: - cout << indent_str << "Device Type= " << "LP-DRAM device type" << endl; + cout << indent_str << "Device Type= " + << "LP-DRAM device type" << endl; break; case 4: - cout << indent_str << "Device Type= " << "COMM-DRAM device type" << endl; + cout << indent_str << "Device Type= " + << "COMM-DRAM device type" << endl; break; default: { cout << indent_str << "Unknown Device Type" << endl; diff --git a/src/accelwattch/xmlParser.cc b/src/accelwattch/xmlParser.cc index 9f01ebe..084a50d 100644 --- a/src/accelwattch/xmlParser.cc +++ b/src/accelwattch/xmlParser.cc @@ -3241,7 +3241,7 @@ XMLSTR XMLParserBase64Tool::encode(unsigned char *inbuf, unsigned int inlen, *(curr++) = base64EncodeTable[j >> 18]; *(curr++) = base64EncodeTable[(j >> 12) & 0x3f]; *(curr++) = base64EncodeTable[(j >> 6) & 0x3f]; - *(curr++) = base64EncodeTable[(j) & 0x3f]; + *(curr++) = base64EncodeTable[(j)&0x3f]; if (formatted) { if (!k) { *(curr++) = _CXML('\n'); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 69d1eb7..14ae504 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2782,7 +2782,9 @@ void print_ptxinfo() { } } -struct gpgpu_ptx_sim_info get_ptxinfo() { return g_ptxinfo; } +struct gpgpu_ptx_sim_info get_ptxinfo() { + return g_ptxinfo; +} std::map<unsigned, const char *> get_duplicate() { return g_duplicate; } diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index b1caf0c..21e1ca0 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -101,8 +101,8 @@ class functionalCoreSim : public core_t { bool *m_warpAtBarrier; }; -#define RECONVERGE_RETURN_PC ((address_type) - 2) -#define NO_BRANCH_DIVERGENCE ((address_type) - 1) +#define RECONVERGE_RETURN_PC ((address_type)-2) +#define NO_BRANCH_DIVERGENCE ((address_type)-1) address_type get_return_pc(void *thd); const char *get_ptxinfo_kname(); void print_ptxinfo(); diff --git a/src/cuda-sim/cuda_device_printf.cc b/src/cuda-sim/cuda_device_printf.cc index 69ebc63..ef780e9 100644 --- a/src/cuda-sim/cuda_device_printf.cc +++ b/src/cuda-sim/cuda_device_printf.cc @@ -65,7 +65,7 @@ void my_cuda_printf(const char *fmtstr, const char *arg_list) { fprintf(fp, buf, *((unsigned *)ptr)); arg_offset += 4; } else if (c == 'f') { - if (acc%2) { + if (acc % 2) { arg_offset += 4; ptr = (void *)&arg_list[arg_offset]; } diff --git a/src/cuda-sim/half.h b/src/cuda-sim/half.h index 67e607b..fab1a22 100644 --- a/src/cuda-sim/half.h +++ b/src/cuda-sim/half.h @@ -846,8 +846,10 @@ uint16 int2half_impl(T value) { bits |= 0x7BFF + (R != std::round_toward_zero);
} else if (value) {
unsigned int m = value, exp = 24;
- for (; m < 0x400; m <<= 1, --exp);
- for (; m > 0x7FF; m >>= 1, ++exp);
+ for (; m < 0x400; m <<= 1, --exp)
+ ;
+ for (; m > 0x7FF; m >>= 1, ++exp)
+ ;
bits |= (exp << 10) + m;
if (exp > 24) {
if (R == std::round_to_nearest)
@@ -1272,7 +1274,8 @@ inline double half2float_impl(uint16 value, double, true_type) { int abs = value & 0x7FFF;
if (abs) {
hi |= 0x3F000000 << static_cast<unsigned>(abs >= 0x7C00);
- for (; abs < 0x400; abs <<= 1, hi -= 0x100000);
+ for (; abs < 0x400; abs <<= 1, hi -= 0x100000)
+ ;
hi += static_cast<uint32>(abs) << 10;
}
uint64 bits = static_cast<uint64>(hi) << 32;
@@ -2113,7 +2116,8 @@ struct functions { static half frexp(half arg, int *exp) {
int m = arg.data_ & 0x7FFF, e = -14;
if (m >= 0x7C00 || !m) return *exp = 0, arg;
- for (; m < 0x400; m <<= 1, --e);
+ for (; m < 0x400; m <<= 1, --e)
+ ;
return *exp = e + (m >> 10),
half(binary, (arg.data_ & 0x8000) | 0x3800 | (m & 0x3FF));
}
@@ -2131,7 +2135,8 @@ struct functions { unsigned int mask = (1 << (25 - e)) - 1, m = arg.data_ & mask;
iptr->data_ = arg.data_ & ~mask;
if (!m) return half(binary, arg.data_ & 0x8000);
- for (; m < 0x400; m <<= 1, --e);
+ for (; m < 0x400; m <<= 1, --e)
+ ;
return half(binary, static_cast<uint16>((arg.data_ & 0x8000) | (e << 10) |
(m & 0x3FF)));
}
@@ -2143,7 +2148,8 @@ struct functions { static half scalbln(half arg, long exp) {
unsigned int m = arg.data_ & 0x7FFF;
if (m >= 0x7C00 || !m) return arg;
- for (; m < 0x400; m <<= 1, --exp);
+ for (; m < 0x400; m <<= 1, --exp)
+ ;
exp += m >> 10;
uint16 value = arg.data_ & 0x8000;
if (exp > 30) {
@@ -2185,7 +2191,8 @@ struct functions { if (abs < 0x7C00) {
int exp = (abs >> 10) - 15;
if (abs < 0x400)
- for (; abs < 0x200; abs <<= 1, --exp);
+ for (; abs < 0x200; abs <<= 1, --exp)
+ ;
return exp;
}
if (abs > 0x7C00) return FP_ILOGBNAN;
@@ -2201,11 +2208,13 @@ struct functions { if (abs < 0x7C00) {
int exp = (abs >> 10) - 15;
if (abs < 0x400)
- for (; abs < 0x200; abs <<= 1, --exp);
+ for (; abs < 0x200; abs <<= 1, --exp)
+ ;
uint16 bits = (exp < 0) << 15;
if (exp) {
unsigned int m = std::abs(exp) << 6, e = 18;
- for (; m < 0x400; m <<= 1, --e);
+ for (; m < 0x400; m <<= 1, --e)
+ ;
bits |= (e << 10) + m;
}
return half(binary, bits);
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc index 9a22d02..d0de1ca 100644 --- a/src/cuda-sim/ptx_sim.cc +++ b/src/cuda-sim/ptx_sim.cc @@ -203,8 +203,8 @@ unsigned ptx_thread_info::get_builtin(int builtin_id, unsigned dim_mod) { return (unsigned)(m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle); case CLOCK64_REG: // Change return value to unsigned long long? - // Currently returns 32-bit unsigned, which may cause truncation for large values. - // GPGPUSim clock is 4 times slower - multiply by 4 + // Currently returns 32-bit unsigned, which may cause truncation for large + // values. GPGPUSim clock is 4 times slower - multiply by 4 return (m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle) * 4; case HALFCLOCK_ID: // GPGPUSim clock is 4 times slower - multiply by 4 diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h index 7128f8e..8eec922 100644 --- a/src/cuda-sim/ptx_sim.h +++ b/src/cuda-sim/ptx_sim.h @@ -340,7 +340,9 @@ class ptx_thread_info { dim3 get_ctaid() const { return m_ctaid; } dim3 get_tid() const { return m_tid; } dim3 get_ntid() const { return m_ntid; } - class gpgpu_sim *get_gpu() { return (gpgpu_sim *)m_gpu; } + class gpgpu_sim *get_gpu() { + return (gpgpu_sim *)m_gpu; + } unsigned get_hw_tid() const { return m_hw_tid; } unsigned get_hw_ctaid() const { return m_hw_ctaid; } unsigned get_hw_wid() const { return m_hw_wid; } diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc index 1e1ad3d..80e20d7 100644 --- a/src/gpgpu-sim/dram.cc +++ b/src/gpgpu-sim/dram.cc @@ -684,9 +684,13 @@ bool dram_t::issue_row_command(int j) { } // if mrq is being serviced by dram, gets popped after CL latency fulfilled -class mem_fetch *dram_t::return_queue_pop() { return returnq->pop(); } +class mem_fetch *dram_t::return_queue_pop() { + return returnq->pop(); +} -class mem_fetch *dram_t::return_queue_top() { return returnq->top(); } +class mem_fetch *dram_t::return_queue_top() { + return returnq->top(); +} void dram_t::print(FILE *simFile) const { unsigned i; diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 68bdca7..98e52f4 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -879,7 +879,7 @@ class sst_gpgpu_sim : public gpgpu_sim { * @param dst_start_addr * @param count */ - void perf_memcpy_to_gpu(size_t dst_start_addr, size_t count) {}; + void perf_memcpy_to_gpu(size_t dst_start_addr, size_t count){}; /** * @brief Check if the SST config matches up with the diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h index 9d164d7..65c9c38 100644 --- a/src/gpgpu-sim/l2cache.h +++ b/src/gpgpu-sim/l2cache.h @@ -106,7 +106,9 @@ class memory_partition_unit { unsigned get_mpid() const { return m_id; } - class gpgpu_sim *get_mgpu() const { return m_gpu; } + class gpgpu_sim *get_mgpu() const { + return m_gpu; + } private: unsigned m_id; diff --git a/src/gpgpu-sim/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc index e4cdaa5..ce4415c 100644 --- a/src/gpgpu-sim/local_interconnect.cc +++ b/src/gpgpu-sim/local_interconnect.cc @@ -190,7 +190,7 @@ void xbar_router::iSLIP_Advance() { input_nodes.insert(i); unsigned out_node = in_buffers[i].front().output_deviceID; - if(destination_set.find(out_node) != destination_set.end()) { + if (destination_set.find(out_node) != destination_set.end()) { conflict_sub++; } destination_set.insert(out_node); @@ -207,40 +207,40 @@ void xbar_router::iSLIP_Advance() { for (auto dest : destination_set) { if (Has_Buffer_Out(dest, 1)) { unsigned start_node = next_node[dest]; - auto it = std::upper_bound(input_nodes.begin(), input_nodes.end(), - start_node); + auto it = + std::upper_bound(input_nodes.begin(), input_nodes.end(), start_node); for (unsigned j = 0; j < input_nodes.size(); j++, it++) { if (it == input_nodes.end()) { it = input_nodes.begin(); } unsigned node_id = *it; assert(!in_buffers[node_id].empty()); - Packet _packet = in_buffers[node_id].front(); - if (_packet.output_deviceID == dest) { - out_buffers[_packet.output_deviceID].push(_packet); - in_buffers[node_id].pop(); - input_nodes.erase(node_id); //can only be used once - if (verbose) - printf("%d : cycle %llu : send req from %d to %d\n", m_id, cycles, - node_id, dest - _n_shader); - if (grant_cycles_count == 1) - next_node[dest] = (++node_id % total_nodes); - if (verbose) { - for (unsigned k = j + 1; k < total_nodes; ++k) { - unsigned node_id2 = (k + next_node[dest]) % total_nodes; - if (!in_buffers[node_id2].empty()) { - Packet _packet2 = in_buffers[node_id2].front(); + Packet _packet = in_buffers[node_id].front(); + if (_packet.output_deviceID == dest) { + out_buffers[_packet.output_deviceID].push(_packet); + in_buffers[node_id].pop(); + input_nodes.erase(node_id); // can only be used once + if (verbose) + printf("%d : cycle %llu : send req from %d to %d\n", m_id, cycles, + node_id, dest - _n_shader); + if (grant_cycles_count == 1) + next_node[dest] = (++node_id % total_nodes); + if (verbose) { + for (unsigned k = j + 1; k < total_nodes; ++k) { + unsigned node_id2 = (k + next_node[dest]) % total_nodes; + if (!in_buffers[node_id2].empty()) { + Packet _packet2 = in_buffers[node_id2].front(); - if (_packet2.output_deviceID == dest) - printf("%d : cycle %llu : cannot send req from %d to %d\n", - m_id, cycles, node_id2, dest - _n_shader); - } + if (_packet2.output_deviceID == dest) + printf("%d : cycle %llu : cannot send req from %d to %d\n", + m_id, cycles, node_id2, dest - _n_shader); } } - - reqs++; - break; } + + reqs++; + break; + } } } else { out_buffer_full++; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 1ea011e..730cb6d 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -891,7 +891,7 @@ void shader_core_ctx::decode() { // decode 1 or 2 instructions and place them into ibuffer address_type pc = m_inst_fetch_buffer.m_pc; const warp_inst_t *pI1 = get_next_inst(m_inst_fetch_buffer.m_warp_id, pc); - if (pI1) { + if (pI1) { m_warp[m_inst_fetch_buffer.m_warp_id]->ibuffer_fill(0, pI1); m_warp[m_inst_fetch_buffer.m_warp_id]->inc_inst_in_pipeline(); m_stats->m_num_decoded_insn[m_sid]++; diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index b90ce2f..5922d24 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -189,7 +189,7 @@ class shd_warp_t { m_active_threads.reset(lane); n_completed++; } - bool test_active(unsigned lane) {return m_active_threads.test(lane);} + bool test_active(unsigned lane) { return m_active_threads.test(lane); } void set_last_fetch(unsigned long long sim_cycle) { m_last_fetch = sim_cycle; @@ -274,7 +274,9 @@ class shd_warp_t { unsigned get_dynamic_warp_id() const { return m_dynamic_warp_id; } unsigned get_warp_id() const { return m_warp_id; } - class shader_core_ctx *get_shader() { return m_shader; } + class shader_core_ctx *get_shader() { + return m_shader; + } private: static const unsigned IBUFFER_SIZE = 2; diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index be49229..4462dc3 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -58,7 +58,8 @@ class stream_manager *g_stream_manager() { extern void SST_callback_cudaThreadSynchronize_done(); extern void SST_callback_cudaStreamSynchronize_done(cudaStream_t stream); __attribute__((weak)) void SST_callback_cudaThreadSynchronize_done() {} -__attribute__((weak)) void SST_callback_cudaStreamSynchronize_done(cudaStream_t stream) {} +__attribute__((weak)) void SST_callback_cudaStreamSynchronize_done( + cudaStream_t stream) {} void *gpgpu_sim_thread_sequential(void *ctx_ptr) { gpgpu_context *ctx = (gpgpu_context *)ctx_ptr; @@ -102,7 +103,8 @@ void *gpgpu_sim_thread_concurrent(void *ctx_ptr) { fflush(stdout); } while (ctx->the_gpgpusim->g_stream_manager->empty_protected() && - !ctx->the_gpgpusim->g_sim_done); + !ctx->the_gpgpusim->g_sim_done) + ; if (g_debug_execution >= 3) { printf("GPGPU-Sim: ** START simulation thread (detected work) **\n"); ctx->the_gpgpusim->g_stream_manager->print(stdout); @@ -191,7 +193,8 @@ bool SST_Cycle() { // Check if Synchronize is done when SST previously requested // cudaThreadSynchronize if (GPGPU_Context()->requested_synchronize && - ((g_stream_manager()->empty_protected() && !GPGPUsim_ctx_ptr()->g_sim_active) || + ((g_stream_manager()->empty_protected() && + !GPGPUsim_ctx_ptr()->g_sim_active) || GPGPUsim_ctx_ptr()->g_sim_done)) { SST_callback_cudaThreadSynchronize_done(); GPGPU_Context()->requested_synchronize = false; @@ -199,19 +202,20 @@ bool SST_Cycle() { // Polling to check for each stream if it is marked for requested with sync if (g_stream_manager()->get_stream_zero()->requested_synchronize() && - ((g_stream_manager()->empty_protected() && !GPGPUsim_ctx_ptr()->g_sim_active) || - GPGPUsim_ctx_ptr()->g_sim_done)) { + ((g_stream_manager()->empty_protected() && + !GPGPUsim_ctx_ptr()->g_sim_active) || + GPGPUsim_ctx_ptr()->g_sim_done)) { SST_callback_cudaStreamSynchronize_done(0); g_stream_manager()->get_stream_zero()->reset_request_synchronize(); } // Iterate through each stream to check if SST is waiting on // it and it does not have any operation - std::list<CUstream_st *>& streams = g_stream_manager()->get_concurrent_streams(); + std::list<CUstream_st *> &streams = + g_stream_manager()->get_concurrent_streams(); for (auto it = streams.begin(); it != streams.end(); it++) { CUstream_st *stream = *it; - if (stream->requested_synchronize() && - stream->empty()) { + if (stream->requested_synchronize() && stream->empty()) { // This stream is ready SST_callback_cudaStreamSynchronize_done(stream); stream->reset_request_synchronize(); diff --git a/src/stream_manager.cc b/src/stream_manager.cc index d43964a..ab5b74f 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -34,21 +34,30 @@ unsigned CUstream_st::sm_next_stream_uid = 0; -// SST memcpy callbacks, called after a stream operation is done via record_next_done() -extern void SST_callback_memcpy_H2D_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream); -extern void SST_callback_memcpy_D2H_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream); +// SST memcpy callbacks, called after a stream operation is done via +// record_next_done() +extern void SST_callback_memcpy_H2D_done(uint64_t dst, uint64_t src, + size_t count, cudaStream_t stream); +extern void SST_callback_memcpy_D2H_done(uint64_t dst, uint64_t src, + size_t count, cudaStream_t stream); extern void SST_callback_memcpy_to_symbol_done(); extern void SST_callback_memcpy_from_symbol_done(); extern void SST_callback_cudaEventSynchronize_done(cudaEvent_t event); extern void SST_callback_kernel_done(cudaStream_t stream); -__attribute__((weak)) void SST_callback_memcpy_H2D_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream) {} -__attribute__((weak)) void SST_callback_memcpy_D2H_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream) {} +__attribute__((weak)) void SST_callback_memcpy_H2D_done(uint64_t dst, + uint64_t src, + size_t count, + cudaStream_t stream) {} +__attribute__((weak)) void SST_callback_memcpy_D2H_done(uint64_t dst, + uint64_t src, + size_t count, + cudaStream_t stream) {} __attribute__((weak)) void SST_callback_memcpy_to_symbol_done() {} __attribute__((weak)) void SST_callback_memcpy_from_symbol_done() {} -__attribute__((weak)) void SST_callback_cudaEventSynchronize_done(cudaEvent_t event); +__attribute__((weak)) void SST_callback_cudaEventSynchronize_done( + cudaEvent_t event); __attribute__((weak)) void SST_callback_kernel_done(cudaStream_t stream); - CUstream_st::CUstream_st() { m_pending = false; m_uid = sm_next_stream_uid++; @@ -79,9 +88,7 @@ void CUstream_st::synchronize() { } while (!done); } -bool CUstream_st::synchronize_check() { - return m_operations.empty(); -} +bool CUstream_st::synchronize_check() { return m_operations.empty(); } void CUstream_st::push(const stream_operation &op) { // called by host thread @@ -142,14 +149,19 @@ bool stream_operation::do_operation(gpgpu_sim *gpu) { gpu->memcpy_to_gpu(m_device_address_dst, m_host_address_src, m_cnt); m_stream->record_next_done(); if (gpu->is_SST_mode()) { - SST_callback_memcpy_H2D_done((uint64_t) m_device_address_dst, (uint64_t) m_host_address_src, m_cnt, m_stream->is_stream_zero_stream() ? 0 : m_stream); + SST_callback_memcpy_H2D_done( + (uint64_t)m_device_address_dst, (uint64_t)m_host_address_src, m_cnt, + m_stream->is_stream_zero_stream() ? 0 : m_stream); } break; case stream_memcpy_device_to_host: if (g_debug_execution >= 3) printf("memcpy device-to-host\n"); gpu->memcpy_from_gpu(m_host_address_dst, m_device_address_src, m_cnt); m_stream->record_next_done(); - if (gpu->is_SST_mode()) SST_callback_memcpy_D2H_done((uint64_t) m_host_address_dst, (uint64_t) m_device_address_src, m_cnt, m_stream->is_stream_zero_stream() ? 0 : m_stream); + if (gpu->is_SST_mode()) + SST_callback_memcpy_D2H_done( + (uint64_t)m_host_address_dst, (uint64_t)m_device_address_src, m_cnt, + m_stream->is_stream_zero_stream() ? 0 : m_stream); break; case stream_memcpy_device_to_device: if (g_debug_execution >= 3) printf("memcpy device-to-device\n"); @@ -419,7 +431,8 @@ void stream_manager::add_stream(struct CUstream_st *stream) { void stream_manager::destroy_stream(CUstream_st *stream) { // called by host thread pthread_mutex_lock(&m_lock); - while (!stream->empty()); + while (!stream->empty()) + ; std::list<CUstream_st *>::iterator s; for (s = m_streams.begin(); s != m_streams.end(); s++) { if (*s == stream) { diff --git a/src/stream_manager.h b/src/stream_manager.h index 55cfb8d..c6d2fc3 100644 --- a/src/stream_manager.h +++ b/src/stream_manager.h @@ -82,7 +82,7 @@ struct CUevent_st { unsigned int m_issued; time_t m_wallclock; double m_gpu_tot_sim_cycle; - + // SST related bool m_requested_synchronize = false; @@ -206,7 +206,9 @@ class stream_operation { kernel_info_t *get_kernel() { return m_kernel; } bool do_operation(gpgpu_sim *gpu); void print(FILE *fp) const; - struct CUstream_st *get_stream() { return m_stream; } + struct CUstream_st *get_stream() { + return m_stream; + } void set_stream(CUstream_st *stream) { m_stream = stream; } private: @@ -284,7 +286,7 @@ class stream_manager { unsigned size() { return m_streams.size(); }; bool is_blocking() { return m_cuda_launch_blocking; }; CUstream_st *get_stream_zero() { return &m_stream_zero; }; - std::list<CUstream_st *>& get_concurrent_streams() { return m_streams; }; + std::list<CUstream_st *> &get_concurrent_streams() { return m_streams; }; private: void print_impl(FILE *fp); |